Provable PCR extensions

Abstract

The present disclosure relates to provable PCR extensions. Processors, systems, methods, and computer program products for detecting unauthorized modifications to security information in a security device. In at least one embodiment, a processor includes one or more circuits that use information stored in one or more storage locations of one or more security devices to detect whether the security devices have been restarted.

Description

TECHNICAL FIELD

[0001] At least one embodiment relates to computer security. For example, at least one embodiment relates to protecting a security device in a computer from tampering. BACKGROUND

[0002] Security devices (e.g., Trusted Platform Modules (TPMs)) store information that records the state of computer firmware and software. However, this information is vulnerable to tampering by attacks. Techniques for detecting these attacks can be improved. BRIEF DESCRIPTION OF DRAWINGS

[0003] Figure 1 An exemplary system for detecting unauthorized PCR extensions is shown, in accordance with at least one embodiment;

[0004] Figure 2 A process for creating a canary object in a TPM is shown, in accordance with at least one embodiment;

[0005] Figure 3 A secure cryptographic binding between a CPU 306 and a TPM 308 in a computer 302 is shown, in accordance with at least one embodiment;

[0006] Figure 4 is a sequence diagram showing operations performed during a secure boot of a computer, in accordance with at least one embodiment;

[0007] Figure 5A and Figure 5B is a block diagram showing interactions between a CPU and a TPM acting as an SPDM requester and an SPDM responder, respectively, in accordance with at least one embodiment;

[0008] Figure 6 is a flow diagram showing a process for verifying a TPM and a CPU, in accordance with an embodiment;

[0009] Figure 7 is a flow diagram showing a process for creating an attestable PCR extension mechanism, in accordance with at least one embodiment;

[0010] Figure 8 is a flow diagram showing a process for detecting unauthorized modifications by an in-the-field attacker, in accordance with at least one embodiment;

[0011] Figure 9 An exemplary data center is shown, in accordance with at least one embodiment;

[0012] Figure 10 A processing system is shown, in accordance with at least one embodiment;

[0013] Figure 11 A computer system, in accordance with at least one embodiment, is shown;

[0014] Figure 12 A system, in accordance with at least one embodiment, is shown;

[0015] Figure 13 An exemplary integrated circuit, in accordance with at least one embodiment, is shown;

[0016] Figure 14 A computing system, in accordance with at least one embodiment, is shown;

[0017] Figure 15 An APU, in accordance with at least one embodiment, is shown;

[0018] Figure 16 A CPU, in accordance with at least one embodiment, is shown;

[0019] Figure 17 An exemplary accelerator integration slice, in accordance with at least one embodiment, is shown;

[0020] Figure 18A and Figure 18B An exemplary graphics processor, in accordance with at least one embodiment, is shown;

[0021] Figure 19A A graphics core, in accordance with at least one embodiment, is shown;

[0022] Figure 19B A GPGPU, in accordance with at least one embodiment, is shown;

[0023] Figure 20A A parallel processor, in accordance with at least one embodiment, is shown;

[0024] Figure 20B A processing cluster, in accordance with at least one embodiment, is shown;

[0025] Figure 20C A graphics multiprocessor, in accordance with at least one embodiment, is shown;

[0026] Figure 21 A graphics processor, in accordance with at least one embodiment, is shown;

[0027] Figure 22 A processor, in accordance with at least one embodiment, is shown;

[0028] Figure 23 A processor, in accordance with at least one embodiment, is shown;

[0029] Figure 24 A graphics processor core, in accordance with at least one embodiment, is shown;

[0030] Figure 25 A PPU according to at least one embodiment is shown;

[0031] Figure 26 A GPC according to at least one embodiment is shown;

[0032] Figure 27 A streaming multiprocessor according to at least one embodiment is illustrated;

[0033] Figure 28 A software stack of a programming platform according to at least one embodiment is shown;

[0034] Figure 29 The illustration shows an embodiment according to at least one of the embodiments. Figure 28 The CUDA implementation of the software stack;

[0035] Figure 30 The illustration shows an embodiment according to at least one of the embodiments. Figure 28 The ROCm implementation of the software stack;

[0036] Figure 31 The illustration shows an embodiment according to at least one of the embodiments. Figure 28 The OpenCL implementation of the software stack;

[0037] Figure 32 Software supported by a programming platform according to at least one embodiment is shown;

[0038] Figure 33 The illustration shows an embodiment of at least one of the following: Figures 28-31 Compiled code executed on the programming platform;

[0039] Figure 34 The illustration shows an embodiment of at least one of the following: Figures 28-31 More detailed compiled code executed on the programming platform;

[0040] Figure 35 This illustrates the transformation of source code before compilation, according to at least one embodiment;

[0041] Figure 36A A system configured to compile and execute CUDA source code using different types of processing units, according to at least one embodiment, is shown;

[0042] Figure 36B The diagram illustrates a configuration, according to at least one embodiment, to compile and execute using a CPU and a CUDA-enabled GPU. Figure 36A The system of CUDA source code;

[0043] Figure 36C The diagram illustrates a configuration, according to at least one embodiment, to compile and execute using a CPU and a GPU with CUDA disabled.Figure 36A a system of CUDA source code;

[0044] Figure 37 shows an exemplary kernel converted by a CUDA to HIP conversion tool of Figure 36C

[0045] Figure 38 shows in more detail a CUDA-enabled GPU of Figure 36C

[0046] Figure 39 shows how threads of an exemplary CUDA grid are mapped to different compute units of Figure 38

[0047] Figure 40 shows how to migrate existing CUDA code to data-parallel C++ code, according to at least one embodiment; and

[0048] Figure 41 shows components of a system for accessing large language models, according to at least one embodiment.DETAILED DESCRIPTION

[0049] In the following description, numerous specific details are set forth to provide a more thorough understanding of at least one embodiment. However, it will be apparent to one skilled in the art that the inventive concept can be practiced without one or more of these specific details.

[0050] In at least one embodiment, a computer has a secure device that stores information that can be read to detect whether a hacker has not modified firmware / software of the computer. In at least one embodiment, this secure device is referred to as a Trusted Platform Module (TPM). In at least one embodiment, the information includes Platform Configuration Register (PCR) values, which are hash values stored in registers in the secure device. In at least one embodiment, if an onsite hacker has access to a bus between a processor and the secure device, the onsite hacker can use a hacking device to intercept information transmitted on the bus and use the intercepted information to cause the secure device to reboot, and during the reboot, emulate the processor to modify these PCR values (e.g., by replaying its own version of the PCR values to cause these values to be stored in the secure device), where the modified PCR values deceive security monitoring software so that it cannot detect that firmware / software of the computer has been modified.

[0051] ​​​In at least one embodiment, described herein are processors, systems, methods, and computer program products for detecting an unauthorized modification to a PCR value in a secure device. In at least one embodiment, a CPU of a computer stores a canary object, such as a sequence of random bytes or a random number, in volatile memory of the secure device during a boot time period. In at least one embodiment, since a live hacker would need to reboot the secure device to initiate such an attack, and since the volatile memory is reset during a reboot, the attack results in the canary object being erased. In at least one embodiment, a security program (e.g., a remote attester) can use the absence of the canary object as an indication that an attack has occurred, and then take mitigating measures, such as by shutting down the computer, sounding an alarm, erasing a key, etc.

[0052] In at least one embodiment, a processor includes one or more circuits to detect whether a secure device has been rebooted using information stored in one or more storage locations of the secure device. In at least one embodiment, the information is an object. In at least one embodiment, the object is a canary object. In at least one embodiment, the object is one of a sequence of random bytes, a string, a digital file appended with a unique identifier, or a unique identifier. In at least one embodiment, the one or more storage locations are volatile memory locations in the secure device, such as a TPM. In at least one embodiment, the processor is a CPU of the computer.

[0053] In at least one embodiment, the canary object on the TPM is created in a manner that a security program can recognize that it was created by the CPU. In at least one embodiment, the CPU stores a copy of the canary object in secure storage on the CPU. In at least one embodiment, both the copy on the CPU and the canary object on the TPM include the same unique identifier. In at least one embodiment, the unique identifier is a specific, one-of-a-kind label or number. In at least one embodiment, the identifier enables to specifically and distinctly identify the canary object to allow a security program to attest that the canary object was indeed created by the CPU, and that the canary object was not tampered with or replaced by another canary object created by an attacker after a reboot of the TPM or the CPU.

[0054] In at least one embodiment, for example, when a live attacker resets the TPM to erase the canary object created by the CPU and subsequently creates a different canary object, the newly created canary object will not have such a unique identifier. In at least one embodiment, the security program can identify the newly created canary object as a fraudulent object by comparing the newly created canary object to the copy canary object stored on the CPU.

[0055] In at least one embodiment, an authorization policy and configuration is established to only allow the CPU to extend or otherwise update sensitive information (e.g., PCR values) of the TPM without requiring a reboot of the TPM. In at least one embodiment, during manufacturing time of a computer, both the CPU and TPM of the computer are provisioned with special keys or certificates that can be used to prove who they are so that they can recognize and trust each other and establish a secure connection between the CPU and the TPM.

[0056] In at least one embodiment, during boot of the computer, the CPU verifies the integrity of the TPM to ensure that it is the correct device associated with that CPU and has not been tampered with. After this verification, the CPU initiates a secure communication session with the TPM to protect all subsequent interactions. Additionally, a setting is configured to only allow the CPU to modify PCR values and only within this secure session.

[0057] In at least one embodiment, a live attacker cannot modify sensitive information in the TPM without rebooting the TPM, and rebooting the TPM will erase the original canary object created by the CPU, which allows the security program to detect that the TPM has been rebooted.

[0058] In at least one embodiment, this method of detecting unauthorized extensions (modifications) of sensitive information does not require complex setup and periodic updates, making it more resource-efficient, easier to implement, and more reliable than using attestation reports for such detection.

[0059] Figure 1 An exemplary system 100 for detecting unauthorized PCR extensions is shown in accordance with at least one embodiment. In at least one embodiment, system 100 includes a computer 102, such as computer system 1100 described in Figure 11 “System Architecture,” pages 1-12, which is incorporated by reference herein in its entirety. In at least one embodiment, computer system 102 includes a central processing unit (CPU) 106, a trusted platform module (TPM) 108, boot software 128, and a verifier 130.

[0060] In at least one embodiment, CPU 106 corresponds to CPU 1600 described in Figure 16 In at least one embodiment, TPM 108 is a hardware-based security device that is a specialized processor that enhances the security of a host system (e.g., computer 102) by integrating hardware-based security measures. In at least one embodiment, TPM 108 includes PCRs 110, which are specialized registers for securely storing hash values (e.g., cryptographic hashes representing the state of various components of computer 102). In at least one embodiment, each hash value is derived from a measurement (e.g., checksum or hash) of one of the components on computer 102 (e.g., BIOS, bootloader, operating system, and installed applications).

[0061] In at least one embodiment, PCRs 110 provide secure storage for these hash values within TPM 108. In at least one embodiment, these values (PCR values) record the sequence of software and configuration states that computer 102 undergoes during a reboot. In at least one embodiment, unauthorized software or users cannot directly change these values using commands without rebooting the TPM.

[0062] In at least one embodiment, at system startup, CPU 106 measures a hash of a software component using a hash function (e.g., SHA-256). In at least one embodiment, this function takes binary data of the software component as input to produce a fixed-size output (hash). In at least one embodiment, subsequently, CPU 106 securely transfers this hash value to TPM 108 over bus 124, which stores it in one or more PCRs 110. In at least one embodiment, PCRs 110 are configured to accept new values only through certain secure operations, such as extending an existing hash with a new hash in a cryptographically secure manner.

[0063] In at least one embodiment, extending a current hash in a PCR is taking the current value and the hash of the component being measured, combining them, and then rehashing the combination to produce a new value to store back in the PCR. In at least one embodiment, this process creates a chain of trust, where each new measurement depends on the previous state, ensuring that any unauthorized changes to software can be detected.

[0064] In at least one embodiment, the PCR values are stored in volatile memory and are reset to a known secure state upon a reboot of TPM 108 or computer 102 or CPU 106. In at least one embodiment, upon system startup, CPU 106 remeasures components in computer 102 (e.g., BIOS, bootloader, operating system, installed applications, etc.) and extends these values into the PCRs. In at least one embodiment, the components are remeasured to recalculate hash values of the components within computer 102.

[0065] In at least one embodiment, these PCR values are stored in volatile memory and are reset to their default values (e.g., zero or manufacturer-specific baseline). Thus, while a hacker cannot directly modify these PCR values by extending them using a command such as “TPM2_PCR_Extend,” they can change the values by resetting or rebooting TPM 108.

[0066] In at least one embodiment, TPM 108 also includes a null hierarchy 122, which is volatile memory used to store transient objects. In at least one embodiment, management of null hierarchy 122 is handled internally by TPM 108 using strict protocols such that all operations are performed securely and all transient data in null hierarchy 122 is cleared upon reset or power down.

[0067] In at least one embodiment, startup software 128 initiates extension of PCR values. In at least one embodiment, startup software 128 includes one or more of: startup firmware; operating system software; specific security software that interfaces with TPM 108; Unified Extensible Firmware Interface (UEFI); or Basic Input / Output System (BIOS).

[0068] In at least one embodiment, startup software 128 cooperates with CPU 106 to initiate software instructions to extend PCR values in TPM 108. In at least one embodiment, an example of such software instructions is “TPM2_PCR_Extend,” which is used to update the value of a specified PCR to include a hash of new data. In at least one embodiment, the software instructions are executed or otherwise implemented by a TPM driver or a direct interface layer that communicates between an operating system (or other software) of computer 102 and TPM 108. In at least one embodiment, CPU 106 securely sends these software instructions (e.g., commands) to TPM 108 via the operating system and the TPM driver. In at least one embodiment, TPM 108 subsequently executes these software instructions and cryptographic operations and stores updated PCR values in TPM 108.

[0069] In at least one embodiment, CPU 106 is connected with TPM 108 over bus 124. In at least one embodiment, bus 124 can utilize one of a number of bus interfaces, such as low pin count (LPC), serial peripheral interface (SPI), and inter-integrated circuit (I2C). In embodiments, bus 124 is protected using cryptographic and security protocols to protect data transmitted over the bus. In at least one embodiment, an example of the security protocol is secure protocol and data model (SPDM).

[0070] In at least one embodiment, when computer 102 is restarted, boot software 128 issues a command via a secure SPDM session to create canary object 116 in TMP 108 and return a value of canary object 116 to CPU 106. In at least one embodiment, the value of canary object 116 is a hash value of the canary object 116. In at least one embodiment, boot software 128 creates canary object 116 and stores canary object 116 into secure enclave 120 on CPU 106. In at least one embodiment, canary object 116 is a random sequence of bytes, a string, a numeric file, a random number, a counter, or just a piece of information. In at least one embodiment, a unique identifier is attached to canary object 116 to indicate that the canary object was created by CPU 106.

[0071] In at least one embodiment, secure enclave 120 is a dedicated area within CPU 106 for storing sensitive data and is isolated from other parts of CPU 106, enabling it to resist attacks, including attacks from privileged software. In at least one embodiment, data stored within secure enclave 120 (e.g., canary object 116) is encrypted, and the decryption key is only available within the enclave itself. In at least one embodiment, the data cannot be decrypted by any external process, even the operating system, reducing the risk of sensitive information being compromised.

[0072] In at least one embodiment, CPU 106 creates a copy 112 of canary object 116 and stores copy 112 in empty hierarchy 122. In at least one embodiment, live attacker 126 uses a hacker device that impersonates CPU 106 to create a separate secure session (e.g., SPDM session) to issue a command to reset TPM 108, which causes existing PCR values in TPM 108 to be reset to default values. In at least one embodiment, live attacker 126 can issue a command (e.g., “TPM2_Clear”) to reset TPM 108. In at least one embodiment, after these PCR values are reset, live attacker 126 replays its own PCR values by simulating a boot process to match a state (e.g., a system state including malware) that live attacker 126 has prepared. In at least one embodiment, when live attacker 126 resets TPM 108, system 100 provides identifiable information. In at least one embodiment, the identifiable information is the absence of canary object 112 or incorrect values for canary object 112 in empty hierarchy 122.

[0073] In at least one embodiment, attester 130 is a security program for attesting to an attestation report from computer 102. In at least one embodiment, attester 130 can also be a dedicated security server, a trusted processor, a network appliance, or a hybrid server. In at least one embodiment, CPU 106 can act as attester 130. In at least one embodiment, CPU 106 can include one or more software programs or one or more circuits or one or more firmware that function as an attester (e.g., attester 130).

[0074] In at least one embodiment, an attestation report is a document generated by computer 102 to provide attested information about a state of computer 102. In at least one embodiment, the report includes measurements (e.g., PCR values) and operations of computer 102. In at least one embodiment, the report is signed using a key and includes a nonce or a timestamp indicating when the attestation was performed. In at least one embodiment, attester 130 first authenticates the report by attesting to its digital signature, and then compares PCR values in the report to predefined baseline values. In at least one embodiment, if a discrepancy is found, attester 130 seeks additional information to determine whether the change was legitimate; if validated, the baseline values are updated accordingly. In at least one embodiment, if the change cannot be proven to be legitimate, computer 102 can be deemed compromised.

[0075] In at least one embodiment, using attestation reports to verify the state of computer 102 involves complex setup and ongoing updates to maintain accuracy, which requires significant time and resources. Moreover, in at least one embodiment, a sophisticated attacker can manipulate or mimic PCR values to avoid detection. In at least one embodiment, for example, an attacker (e.g., live attacker 126) can use a legitimate software measurement to inject a PCR extension command while malware is executing.

[0076] In at least one embodiment, system 100 provides a solution to detect attackers that does not rely on attestation reports and PCR values, thereby avoiding the above-mentioned drawbacks of using attestation reports.

[0077] Figure 2 Process 200 illustrates creating a canary object in a TPM, according to at least one embodiment. In at least one embodiment, process 200 is performed to prevent a live attacker (e.g., live attacker 126) from creating a separate canary object after rebooting TPM 108, which would erase the canary object created by CPU. In at least one embodiment, process 200 is performed by a CPU, such as in conjunction with Figure 1 CPU 106 described.

[0078] In at least one embodiment, some or all of process 200 (or any other process described herein, or variations and / or combinations thereof) is performed under the control of one or more computer systems configured with computer-executable instructions to perform the process, and is embodied in computer- executable instructions (e.g., in one or more computer programs or in one or more applications) that are stored in computer-readable storage media, which are In at least one embodiment, the computer-readable storage media is a non-transitory computer-readable medium. In at least one embodiment, computer-readable instructions used to perform at least some of process 200 are not stored solely using transitory signals (e.g., a propagating transient electric or electromagnetic transmission). In at least one embodiment, a non-transitory computer-readable medium does not include a transitory signal per se. In at least one embodiment, process 200 is performed, at least in part, on a computer system, such as a computer system described elsewhere in this disclosure. In at least one embodiment, logic (e.g., hardware, software, or a combination of both) performs process 200.

[0079] In at least one embodiment, at step 202, a CPU (e.g., CPU 106) creates a canary object in a TPM (e.g., TPM 108) that is associated with the CPU. In at least one embodiment, the canary object is created in a TPM register (e.g., a TPM register associated with the CPU). In at least one embodiment, the canary object is created in a TPM register that is not used by the operating system (e.g., operating system 110) or other software running on the computer (e.g., computer 102). Figure 1The CPU 106 described in FIG. 1) issues a command via a secure session (e.g., an SPDM session) between the CPU and the TPM to create a canary object in the TPM. In at least one embodiment, the canary object corresponds to Figure 1 The canary object 116 described in FIG. 1. In at least one embodiment, the TPM corresponds to Figure 1 The TPM 108 described in FIG. 1.

[0080] In at least one embodiment, at step 204, the TPM stores the canary object in the TPM. In at least one embodiment, the canary object is a random sequence of bytes, a string, a numerical file, a random number, a counter, or just a piece of information. In at least one embodiment, the canary object is stored in a PCR as a measurement of the state of the CPU.

[0081] In at least one embodiment, at step 206, the TPM returns the value of the object to the CPU via a secure session (e.g., an SPDM session). In at least one embodiment, the CPU stores the returned value in a secure enclave (e.g., the secure enclave 120), where the returned value is used by an attester (e.g., the attester 130) to attest whether the canary object in the TPM has been replaced by a new canary object created by an in-the-wild attacker. In at least one embodiment, the value of the canary object is a hash value, a unique number, or another piece of information that can uniquely identify the canary object.

[0082] Figure 3 A secure cryptographic binding between a CPU 306 and a TPM 308 in a computer 302 is shown, in accordance with at least one embodiment. In at least one embodiment, this secure cryptographic binding is established when the computer 302 is manufactured in a secure, trusted production environment. In at least one embodiment, the computer 302 corresponds to Figure 1 the computer 102 in FIG. 1, the CPU 306 corresponds to Figure 1 the CPU 106 in FIG. 1, and the TPM 308 corresponds to Figure 1 the TPM 108 in FIG. 1.

[0083] In at least one embodiment, the CPU 306 is equipped with a public key 333 that the TPM 308 uses as a responder to participate in SPDM session establishment.

[0084] In at least one embodiment, the TPM 308 is equipped with a private key 320 that corresponds to the public key 333 on the CPU 306. In at least one embodiment, the public key 333 is stored in a secure storage 316 (corresponding to Figure 1the secure enclave 120 in FIG. 1). In at least one embodiment, the private key 320 is stored in a secure storage 318, which is a tamper-resistant non-volatile memory. In at least one embodiment, as part of an SPDM session establishment during a computer 302 boot-up, the CPU 306 uses the public key 333 and the private key 320 to authenticate the TPM 308 and confirm that the TPM 308 is a secure device that the CPU 306 intends to extend PCR values into.

[0085] Figure 4 is a sequence diagram 400 illustrating operations performed during a secure boot-up of a computer, according to at least one embodiment. In at least one embodiment, this secure boot-up can prevent unauthorized modifications to firmware and configuration of the computer. In at least one embodiment, a CPU 402 corresponds to the CPU 106 in FIG. 1, while a TPM 404 corresponds to the TPM 108 in FIG. 1. Figure 1 Figure 1

[0086] In at least one embodiment, at step 406, a firmware stack (e.g., the boot software 128 in FIG. 1) of the CPU 402 initiates a secure communication session with the TPM 404 using an SPDM protocol. In at least one embodiment, all subsequent commands are issued within this SPDM secure session. Figure 1

[0087] In at least one embodiment, at step 408, the firmware randomization platform authorization value (e.g., platformAuth) of the CPU 402 is authorized. In at least one embodiment, this randomization platform authorization value can prevent unauthorized access to a platform hierarchy of the TPM 404.

[0088] In at least one embodiment, at step 410, the firmware of the CPU 402 configures an authorization policy for a PCR within the TPM 404 using the randomization platform authorization value. In at least one embodiment, this authorization policy only allows the CPU 402 to extend a PCR value in the TPM 404 only within an SPDM session between the CPU 402 and the TPM 404. In at least one embodiment, this authorization policy is set for the TPM 404 in such a way that an attacker (e.g., the field attacker 126) cannot change any authorization policy for a PCR in the TPM 404, and a command (e.g., TPM2_PCR_Extension) for a given PCR set is only authorized when it is executed within an SPDM session between the TPM 404 and the CPU 402.

[0089] In at least one embodiment, at step 412, the firmware of the CPU 402 uses the randomization platform authorization value to extend a PCR value in the TPM 404. In at least one embodiment, this extension of the PCR value is only allowed within the SPDM session between the CPU 402 and the TPM 404. Figure 2 ​​​The process described in the middle creates a canary object within the empty hierarchy of the TPM 404.

[0090] In at least one embodiment, at step 414, the TPM 404 generates a value for this canary object. In at least one embodiment, this value is only accessible through a specific SPDM session used to extend a PCR value. In at least one embodiment, at step 416, the TPM 404 then extends this generated value into a PCR, such as a manufacturer specific PCR. In at least one embodiment, before being extended, the generated canary value is returned to the CPU 402. In at least one embodiment, the CPU 402 then stores the generated canary value in the CPU 402.

[0091] In at least one embodiment, at step 418, after the configuration that requires an authorization policy is completed, the firmware of the CPU 402 discards the randomized platform authorization value of the TPM 404.

[0092] In at least one embodiment, at step 420, as part of the secure boot process, the firmware of the CPU 402 continues to extend measurements into the PCR of the TPM 404.

[0093] Figure 5A and Figure 5B are block diagrams illustrating interactions between a CPU and a TPM acting as a SPDM Requester and a SPDM Responder, respectively, in accordance with at least one embodiment. In at least one embodiment, this TPM exposes special non-volatile (NV) indices that are only accessible by the TPM caller when exchanging commands through a SPDM session. In at least one embodiment, when issuing a TPM command within a SPDM session, this TPM populates these NV indices with the Requester (CPU) and Responder (TPM) public keys used during the SPDM session setup. In at least one embodiment, an NV index is a protected area within the TPM used to securely store data in a non-volatile manner.

[0094] In at least one embodiment, TPM2_PolicyNV command allows a policy to be created that mandates the contents of a particular NV index to be a particular value. In at least one embodiment, two PolicyNV items 530 and 532 included in policy NV 526 are chained together to ensure that an SPDM session is established between a particular TPM (e.g., 508) and a particular CPU (e.g., 502). In at least one embodiment, a requestor (CPU) 502 requests access to data or functionality in a responder (TPM) 504. In at least one embodiment, the requestor 502 includes an SPDM requestor key 506 that is used to establish a secure session with the responder 504. In at least one embodiment, the responder 504 responds to the request of the requestor and includes an SPDM responder key 508 that corresponds to the requestor key 506 for establishing secure communication between the requestor 502 and the responder 504.

[0095] In at least one embodiment, an initial handshake operation 510 is performed in which the requestor 502 and the responder 504 establish a secure communication channel using the SPDM protocol. In at least one embodiment, the requestor 502 then initiates an authenticated session 512 with the responder 504 to perform an operation that requires authorization, such as accessing sealed data or using an NV index. In at least one embodiment, the authenticated session 512 started within a particular SPDM session can only be used within that SPDM session.

[0096] In at least one embodiment, the requestor 502 implements an authorization policy 514 using a TPM2_PolicyNV command that specifies that the contents of a particular NV index 530 are a particular value. In at least one embodiment, the command is issued within an authenticated session, thereby binding the authorization policy to the particular session context. In at least one embodiment, the responder 504 checks whether the SPDM requestor key 506 stored in the NV index 530 matches an expected SPDM requestor key encoded in the authorization policy 524 of object 522.

[0097] In at least one embodiment, the requestor 502 performs an unseal operation 516 within the authenticated session that is only allowed if the policy requirements set by the previous policy are satisfied. In at least one embodiment, unsealing makes the sealed data accessible depending on whether the policy conditions of the session are satisfied. In at least one embodiment, the unseal operation binds a value for extending a PCR value in the responder 504 to the particular SPDM session.

[0098] In at least one embodiment, once the authorization policy is verified and data is unsealed, sensitive data is retrieved from responder 504 and sent back 518 to requester 502. In at least one embodiment, the sensitive data is as described above with respect to FIG. 5B. Figure 4 the generated canary object.

[0099] In at least one embodiment, object 522 (e.g., a canary object 116 as described above with respect to FIG. 5B) is bound to SPDM requester key 520, so any operation on object 522 requires verification according to the policy and key defined in NV index 530. In at least one embodiment, policy enforcement using NV indexes 530 and 532 adds an additional layer of security by ensuring that both requester 502 and responder 504 agree and adhere to certain policies before any sensitive operation can be performed. Figure 1

[0100] In at least one embodiment, object 522 is a piece of data, a certificate, or some other digitally representable item, and is bound to SPDM requester key 520. In at least one embodiment, the binding ensures that object 522 is associated with a particular key belonging to requester 502. In at least one embodiment, object 522 is subject to a security policy and is associated with attestation key 528.

[0101] Figure 6 ​is a flowchart illustrating a process 600 of attesting to a TPM and a CPU, according to an embodiment. In at least one embodiment, the process 600 attests to whether an intended TPM is used to store PCR values and whether an intended CPU is used to extend PCR values. In at least one embodiment, part or all of the process 600 (or any other process described herein, or variations and / or combinations thereof) is performed under the control of one or more computer systems configured with computer-executable instructions and is implemented as code (e.g., computer-executable instructions, one or more computer programs, or one or more applications) executing collectively on one or more processors, hardware, software, or combinations thereof. In at least one embodiment, the code is stored on a computer-readable storage medium in a computer program comprising a plurality of computer-readable instructions executable by the one or more processors. In at least one embodiment, a computer-readable storage medium is a non-transitory computer-readable medium. In at least one embodiment, at least some of the computer-readable instructions to perform the process 600 are not stored solely using transitory signals (e.g., a propagating transient electric or electromagnetic transmission). In at least one embodiment, a non-transitory computer-readable medium does not encompass transitory signals per se. In at least one embodiment, a non-transitory computer-readable medium does not necessarily encompass transitory signals’ in-transit circuitry (e.g., a buffer, cache, and queue) within a transceiver. In at least one embodiment, the process 600 is performed at least in part on a computer system (such as a computer system described elsewhere in this disclosure). In at least one embodiment, logic (e.g., hardware, software, or combinations of hardware and software) performs the process 600. In at least one embodiment, the process 600 is performed by the attester 130 (e.g., Figure 1 described elsewhere in this disclosure).

[0102] In at least one embodiment, at step 602, the attester confirms the SPDM identity of the CPU by attesting that the public key of the CPU is a public key that has been certified by the platform manufacturer. In at least one embodiment, the CPU corresponds to the CPU 106 in Figure 1 .

[0103] In at least one embodiment, at step 604, the attester verifies whether the expected firmware of the CPU has extended the PCR of the TPM. In at least one embodiment, the TPM corresponds to the TPM 104 in Figure 1The TPM 108 in the example. In at least one embodiment, the verifier first verifies whether the CPU holds a private key for establishing an SPDM session identified in the authorization policy of the PCR via an NV index containing the public key. In at least one embodiment, the verifier challenges the canary object (e.g., by verifying the canary object using a command such as “TPM2_Certify” and a long-lived proof key (e.g., IAK) stored in the TPM). Figure 1 The canary object 112 in the PCR is then compared with the values ​​in the PCR and associated event logs. In at least one embodiment, this step is to ensure that the public key of the CPU listed in the authorization policy of the canary object matches the public key listed in the authorization policy of the PCR.

[0104] In at least one embodiment, in step 606, the verifier checks whether the PCR authorization policy is correct. In at least one embodiment, the verifier verifies the SPDM public key of the TPM, for example, through a certificate issued by the platform manufacturer binding the TPM to the platform. In at least one embodiment, the security procedure checks whether the same SPDM session is identified in the authorization policy of the PCR and the authorization policy of the canary object.

[0105] In at least one embodiment, in step 608, the verifier continues the normal PCR verification process. In at least one embodiment, for example, after successfully verifying the aforementioned specific point, the verifier checks the platform (e.g., a computer, such as...). Figure 1 The integrity of the software and hardware configuration of computer 102.

[0106] Figure 7 This is a flowchart illustrating a process 700 for creating a provable PCR amplification mechanism according to at least one embodiment. In at least one embodiment, this provable PCR amplification mechanism is used to create a computing system or platform having the following properties (e.g., Figure 1 Computer 102 in the computing system: uses a pre-defined CPU; the pre-defined CPU expands the measurement results of the computing system, the measurement results being untampered with (e.g., deleted, added, and modified) during transmission between the CPU and the TPM of the computing system; and a verifier can verify these attributes. In at least one embodiment, process 700 is performed by... Figures 1-8 The process 700 is performed by one or more components described herein. In at least one embodiment, process 700 is performed by... Figure 1 One or more of the CPU 106, TPM 108, or verifier 130 described herein shall be executed.

[0107] In at least one embodiment, at step 702, the CPU and TPM bootstrap their cryptographic bindings during computer / platform manufacturing. In at least one embodiment, at this stage, both the CPU and the TPM are set up in a way that establishes a cryptographic binding between them to create a secure communication channel and ensure that both the CPU and the TPM can trust each other’s cryptographic identity. In at least one embodiment, the CPU is provisioned with a public key (or a certificate thereof) that the TPM can use. In at least one embodiment, this public key is used in a SPDM session in which the TPM acts as a responder. In at least one embodiment, the CPU uses this public key to attest the identity of the TPM and establish a secure communication channel with it. In at least one embodiment, on the other hand, the TPM is provisioned with an Initial Attestation Key (IAK). In at least one embodiment, the IAK is a private key. In at least one embodiment, the TPM also receives a certificate from the platform / computer manufacturer that verifies the IAK. In at least one embodiment, this certificate includes a public key corresponding to the IAK, and it links the IAK to the specific TPM and platform (e.g., computer, such as computer 102 in FIG. 1). Figure 1

[0108] In at least one embodiment, both the CPU and the TPM have their own public keys. In at least one embodiment, the public key on the CPU is a pre-stored public key in secure storage that is used to attest the public key of the TPM used in a SPDM session so that the CPU and TPM communicate as originally intended.

[0109] In at least one embodiment, at step 704, the CPU checks whether the TPM is the expected TPM and has not been tampered with during platform boot-up. In at least one embodiment, the CPU extends its measurements into the TPM, creating a cryptographic binding to ensure that these measurements originated from the CPU and have not been intercepted or altered.

[0110] In at least one embodiment, at step 706, a secure program checks whether the TPM is the expected TPM and has not been tampered with.

[0111] Figure 8 ​is a flowchart illustrating a process 800 of detecting unauthorized modifications by live attackers, in accordance with at least one embodiment. In at least one embodiment, part or all of process 800 (or any other process described herein, or variations and / or combinations thereof) is performed under control of one or more computer systems configured with computer-executable instructions, and is implemented as code (for example, computer-executable instructions, one or more computer programs, or one or more applications) executed by a processor or processors of a hardware computer system. In at least one embodiment, code is stored on a computer-readable storage medium of a computer-readable storage medium. In at least one embodiment, a computer-readable storage medium is a non-transitory computer-readable medium. In at least one embodiment, at least some computer-readable instructions used to perform process 800 are not stored only using transitory signals (for example, a propagating transitory electromagnetic transmission). In at least one embodiment, a non-transitory computer-readable medium does not necessarily encompass transitory signals per se. In at least one embodiment, process 800 is performed at least in part on a computer system (such as a computer system described elsewhere in this disclosure). In at least one embodiment, logic (for example, hardware, software, or a combination of hardware and software) performs process 800. In at least one embodiment, process 800 is performed by a verifier 130 (for example, a security program) as described elsewhere in this disclosure. Figure 1

[0112] In at least one embodiment, at step 802, the verifier accesses a secure storage on a CPU (for example, CPU 106 in Figure 1 ) to retrieve a canary object. In at least one embodiment, the verifier is authorized by a particular policy and access control to access a protected region within the CPU storing the canary object.

[0113] In at least one embodiment, at step 804, the verifier accesses a volatile memory location (for example, a location in an empty hierarchy 122 in Figure 1 ) on a TPM (for example, TPM 108 in Figure 1 ). In at least one embodiment, this volatile memory location is in an empty hierarchy, and thus is accessible by any command that does not explicitly require hierarchy authorization.

[0114] In at least one embodiment, at step 806, the verifier checks whether a canary object matching the expected canary from the CPU is present in the volatile memory location.

[0115] ​In at least one embodiment, at step 808, if the canary object in the volatile memory location does not match the canary object from the CPU, the verifier signals that the TPM can have been reset, indicating potential tampering. In at least one embodiment, for example, a live attacker can initiate a separate SPDM session and use that session to restart the TPM to erase the original canary object stored there by the CPU, leaving no object in the volatile memory. In at least one embodiment, if the attacker then creates a new object in the volatile memory after the restart, that new object will not carry the unique identifier that the canary object from the CPU has, and thus will not be a matching canary object.

[0116] In at least one embodiment, at step 810, if a matching canary object exists in the volatile memory location, the verifier signals that the TPM has not been tampered with, and process 800 returns to step 804. In at least one embodiment, the verifier waits a fixed interval (e.g., 5 seconds) before performing step 804 again.

[0117] In the following description, numerous specific details are set forth to provide a more thorough understanding of at least one embodiment. However, it will be apparent to one skilled in the art that the inventive concept can be practiced without one or more of these specific details.

[0118] Data Center

[0119] Figure 9 An example data center 900 is shown in accordance with at least one embodiment. In at least one embodiment, data center 900 includes, without limitation, a data center infrastructure layer 910, a framework layer 920, a software layer 930, and an application layer 940.

[0120] In at least one embodiment, with respect to Figure 9 At least one component shown or described is used to implement a technique and / or functionality described in connection with at least one embodiment. Figures 1-8 In at least one embodiment, the at least one embodiment includes one or more secure devices having information that will be used to detect whether the secure device has been restarted, indicating that the one or more secure devices have been tampered with.

[0121] In at least one embodiment, as Figure 9As shown, the data center infrastructure layer 910 can include a resource orchestrator 912, grouped computing resources 914, and node computing resources (“node C.R.s”) 916(1)-916(N), where “N” represents any whole, positive integer. In at least one embodiment, node C.R.s 916(1)-916(N) can include, but are not limited to, any number of central processing units (“CPUs” or “processors”), including accelerators, field programmable gate arrays (“FPGAs”), data processing units in network devices (“DPUs”), graphics processors, etc., memory devices such as dynamic read-only memory, storage devices such as solid state or disk drives, network input / output (“NW I / O”) devices, network switches, virtual machines (“VMs”), power modules, and cooling modules, etc. In at least one embodiment, one or more node C.R.s of node C.R.s 916(1)-916(N) can be a server having one or more of the above-described computing resources.

[0122] In at least one embodiment, grouped computing resources 914 can include individual groups of node C.R.s housed within one or more racks (not shown), or housed within a number of racks (also not shown) within various geographic locations. Individual groups of node C.R.s within grouped computing resources 914 can include grouped computing, network, memory, or storage resources that can be configured or allocated to support one or more workloads. In at least one embodiment, several node C.R.s including CPUs or processors can be grouped within one or more racks to provide computing resources to support one or more workloads. In at least one embodiment, one or more racks can also include any number of power modules, cooling modules, and network switches, in any combination.

[0123] In at least one embodiment, resource orchestrator 912 can configure or otherwise control one or more node C.R.s 916(1)-916(N) and / or grouped computing resources 914. In at least one embodiment, resource orchestrator 912 can include a software design infrastructure (“SDI”) management entity for data center 900. In at least one embodiment, resource orchestrator 912 can comprise hardware, software, or some combination thereof.

[0124] In at least one embodiment, as Figure 9As shown, framework layer 920 includes, without limitation, a job scheduler 932, a configuration manager 934, a resource manager 936, and a distributed file system 938. In at least one embodiment, framework layer 920 can include a framework that supports software 952 of software layer 930 and / or one or more applications 942 of application layer 940. In at least one embodiment, software 952 or applications 942 can include, respectively, web-based service software or applications, such as services or applications provided by Amazon Web Services, Google Cloud, and Microsoft Azure. In at least one embodiment, framework layer 920 can be, without limitation, a type of free and open-source software web application framework such as Apache Spark™ (hereinafter “Spark”) that can utilize distributed file system 938 for large-scale data processing (e.g., “big data”). In at least one embodiment, job scheduler 932 can include a Spark driver to facilitate scheduling workloads supported by various layers of data center 900. In at least one embodiment, configuration manager 934 can be capable of configuring different layers, such as software layer 930 and framework layer 920 including Spark and distributed file system 938 for supporting large-scale data processing. In at least one embodiment, resource manager 936 can be capable of managing clustered or grouped computing resources mapped to or allocated for supporting distributed file system 938 and job scheduler 932. In at least one embodiment, clustered or grouped computing resources can include grouped computing resources 914 on data center infrastructure layer 910. In at least one embodiment, resource manager 936 can coordinate with resource orchestrator 912 to manage these mapped or allocated computing resources.

[0125] In at least one embodiment, software 952 included in software layer 930 can include software used by at least portions of node C.R.s 916(1)-916(N), grouped computing resources 914, and / or distributed file system 938 of framework layer 920. One or more types of software can include, without limitation, Internet web page search software, email virus scanning software, database software, and streaming video content software.

[0126] In at least one embodiment, one or more applications 942 included in application layer 940 can include one or more types of applications used by at least portions of node C.R.s 916(1)-916(N), grouped computing resources 914, and / or distributed file system 938 of framework layer 920. One or more types of applications can include, without limitation, CUDA applications.

[0127] In at least one embodiment, any of the configuration manager 934, resource manager 936, and resource coordinator 912 can implement any number and type of self-modification actions based on any amount and type of data acquired in any technically feasible manner. In at least one embodiment, self-modification actions can mitigate potentially poor configuration decisions by data center operators of data center 900 and can prevent underutilization and / or poor performance of the data center.

[0128] Computer-based systems

[0129] The following figures present, but are not limited to, exemplary computer-based systems that can be used to implement at least one embodiment.

[0130] Figure 10 A processing system 1000 according to at least one embodiment is illustrated. In at least one embodiment, the system 1000 includes one or more processors 1002 and one or more graphics processors 1008, and may be a single-processor desktop system, a multi-processor workstation system, or a server system having a large number of processors 1002 or processor cores 1007. In at least one embodiment, the processing system 1000 is a processing platform incorporated within a system-on-a-chip (SoC) integrated circuit for use in mobile, handheld, or embedded devices. In at least one embodiment, the processor core 1007 is referred to as a computing unit or arithmetic unit.

[0131] In at least one embodiment, regarding Figure 10 At least one component shown or described is used to implement the combination. Figures 1-8 The described technologies and / or functions. In at least one embodiment, the at least one embodiment includes one or more security devices having information for detecting whether the security device has been restarted, thereby indicating that the one or more security devices have been tampered with.

[0132] In at least one embodiment, the processing system 1000 may include or be integrated into a server-based gaming platform, including a game console, mobile game console, handheld game console, or online game console, which are game and media consoles. In at least one embodiment, the processing system 1000 is a mobile phone, smartphone, tablet computing device, or mobile internet device. In at least one embodiment, the processing system 1000 may also include components coupled to or integrated into a wearable device, such as a smartwatch, smart glasses, augmented reality, or virtual reality device. In at least one embodiment, the processing system 1000 is a television or set-top box device having one or more processors 1002 and a graphical interface generated by one or more graphics processors 1008.

[0133] In at least one embodiment, one or more processors 1002 each include one or more processor cores 1007 to process instructions which, when executed, perform operations such as operations for systems and user software. In at least one embodiment, each of the one or more processor cores 1007 is configured to process a specific instruction set 1009. In at least one embodiment, instruction set 1009 can facilitate complex instruction set computing (CISC), reduced instruction set computing (RISC), or computing via a very long instruction word (VLIW). In at least one embodiment, multiple processor cores 1007 can each process a different instruction set 1009, which can include instructions to facilitate emulation of other instruction sets. In at least one embodiment, processor core 1007 can also include other processing devices, such a digital signal processor (DSP).

[0134] In at least one embodiment, processor 1002 includes cache memory 1004. In at least one embodiment, processor 1002 can have single-level or multi-level internal caches. In at least one embodiment, cache memory is shared among multiple components of processor 1002. In at least one embodiment, processor 1002 also uses an external cache (e.g., a level three (L3) cache or last level cache (LLC)) (not shown), which can be shared among processor cores 1007 using known cache coherency techniques. In at least one embodiment, register file 1006 is additionally included in processor 1002, which can include different types of registers such as integer registers, floating point registers, status registers, and instruction pointer registers. In at least one embodiment, register file 1006 can include a general register file or another register.

[0135] In at least one embodiment, one or more processors 1002 are coupled with one or more interface buses 1010 for passing control information and receiving data signals between processor 1002 and other components of system 1000. In at least one embodiment, interface bus 1010 can be versions of a processor bus, such as Direct Media Interface (DMI) bus. In at least one embodiment, interface bus 1010 is not limited to DMI bus, and can include one or more Peripheral Component Interconnect buses (e.g., PCI, PCI Express), memory buses, or other types of interface buses. In at least one embodiment, processor 1002 includes integrated memory controller 1016 and platform controller hub 1030. In at least one embodiment, memory controller 1016 facilitates communication between memory 1020 and other components of processing system 1000, while platform controller hub 1030 provides connections between local I / O devices and other components within the processing system. In at least one embodiment, one or more peripheral component interconnect buses include PCIe Gen 5, which provides an interface for a processor.

[0136] In at least one embodiment, storage devices 1020 can be dynamic random access memory (DRAM) devices, static random access memory (SRAM) devices, flash memory devices, phase-change memory devices, or some other memory devices having suitable performance to serve as external memory for processor 1002. In at least one embodiment, storage devices 1020 can be used as a system memory for processing system 1000, to store data 1022 and instructions 1021 for use when one or more processors 1002 executes an application or process. In at least one embodiment, memory controller 1016 also couples with an optional external graphics processor 1012, which can communicate with one or more graphics processors 1008 within processor 1002 to perform graphics and media operations. In at least one embodiment, display device 1011 can connect to processor 1002. In at least one embodiment, display device 1011 can include one or more internal display devices, such as a display device within a mobile electronic device or a laptop computer or an external display device connected to processor 1002 through a display interface, such as DisplayPort or the like. In at least one embodiment, display device 1011 can include a head-mounted display (HMD) such as a stereoscopic display device for use in virtual reality (VR) applications or augmented reality (AR) applications.

[0137] In at least one embodiment, platform controller hub 1030 enables peripherals to connect to memory device 1020 and processor 1002 via a high-speed I / O bus. In at least one embodiment, I / O peripherals include, but are not limited to, audio controller 1046, network controller 1034, firmware interface 1028, wireless transceiver 1026, touch sensors 1025, data storage device 1024 (e.g., hard disk drive, flash memory, etc.). In at least one embodiment, data storage device 1024 can connect to the system processor 1002 through a storage interface (e.g., SATA) or through a peripheral bus, such as a Peripheral Component Interconnect bus (e.g., PCI, PCIe). In at least one embodiment, touch sensors 1025 can include touch screen sensors, pressure sensors, or fingerprint sensors. In at least one embodiment, wireless transceiver 1026 can be a Wi-Fi transceiver, a Bluetooth transceiver, or a mobile network transceiver such as a 3G, 4G, or Long Term Evolution (LTE) transceiver. In at least one embodiment, firmware interface 1028 enables communication with system firmware, and can be, for example, a unified extensible firmware interface (UEFI). In at least one embodiment, network controller 1034 can enable network connectivity to one or more private or public networks. In at least one embodiment, a high-performance network controller (not shown) is coupled to interface bus 1010. In at least one embodiment, audio controller 1046 is a multi-channel high definition audio controller. In at least one embodiment, processing system 1000 includes an optional legacy I / O controller 1040 for coupling legacy (e.g., Personal System 2 (PS / 2)) devices to processing system 1000. In at least one embodiment, platform controller hub 1030 can also connect to one or more Universal Serial Bus (USB) controllers 1042 connect input devices, such as keyboard and mouse 1043 combinations, camera 1044, or other USB input devices.

[0138] In at least one embodiment, memory controller 1016 and instances of platform controller hub 1030 can be integrated into a discrete external graphics processor, such as external graphics processor 1012. In at least one embodiment, platform controller hub 1030 and / or memory controller 1016 can be external to one or more processor(s) 1002. For example, in at least one embodiment, processing system 1000 can include an external memory controller 1016 and platform controller hub 1030, which can be configured as a memory controller hub and a peripheral controller hub in a system chipset that is in communication with processor(s) 1002.

[0139] Figure 11A computer system 1100 according to at least one embodiment is shown. In at least one embodiment, computer system 1100 can be a system-on-a-chip, SOC, or some combination of a system and a processor. In at least one embodiment, computer system 1100 is formed from a processor 1102 that can include execution units to execute an instruction, in at least one embodiment, computer system 1100 can include, without limitation, a component, such as a processor 1102, to employ execution units including logic to perform algorithms for process data. In at least one embodiment, computer system 1100 can include processors, such as a Pentium® XeonTM, XScaleTMand / or StrongARMTM, Core TM or Nervana TM microprocessors, though other systems (including PCs, workstations, set-top boxes, etc. with other microprocessors) can also be used. In at least one embodiment, computer system 1100 can execute a version of the WINDOWS operating system available from Microsoft Corporation of Redmond, Wash., although other operating systems (UNIX and Linux, for example), embedded software, and / or graphical user interfaces, can also be used.

[0140] In at least one embodiment, at least one component shown or described is used to implement techniques and / or functionality described in conjunction with Figure 11 at least one embodiment. In at least one embodiment, the at least one embodiment includes one or more secure devices having information that will be used to detect whether the secure device has been restarted indicating that the one or more secure devices have been tampered with. Figures 1-8

[0141] In at least one embodiment, the computer system 1100 can be used in other devices, such as handheld devices and embedded applications. Some examples of handheld devices include cellular phones, Internet Protocol (IP) devices, digital cameras, personal digital assistants (“PDAs”), and handheld PCs. In at least one embodiment, the embedded application may include a microcontroller, a digital signal processor (“DSP”), a system-on-a-chip (SoC), a network computer (“NetPC”), a set-top box, a network hub, a wide area network (“WAN”) switch, or any other system capable of executing one or more instructions according to at least one embodiment.

[0142] In at least one embodiment, the computer system 1100 may include, but is not limited to, a processor 1102, which may include, but is not limited to, one or more execution units 1108 configured to execute a Computational Unified Device Architecture (“CUDA”). The program is developed by NVIDIA Corporation in Santa Clara, California. In at least one embodiment, the CUDA program is at least a part of a software application written in the CUDA programming language. In at least one embodiment, the computer system 1100 is a single-processor desktop or server system. In at least one embodiment, the computer system 1100 may be a multiprocessor system. In at least one embodiment, the processor 1102 may include, but is not limited to, a CISC microprocessor, a RISC microprocessor, a VLIW microprocessor, a processor implementing instruction set combinations, or any other processor device, such as a digital signal processor. In at least one embodiment, the processor 1102 may be coupled to a processor bus 1110, which can transmit data signals between the processor 1102 and other components in the computer system 1100.

[0143] In at least one embodiment, processor 1102 may include, but is not limited to, a Level 1 (“L1”) internal cache memory (“cache”) 1104. In at least one embodiment, processor 1102 may have a single internal cache or multiple levels of internal cache. In at least one embodiment, the cache memory may reside external to processor 1102. In at least one embodiment, processor 1102 may include a combination of internal and external caches. In at least one embodiment, register file 1106 may store different types of data in various registers, including but not limited to integer registers, floating-point registers, status registers, and instruction pointer registers.

[0144] In at least one embodiment, execution unit 1108 includes, without limitation, logic to perform integer and floating point operations, including the execution of int and float instructions, in at least one embodiment. Processor 1102 can also include microcode ("ucode") read-only memory ("ROM"), in at least one embodiment, that stores microcode for certain macro instructions. In at least one embodiment, execution unit 1108 can also include logic to handle a packed instruction set 1109. In at least one embodiment, by including the packed instruction set 1109 in an instruction set of a general- purpose processor 1102, along with associated circuitry to execute the instructions, the general purpose processor 1102 can be used to perform the operations of many multimedia applications using the packed data of the processor in a manner that is more efficient than using a specialized graphics or DSP processor.

[0145] In at least one embodiment, execution unit 1108 can also be used in microcontrollers, embedded processors, graphics devices, DSPs, and other types of logic circuits. In at least one embodiment, computer system 1100 can include, without limitation, memory 1120. In at least one embodiment, memory 1120 can be implemented as a DRAM device, SRAM device, flash memory device, or other memory device. Memory 1120 can store instructions 1119 and / or data 1121 that can be executed by processor 1102, in at least one embodiment.

[0146] In at least one embodiment, a system logic chip can be coupled to processor bus 1110 and memory 1120. In at least one embodiment, system logic chip can include, without limitation, a memory controller hub (“MCH”) 1116 and processor 1102 can communicate with MCH 1116 via processor bus 1110. In at least one embodiment, MCH 1116 can provide a high bandwidth memory path 1118 to memory 1120 for instruction and data storage and for storage of graphics commands, data, and textures. In at least one embodiment, MCH 1116 can direct data signals between processor 1102, memory 1120, and other components in computer system 1100, and communicate data signals between processor bus 1110, memory 1120, and system I / O 1122. In at least one embodiment, system logic chip can provide a graphics port for coupling to a graphics controller. In at least one embodiment, MCH 1116 can be coupled to memory 1120 through a high bandwidth memory path 1118 and to graphics / video card 1112 through an Accelerated Graphics Port (“AGP”) interconnect 1114.

[0147] In at least one embodiment, computer system 1100 can use system I / O 1122 as a proprietary hub interface bus to couple MCH 1116 to I / O controller hub (“ICH”) 1130. In at least one embodiment, ICH 1130 can provide a direct connection to some I / O devices and

[0148] In at least one embodiment, Figure 11 A system including interconnected hardware devices or “chips” is shown. In at least one embodiment, Figure 11 An exemplary SoC can be shown. In at least one embodiment, Figure 11The devices shown can be interconnected with proprietary interconnects, standardized interconnects (e.g., PCIe), or some combination thereof. In at least one embodiment, one or more components of system 1100 are interconnected using a compute fast link (CXL) interconnect.

[0149] Figure 12 A system 1200 according to at least one embodiment is illustrated. In at least one embodiment, system 1200 is an electronic device utilizing processor 1210. In at least one embodiment, system 1200 may be, for example, but not limited to, a laptop computer, tower server, rack server, blade server, edge device communicatively coupled to one or more local or cloud service providers, laptop computer, desktop computer, tablet computer, mobile device, telephone, embedded computer, or any other suitable electronic device.

[0150] In at least one embodiment, regarding Figure 12 At least one component shown or described is used to implement the combination. Figures 1-8 The described technologies and / or functions. In at least one embodiment, the at least one embodiment includes one or more security devices having information for detecting whether the security device has been restarted, thereby indicating that the one or more security devices have been tampered with.

[0151] In at least one embodiment, system 1200 may include, but is not limited to, processor 1210 communicatively coupled to any suitable number or type of components, peripherals, modules, or devices. In at least one embodiment, processor 1210 is coupled using a bus or interface, such as I... 2 C-bus, System Management Bus (“SMBus”), Low Pin Count (LPC) bus, Serial Peripheral Interface (“SPI”), High Definition Audio (“HDA”) bus, Serial Advanced Technology Accessory (“SATA”) bus, USB (versions 1, 2, and 3) or Universal Asynchronous Receiver / Transmitter (“UART”) bus. In at least one embodiment, Figure 12 A system is illustrated, comprising interconnected hardware devices or "chips". In at least one embodiment, Figure 12 An exemplary SoC can be shown. In at least one embodiment, Figure 12 The device shown can be interconnected with proprietary interconnects, standardized interconnects (e.g., PCIe), or some combination thereof. In at least one embodiment, Figure 12 One or more components are interconnected using Computational Fast Link (CXL) interconnects.

[0152] In at least one embodiment, Figure 12The components can include a display 1224, a touch screen 1225, a touch pad 1230, a near field communication unit ("NFC") 1245, a sensor hub 1240, a thermal sensor 1246, an express chipset ("EC") 1235, a trusted platform module ("TPM") 1238, a BIOS / firmware / flash memory ("BIOS, FW Flash") 1222, a DSP 1260, a solid state disk ("SSD") or hard drive ("HDD") 1220, a wireless local area network unit ("WLAN") 1250, a Bluetooth unit 1252, a wireless wide area network unit ("WWAN") 1256, a global positioning system (GPS) 1255, a camera ("USB 3.0 camera") 1254 (e.g., a USB 3.0 camera), or a low power double data rate ("LPDDR") memory unit ("LPDDR3") 1215 implemented in, for example, the LPDDR3 standard. These components can each be implemented in any suitable manner.

[0153] In at least one embodiment, other components can be communicatively coupled to processor 1210 by components discussed above. In at least one embodiment, an accelerometer 1241, an ambient light sensor ("ALS") 1242, a compass 1243, and a gyroscope 1244 can be communicatively coupled to sensor hub 1240. In at least one embodiment, a thermal sensor 1239, a fan 1237, a keyboard 1236, and a touch pad 1230 can be communicatively coupled to EC 1235. In at least one embodiment, a speaker 1263, a headphone 1264, and a microphone ("mic") 1265 can be communicatively coupled to an audio unit ("audio codec and class D amplifier") 1262, which can in turn be communicatively coupled to DSP 1260. In at least one embodiment, audio unit 1262 can include, for example and without limitation, an audio coder / decoder ("codec") and a class D amplifier. In at least one embodiment, a SIM card ("SIM") 1257 can be communicatively coupled to WWAN unit 1256. In at least one embodiment, components such as WLAN unit 1250 and Bluetooth unit 1252, as well as WWAN unit 1256, can be implemented as a next generation form factor (NGFF).

[0154] Figure 13An exemplary integrated circuit 1300 is shown, in accordance with at least one embodiment. In at least one embodiment, the exemplary integrated circuit 1300 is a SoC, which can be fabricated using one or more IP cores. In at least one embodiment, integrated circuit 1300 includes one or more application processors 1305 (e.g., CPUs, DPUs), at least one graphics processor 1310, and can additionally include an image processor 1315 and / or a video processor 1320, any of which can be a modular IP core. In at least one embodiment, integrated circuit 1300 includes peripheral or bus logic including USB controller 1325, UART controller 1330, SPI / SDIO controller 1335, and I2S / I2C controller 1340. In at least one embodiment, integrated circuit 1300 can include a display device 1345 coupled to one or more of a high-definition multimedia interface (HDMI) controller 1350 and a mobile industry processor interface (MIPI) display interface 1355. In at least one embodiment, storage can be provided by flash storage subsystem 1360, including flash memory and a flash memory controller. In at least one embodiment, memory interfaces can be provided via a memory controller 1365 for access to SDRAM or SRAM memory devices. In at least one embodiment, some integrated circuits also include an embedded security engine 1370. 2 S / I 2 C controller 1340. In at least one embodiment, integrated circuit 1300 can include a display device 1345 coupled to one or more of a high-definition multimedia interface (HDMI) controller 1350 and a mobile industry processor interface (MIPI) display interface 1355. In at least one embodiment, storage can be provided by flash storage subsystem 1360, including flash memory and a flash memory controller. In at least one embodiment, memory interfaces can be provided via a memory controller 1365 for access to SDRAM or SRAM memory devices. In at least one embodiment, some integrated circuits also include an embedded security engine 1370.

[0155] In at least one embodiment, with respect to Figure 13 At least one component shown or described is used to implement techniques and / or functionality described in connection with Figures 1-8 described. In at least one embodiment, the at least one embodiment includes one or more secure devices having information that will be used to detect whether the secure device has been restarted indicating that the one or more secure devices have been tampered with.

[0156] Figure 14A computing system 1400 is shown in accordance with at least one embodiment. In at least one embodiment, the computing system 1400 includes a processing subsystem 1401 having one or more processor(s) 1402 and a system memory 1404 communicating via an interconnection path 1405 that can include a memory hub 1405. In at least one embodiment, the memory hub 1405 can be a separate component coupled with one or more processors 1402 via individual communication links 1407A to 1407N. In at least one embodiment, memory hub 1405 can be integrated into a processor or with one or more processors 1402.

[0157] In at least one embodiment, with respect to Figure 14 At least one component shown or described is used to implement a method Figures 1-8 described. In at least one embodiment, the at least one embodiment includes one or more secure devices having information that will be used to detect whether the secure device has been restarted indicating that the one or more secure devices have been tampered with.

[0158] In at least one embodiment, processing subsystem 1401 includes one or more parallel processor(s) 1412 coupled to memory hub 1405 via a bus or other communication link 1413. In at least one embodiment, communication link 1413 can be implemented with one or more of standard computer-based bus protocols, such as via a PCI Express bus, a

[0159] In at least one embodiment, system storage 1414 can connect to I / O hub 1407 to provide storage mechanisms for computing system 1400. In at least one embodiment, I / O switches 1416 can be used to provide interface mechanisms to enable connections between I / O hub 1407 and other components such as network adapters 1418 and / or wireless network adapters 1419 that can be integrated into a platform, as well as various other devices that can be added via one or more add-in devices 1420. In at least one embodiment, network adapters 1418 can be Ethernet adapters or another wired network adapters. In at least one embodiment, wireless network adapters 1419 can include one or more of Wi-Fi, Bluetooth, NFC, or other network devices that include one or more radios.

[0160] In at least one embodiment, computing system 1400 can include other components not explicitly shown, including USB or other port connections, optical storage drives, video capture devices, etc., that can also be connected to I / O hub 1407. In at least one embodiment, communication paths interconnecting various Figure 14 Communication paths between various components in FIG. 14 can use any suitable protocol, such as a protocol based on PCI (Peripheral Component Interconnect) such as PCI Express, or another suitable protocol such as a protocol based on NVLink, a high-speed interconnect technology developed by NVIDIA, or other interconnect technologies and / or protocols (e.g., proprietary interconnect technologies developed by other companies).

[0161] In at least one embodiment, parallel processor(s) 1412 include circuitry optimized for graphics and video processing, including for example video output circuitry, and are configured for use in a gaming console, a mobile phone, a personal computer, or other application. In at least one embodiment, parallel processor(s) 1412 include circuitry optimized for general use applications, including for example high-precision floating point, integer and Boolean operations on a wide variety of data. In at least one embodiment, compute system 1400 can include multiple parallel processor(s) 1412. In at least one embodiment, parallel processor(s) 1412 each include independent processing cores that are optimized for use with a graphics processing unit (GPU). In at least one embodiment, parallel processor(s) 1412 are integrated on a common substrate, or in at least one embodiment, one or more of parallel processor(s) 1412 are discrete devices interconnected with a interprocessor interconnect. In at least one embodiment, at least one of parallel processor(s) 1412 includes circuitry for one or more of the following: independent parallel processing cores; register file shared by set of processing cores; memory interface; one or more additional instructions to support general use applications; and / or one or more additional instructions to support graphics applications. In at least one embodiment, one or more of parallel processor(s) 1412 include circuitry to load data into memory; circuitry to distribute data to one or more memory interfaces; circuitry to manage one or more memory

[0162] Processing system

[0163] The following figures illustrate an exemplary processing system that can be used to implement at least one embodiment, but is not limited thereto.

[0164] Figure 15An accelerated processing unit (“APU”) 1500 is shown, in accordance with at least one embodiment. In at least one embodiment, APU 1500 is developed by AMD Corporation of Santa Clara, California. In at least one embodiment, APU 1500 can be configured to execute application programs such as CUDA programs. In at least one embodiment, APU 1500 includes, without limitation, core complex 1510, graphics complex 1540, fabric 1560, I / O interface 1570, memory controllers 1580, display controllers 1592, and multimedia engine 1594. In at least one embodiment, APU 1500 can include, without limitation, any number of core complexes 1510, any number of graphics complexes 1540, any number of display controllers 1592, and any number of multimedia engines 1594, in any combination. For purposes of illustration, multiple instances of like objects are denoted with reference numerals including a letter suffix, in which the reference numeral identifies the object and the letter suffix identifies the instance. In at least one embodiment, APU 1500 is configured to execute a CUDA program.

[0165] In at least one embodiment, with respect to Figure 15 At least one component shown or described is used to implement a technique and / or functionality described Figures 1-8 In at least one embodiment, the at least one embodiment includes one or more secure devices having information that will be used to detect whether the secure device has been restarted indicating that the one or more secure devices have been tampered with.

[0166] In at least one embodiment, core complex 1510 is a CPU, graphics complex 1540 is a GPU, and APU 1500 is a processing unit that will not be limited to integrating 1510 and 1540 onto a single chip. In at least one embodiment, some tasks can be assigned to core complex 1510 while other tasks can be assigned to graphics complex 1540. In at least one embodiment, core complex 1510 is configured to execute host executable code associated with APU 1500, such as an operating system. In at least one embodiment, core complex 1510 is a master processor of APU 1500 that controls and coordinates operations of other processors. In at least one embodiment, core complex 1510 issues commands that control operations of graphics complex 1540. In at least one embodiment, core complex 1510 can be configured to execute host executable code derived from CUDA source code and graphics complex 1540 can be configured to execute device executable code derived from CUDA source code.

[0167] In at least one embodiment, core complex 1510 includes, without limitation, cores 1520(1)-1520(4) and L3 cache 1530. In at least one embodiment, core complex 1510 can include, without limitation, any number of cores 1520 and any combination and number of caches. In at least one embodiment, cores 1520 are configured to execute instructions of a particular instruction set architecture (“ISA”). In at least one embodiment, each core 1520 is a CPU core. In at least one embodiment, cores 1520 are referred to as processing units or execution units.

[0168] In at least one embodiment, each core 1520 includes, without limitation, a fetch / decode unit 1522, an integer execution engine 1524, a floating point execution engine 1526, and an L2 cache 1528. In at least one embodiment, fetch / decode unit 1522 fetches instructions, decodes such instructions, generates micro-operations, and dispatches individual micro-instructions to integer execution engine 1524 and floating point execution engine 1526. In at least one embodiment, fetch / decode unit 1522 can simultaneously dispatch one micro-instruction to integer execution engine 1524 and another micro-instruction to floating point execution engine 1526. In at least one embodiment, integer execution engine 1524 executes, without limitation, integer and memory operations. In at least one embodiment, floating point engine 1526 executes, without limitation, floating point and vector operations. In at least one embodiment, fetch-decode unit 1522 dispatches micro-instructions to a single execution engine in place of both integer execution engine 1524 and floating point execution engine 1526.

[0169] In at least one embodiment, each core 1520(i) can access L2 cache 1528(i) included in core 1520(i), where i is an integer representing a particular instance of core 1520. In at least one embodiment, each core 1520 included in core complex 1510(j) is connected to other cores 1520 included in core complex 1510(j) via L3 cache 1530(j) included in core complex 1510(j), where j is an integer representing a particular instance of core complex 1510. In at least one embodiment, cores 1520 included in core complex 1510(j) can access all L3 caches 1530(j) included in core complex 1510(j), where j is an integer representing a particular instance of core complex 1510. In at least one embodiment, L3 cache 1530 can include, without limitation, any number of slices.

[0170] In at least one embodiment, graphics processing engine 1518 can include, without limitation, a ring interconnect 1515, a register file 1517, one or more SIMD units 1520, a cache 1522, and one or more parallel floating point ALUs 1524. In at least one embodiment, graphics processing engine 1518 can include, without limitation, any number of compute units 1520 and L2 cache 1522. In at least one embodiment, compute units 1520 share L2 cache 1522. In at least one embodiment, L2 cache 1522 is partitioned, for example, between different compute units 1520. In at least one embodiment, graphics processing engine 1518 includes, without limitation, any number of compute units 1520 and any number (including zero) and type of cache. In at least one embodiment, graphics processing engine 1518 includes, without limitation, any number of specialized graphics hardware.

[0171] In at least one embodiment, graphics processing engine 1540 includes, without limitation, any number of compute units 1550 and L2 cache 1542. In at least one embodiment, compute units 1550 share L2 cache 1542. In at least one embodiment, L2 cache 1542 is partitioned. In at least one embodiment, graphics processing engine 1540 includes, without limitation, any number of compute units 1550 and any number (including zero) and type of cache. In at least one embodiment, graphics processing engine 1540 includes, without limitation, any number of specialized graphics hardware.

[0172] In at least one embodiment, each compute unit 1550 includes, without limitation, any number of SIMD units 1552 and shared memory 1554. In at least one embodiment, each SIMD unit 1552 implements a SIMD architecture and is configured to execute operations in parallel. In at least one embodiment, each compute unit 1550 can execute any number of thread blocks, but each thread block executes on a single compute unit 1550. In at least one embodiment, a thread block includes, without limitation, any number of execution threads. In at least one embodiment, a workgroup is a thread block. In at least one embodiment, each SIMD unit 1552 executes a different thread warp. In at least one embodiment, a thread warp is a group of threads (e.g., 16 threads), where each thread in a thread warp belongs to a single thread block and is configured to process a different set of data based on a single instruction set. In at least one embodiment, a predication can be used to disable one or more threads in a thread warp. In at least one embodiment, a lane is a thread. In at least one embodiment, a work item is a thread. In at least one embodiment, a wavefront is a thread warp. In at least one embodiment, different wavefronts in a thread block can synchronize together and communicate via shared memory 1554. In at least one embodiment, each compute unit 1550 includes one or more clusters of thread blocks, where a cluster of thread blocks can enable programmatic control over locality at a greater granularity than a single thread block of a single streaming multi-processor (SM). In at least one embodiment, a cluster of thread blocks (also referred to as a “cluster”) enables multiple thread blocks running concurrently across streaming multi-processors to synchronize and cooperatively fetch, exchange, or otherwise use data.

[0173] In at least one embodiment, fabric 1560 is a system interconnect that facilitates data and control transmissions across core complex 1510, graphics complex 1540, I / O interface 1570, memory controllers 1580, display controller 1592, and multimedia engine 1594. In at least one embodiment, APU 1500 can include, without limitation, any number and type of system interconnects in addition to or instead of fabric 1560 that facilitate data and control transmissions across any number and type of directly or indirectly linked components that can be internal or external to APU 1500. In at least one embodiment, I / O interface 1570 represents any number and type of I / O interface (e.g., PCI, PCI-Extended (“PCI-X”), PCIe, Gigabit Ethernet (“GBE”), USB, etc.). In at least one embodiment, various types of peripheral devices are coupled to I / O interface 1570. In at least one embodiment, peripheral devices coupled to I / O interface 1570 can include, without limitation, a keyboard, a mouse, a printer, a scanner, a joystick or other type of game controller, a media recording device, an external storage device, a network interface card, etc.

[0174] In at least one embodiment, display controller 1592 displays images on one or more display devices, such as liquid crystal display (“LCD”) devices. In at least one embodiment, multimedia engine 1594 includes, without limitation, any number and type of multimedia-related circuitry, such as a video decoder, a video encoder, an image signal processor, etc. In at least one embodiment, memory controllers 1580 facilitate data transfers between APU 1500 and unified system memory 1590. In at least one embodiment, core complex 1510 and graphics complex 1540 share unified system memory 1590.

[0175] In at least one embodiment, APU 1500 implements a memory subsystem that includes, without limitation, any number and type of memory controllers 1580 and memory devices (e.g., shared memory 1554) that can be dedicated to one component or shared among multiple components. In at least one embodiment, APU 1500 implements a cache subsystem that includes, without limitation, one or more cache memories (e.g., L2 cache 1628, L3 cache 1530, and L2 cache 1542), each of which can be private to a component or shared among any number of components (e.g., core 1520, core complex 1510, SIMD unit 1552, compute unit 1550, and graphics complex 1540).

[0176] Figure 16A CPU 1600 is shown, in accordance with at least one embodiment. In at least one embodiment, CPU 1600 is developed by AMD Corporation of Santa Clara, CA. In at least one embodiment, CPU 1600 can be configured to execute application programs. In at least one embodiment, CPU 1600 is configured to execute host executable code derived from CUDA source code, and an external GPU can be configured to execute device executable code derived from such CUDA source code. In at least one embodiment, CPU 1600 includes, without limitation, any number of core complexes 1610, fabric 1660, I / O interfaces 1670, and memory controllers 1680.

[0177] In at least one embodiment, with respect to Figure 16 At least one component shown or described is used to implement a method Figures 1-8 described herein. In at least one embodiment, the at least one embodiment includes one or more secure devices having information that will be used to detect whether the secure device has been re-booted, indicating that the one or more secure devices have been tampered with.

[0178] In at least one embodiment, core complex 1610 includes, without limitation, cores 1620(1)-1620(4) and L3 cache 1630. In at least one embodiment, core complex 1610 can include, without limitation, any number of cores 1620 and any combination and type of caches. In at least one embodiment, cores 1620 are configured to execute instructions of a particular ISA. In at least one embodiment, each core 1620 is a CPU core.

[0179] In at least one embodiment, each core 1620 includes, without limitation, a fetch / decode unit 1622, an integer execution engine 1624, a floating point execution engine 1626, and an L2 cache 1628. In at least one embodiment, fetch / decode unit 1622 fetches instructions, decodes such instructions, generates micro-operations, and dispatches individual micro-instructions to integer execution engine 1624 and floating point execution engine 1626. In at least one embodiment, fetch / decode unit 1622 can simultaneously dispatch one micro-instruction to integer execution engine 1624 and another micro-instruction to floating point execution engine 1626. In at least one embodiment, integer execution engine 1624 executes integer and memory operations, without limitation. In at least one embodiment, floating point engine 1626 executes floating point and vector operations, without limitation. In at least one embodiment, fetch-decode unit 1622 dispatches micro-instructions to a single execution engine in place of both integer execution engine 1624 and floating point execution engine 1626.

[0180] In at least one embodiment, each core 1620(i) has access to an L2 cache 1628(i) included in core 1620(i), where i is an integer representing a particular instance of core 1620. In at least one embodiment, each core 1620 included in core complex 1610(j) is connected to the other cores 1620 in core complex 1610(j) via an L3 cache 1630(j) included in core complex 1610(j), where j is an integer representing a particular instance of core complex 1610. In at least one embodiment, cores 1620 included in core complex 1610(j) have access to all L3 caches 1630(j) included in core complex 1610(j), where j is an integer representing a particular instance of core complex 1610. In at least one embodiment, L3 cache 1630 can include, without limitation, any number of slices.

[0181] In at least one embodiment, fabric 1660 is a system interconnect that facilitates data and control transfers across core complexes 1610(1)-1610(N) (where N is an integer greater than zero), I / O interface 1670, and memory controllers 1680. In at least one embodiment, CPU 1600 can include, without limitation, any number and type of system interconnects in addition to or instead of fabric 1660 that facilitate data and control transfers across any number and type of directly or indirectly linked components that can be internal or external to CPU 1600. In at least one embodiment, I / O interface 1670 represents any number and type of I / O interfaces (e.g., PCI, PCI-X, PCIe, GBE, USB, etc.). In at least one embodiment, various types of peripheral devices are coupled to I / O interface 1670. In at least one embodiment, peripheral devices coupled to I / O interface 1670 can include, without limitation, a display, a keyboard, a mouse, a printer, a scanner, a joystick or other types of game controller, a media recording device, an external storage device, a network interface card, etc.

[0182] In at least one embodiment, memory controllers 1680 facilitate data transfers between CPU 1600 and system memory 1690. In at least one embodiment, core complexes 1610 and graphics complex 1640 share system memory 1690. In at least one embodiment, CPU 1600 implements a memory subsystem that includes, without limitation, any number and type of memory controllers 1680 and memory devices that can be dedicated to one component or shared among multiple components. In at least one embodiment, CPU 1600 implements a cache subsystem that includes, without limitation, one or more cache memories (e.g., L2 cache 1628 and L3 cache 1630), each of which can be private to a component or shared among any number of components (e.g., cores 1620 and core complexes 1610).

[0183] Figure 17An exemplary accelerator integration slice 1790 is shown in accordance with at least one embodiment. As used herein, a “slice” includes a specified portion of processing resources of an accelerator integration circuit. In at least one embodiment, an accelerator integration circuit provides cache management, memory access, environment management, and interrupt management services on behalf of a number of graphics processing engines of a graphics acceleration module. Graphics processing engines can each comprise a separate GPU. Alternatively, graphics processing engines can include different types of graphics processing engines within a GPU, such as graphics execution units, media processing engines (e.g., video encoders / decoders), samplers, and blit engines. In at least one embodiment, a graphics acceleration module can be a GPU with a plurality of graphics processing engines. In at least one embodiment, a graphics processing engine can be an individual GPU integrated on a common package, line card, or chip, as at least one embodiment.

[0184] In at least one embodiment, with respect to Figure 17 At least one component shown or described is used in implementing the techniques and / or functionality described herein. In at least one embodiment, the at least one embodiment includes one or more secure devices having information that will be used to detect whether the secure device has been re-started indicating that the one or more secure devices have been tampered with. Figures 1-8

[0185] Application effective address space 1782 within system memory 1714 stores process elements 1783. In one embodiment, process elements 1783 are stored in response to GPU invocations 1781 from applications 1780 executing on processor 1707. Process elements 1783 contain processing state for a respective application 1780. A work descriptor (WD) 1784 contained in process element 1783 can be a single job requested by an application or can contain pointers to a queue of jobs. In at least one embodiment, WD 1784 is a pointer to a job request queue in application effective address space 1782.

[0186] Graphics acceleration module 1746 and / or individual graphics processing engines can be shared by all or a subset of processes in a system. In at least one embodiment, a process can include infrastructure for setting up processing state and sending WDs 1784 to graphics acceleration module 1746 to start jobs in a virtualized environment.

[0187] In at least one embodiment, a dedicated process programming model is implemented for an accelerator integration circuit. In this model, a single process owns a graphics acceleration module 1746 or individual graphics processing engines. As graphics acceleration module 1746 is owned by a single process, a hypervisor initializes the accelerator integration circuit for the owning partition and an operating system initializes the accelerator integration circuit for the owning partition when assigning graphics acceleration module 1746.​

[0188] During operation, the WD fetch unit 1791 in the accelerator integrated slice 1790 fetches the next WD 1784, which includes instructions for the work to be performed by one or more graphics processing engines of the graphics acceleration module 1746. Data from the WD 1784 can be stored in register 1745 and used by the memory management unit (MMU) 1739, interrupt management circuitry 1747, and / or environment management circuitry 1748, as shown. For example, one embodiment of the MMU 1739 includes segment / page roaming circuitry for accessing segment / page tables 1786 within the OS virtual address space 1785. The interrupt management circuitry 1747 can handle interrupt events (INT) 1792 received from the graphics acceleration module 1746. When performing graph operations, the effective address 1793 generated by the graphics processing engine is translated into an actual address by the MMU 1739.

[0189] In one embodiment, the same register set 1745 is copied for each graphics processing engine and / or graphics acceleration module 1746 and can be initialized by the hypervisor or operating system. Each of these copied registers can be included in the accelerator integration slice 1790. Exemplary registers that can be initialized by the hypervisor are shown in Table 1.

[0190] Table 1 - Registers for Supervisor Initialization

[0191] 1 Slice control register 2 Real address (RA) plan's processing region pointer 3 Privilege mask override register 4 Interrupt vector table input offset 5 Interrupt vector table entry limit 6 Status register 7 Logical partition ID 8 Real address (RA) hypervisor accelerator utilization record pointer 9 Storage description register

[0192] Table 2 shows exemplary registers that can be initialized by the operating system.

[0193] Table 2 - Operating System Initialization Registers

[0194] 1 Process and thread identification 2 Effective address (EA) environment save / restoration pointer 3 Virtual address (VA) accelerator utilization record pointer 4 Virtual address (VA) storage segment table pointer 5 Privilege mask 6 Work descriptor

[0195] In one embodiment, each WD 1784 is specific to a particular graphics acceleration module 1746 and / or a particular graphics processing engine. It contains all the information required for the graphics processing engine to perform its work or to do so, or it may be a pointer to a memory location where the application has established a command queue for the work to be done.

[0196] Figure 18A and Figure 18B An exemplary graphics processor according to at least one embodiment herein is illustrated. In at least one embodiment, any exemplary graphics processor may be manufactured using one or more IP cores. In addition to the illustrations, other logic and circuitry may be included in at least one embodiment, including additional graphics processor / cores, peripheral interface controllers, or general-purpose processor cores. In at least one embodiment, the exemplary graphics processor is used within a System-on-a-Chip (SoC).

[0197] Figure 18A An exemplary graphics processor 1810 of an SoC integrated circuit is shown that can be fabricated using one or more IP cores, in accordance with at least one embodiment. Figure 18B An additional exemplary graphics processor 1840 of an SoC integrated circuit is shown that can be fabricated using one or more IP cores, in accordance with at least one embodiment. In at least one embodiment, Figure 18A The graphics processor 1810 is a low power graphics processor core. In at least one embodiment, Figure 18B The graphics processor 1840 is a higher performance graphics processor core. In at least one embodiment, each graphics processor 1810, 1840 can be Figure 13 A variant of the graphics processor 1310.

[0198] In at least one embodiment, with respect to Figure 18A and Figure 18B At least one component shown or described with respect to Figures 1-8 is used to implement techniques and / or functionality described in connection with

[0199] In at least one embodiment, the graphics processor 1810 includes a vertex processor 1805 and one or more fragment processors 1815A-1815N (e.g., 1815A, 1815B, 1815C, 1815D, through 1815N-1, and 1815N). In at least one embodiment, the graphics processor 1810 can execute different shader programs via separate logical

[0200] In at least one embodiment, the graphics processor 1810 additionally includes one or more MMUs 1820A-1820B, caches 1825A-1825B, and circuit interconnects 1830A-1830B. In at least one embodiment, one or more MMUs 1820A-1820B provide a virtual-to-physical address mapping for the graphics processor 1810, including for the vertex processor 1805 and / or fragment processors 1815A-1815N, which can reference vertex or image / texture data stored in memory, in addition to vertex or image / texture data stored in one or more caches 1825A-1825B. In at least one embodiment, one or more MMUs 1820A-1820B can be synchronized with other MMUs within the system, including with... Figure 13 One or more application processors 1305, graphics processors 1315, and / or video processors 1320 are associated with one or more MMUs, enabling each processor 1305-1320 to participate in a shared or unified virtual memory system. In at least one embodiment, one or more circuit interconnects 1830A-1830B enable the graphics processor 1810 to connect to other IP cores within the SoC via the SoC's internal bus or via a direct connection.

[0201] In at least one embodiment, the graphics processor 1840 includes Figure 18A The graphics processor 1810 includes one or more MMUs 1820A-1820B, caches 1825A-1825B, and circuit interconnects 1830A-1830B. In at least one embodiment, the graphics processor 1840 includes one or more shader cores 1855A-1855N (e.g., 1855A, 1855B, 1855C, 1855D, 1855E, 1855F, to 1855N-1 and 1855N) that provide a unified shader core architecture, wherein a single core or type of core can execute all types of programmable shader code, including shader program code for implementing vertex shaders, fragment shaders, and / or compute shaders. In at least one embodiment, the number of shader cores may vary. In at least one embodiment, the graphics processor 1840 includes an inter-core task manager 1845 that acts as a thread dispatcher to assign execution threads to one or more shader cores 1855A-1855N and a tile unit 1858 to accelerate tile-based rendering operations, wherein rendering operations of a scene are subdivided in image space, for example, to take advantage of local spatial consistency within the scene or to optimize the use of internal caches.

[0202] Figure 19AA graphics core 1900 according to at least one embodiment is shown. In at least one embodiment, graphics core 1900 can be included within graphics processor 1310 of FIG. 13. In at least one embodiment, graphics core 1900 can be a unified shader core 1855A-1855N in Figure 13 Figure 18B In at least one embodiment, graphics core 1900 includes shared instruction cache 1902, texture unit 1918, and cache / shared memory 1920, which are common to execution resources within graphics core 1900. In at least one embodiment, graphics core 1900 can include multiple slices 1901A-1901N or partitions of each core, and graphics processor can include multiple instances of graphics core 1900. Slices 1901A-1901N can include support logic including a local instruction cache 1904A-1904N, a thread scheduler 1906A-1906N, a thread dispatcher 1908A-1908N, and a set of registers 1910A-1910N. In at least one embodiment, slices 1901A-1901N can include a set of additional functional units (AFUs) 1912A-1912N, floating point units (FPUs) 1914A-1914N, integer arithmetic logic units (ALUs) 1916A-1916N, address computation units (ACUs) 1913A-1913N, double precision floating point units (DPFPUs)

[0203] 1915A-1915N, and matrix processing units (MPUs) 1917A-1917N. In at least one embodiment, graphics core 1900 is referred to as a compute unit or processing unit.

[0204] In at least one embodiment, at least one component shown or described with respect to Figure 19A In at least one embodiment, at least one component shown or described with respect to Figures 1-8 In at least one embodiment, at least one component shown or described with respect to

[0205] ​In one embodiment, FPUs 1914A-1914N can perform single-precision (32-bit) and half-precision (16-bit) floating point operations, while DPFPUs 1915A-1915N can perform double-precision (64-bit) floating point operations. In at least one embodiment, ALUs 1916A-1916N can perform variable precision integer operations at 8-bit, 16-bit, and 32-bit precision, and can be configured for mixed precision operations. In at least one embodiment, MPUs 1917A-1917N can also be configured for mixed precision matrix operations, including half-precision floating point and 8-bit integer operations. In at least one embodiment, MPUs 1917A-1917N can perform various matrix operations to accelerate CUDA programs, including enabling support for accelerated General Matrix to Matrix multiplication (GEMM). In at least one embodiment, AFUs 1912A-1912N can perform additional logical operations not supported by floating point or integer units, including trigonometric operations (e.g., Sine, Cosine, etc.).

[0206] Figure 19B A general purpose graphics processing unit (GPGPU) 1930 is shown in at least one embodiment. In at least one embodiment, GPGPU 1930 is highly parallel and suitable for deployment on a multi-chip module. In at least one embodiment, GPGPU 1930 can be configured to enable highly parallel compute operations to be performed by a GPU array. In at least one embodiment, GPGPU 1930 can be directly linked to other instances of GPGPU 1930 to create a multi-GPU cluster to improve execution time for CUDA programs. In at least one embodiment, GPGPU 1930 includes a host interface 1932 to enable connection to a host processor. In at least one embodiment, host interface 1932 is a PCIe interface. In at least one embodiment, host interface 1932 can be a vendor-specific communications interface or communications fabric. In at least one embodiment, GPGPU 1930 receives commands from a host processor to execute via global scheduler 1934 to dispatch execution threads associated with those commands to a group of compute clusters 1936A-1936H. In at least one embodiment, compute clusters 1936A-1936H share a cache memory 1938. In at least one embodiment, cache memory 1938 can be used as an upper level cache for cache memory within compute clusters 1936A-1936H.

[0207] In at least one embodiment, GPGPU 1930 comprises memory 1944A-1944B coupled to compute clusters 1936A-1936H via a set of memory controllers 1942A-1942B. In at least one embodiment, memory 1944A-1944B can include various types of memory devices including dynamic random access memory (DRAM) or graphics random access memory, such as synchronous graphics random access memory (SGRAM), including graphics double data rate (GDDR) memory.

[0208] In at least one embodiment, compute clusters 1936A-1936H each include a group of graphics cores, such as graphics core 1900, which can include multiple types of integer and floating point logic units that can perform computational operations at various precisions, including suitable for computations related to CUDA programs. For example, in at least one embodiment, at least a subset of floating point units in each compute cluster 1936A-1936H can be configured to perform 16- or 32-bit floating point operations, while a different subset of floating point units can be configured to perform 64-bit floating point operations. Figure 19A

[0209] In at least one embodiment, multiple instances of GPGPU 1930 can be configured to operate as compute clusters. Compute clusters 1936A-1936H can implement any technically feasible communication technology for synchronization and data exchange. In at least one embodiment, multiple instances of GPGPU 1930 communicate over host interface 1932. In at least one embodiment, GPGPU 1930 includes I / O hub 1939 that couples GPGPU 1930 with GPU link 1940, enabling a direct connection to other instances of GPGPU 1930. In at least one embodiment, GPU link 1940 is coupled to a specialized GPU-to-GPU bridge that enables communication and synchronization between multiple instances of GPGPU 1930. In at least one embodiment, GPU link 1940 is coupled with a high-speed interconnect to transmit and receive data to other GPGPUs or parallel processors. In at least one embodiment, multiple instances of GPGPU 1930 are located in separate data processing systems and communicate over a network device accessible via host interface 1932. In at least one embodiment, GPU link 1940 can be configured to connect to a host processor, in addition to or in place of host interface 1932. In at least one embodiment, GPGPU 1930 can be configured to execute CUDA programs.

[0210] Figure 20A ​A parallel processor 2000 is shown in accordance with at least one embodiment. In at least one embodiment, various components of parallel processor 2000 can be implemented using one or more integrated circuit devices, such as programmable processor(s), application specific integrated circuits (ASICs), or FPGAs.

[0211] In at least one embodiment, with respect to Figure 20A At least one component shown or described is used to implement a method Figures 1-8 described. In at least one embodiment, the at least one embodiment includes one or more secure devices having information that will be used to detect whether the secure device has been restarted indicating that the one or more secure devices have been tampered with.

[0212] In at least one embodiment, parallel processor 2000 includes a parallel processing unit 2002. In at least one embodiment, parallel processing unit 2002 includes an I / O unit 2004 that enables communication with other devices. In at least one embodiment, I / O unit 2004 is directly connected to other devices. In at least one embodiment, I / O unit 2004 connects with other devices using a hub or switch interface, such as memory hub 2005. In at least one embodiment, connections between memory hub 2005 and I / O unit 2004 form a communication link. In at least one embodiment, I / O unit 2004 connects with a host interface 2006 and a memory crossbar switch 2016, where host interface 2006 receives commands directed to performing processing operations and memory crossbar switch 2016 receives commands directed to performing memory operations.

[0213] In at least one embodiment, when host interface 2006 receives a command buffer via I / O unit 2004, host interface 2006 can direct a work operation to execute those commands to front end 2008. In at least one embodiment, front end 2008 is coupled with scheduler 2010, which is configured to assign commands or other work items to processing array 2012. In at least one embodiment, scheduler 2010 ensures that processing array 2012 is properly configured and in an active state before assigning tasks to processing array 2012 of processing array 2012. In at least one embodiment, scheduler 2010 is implemented by firmware logic executing on a microcontroller. In at least one embodiment, microcontroller- implemented scheduler 2010 is configurable to perform complex scheduling and work distribution operations with both coarse and fine grain, enabling fast preemption and context switching of threads executing on processing array 2012. In at least one embodiment, host software can prove a workload for scheduling on processing array 2012 through one of a number of graphics processing doorbells. In at least one embodiment, workload can then be automatically distributed on processing array 2012 by scheduler 2010 logic within microcontroller including scheduler 2010.

[0214] In at least one embodiment, processing array 2012 can include up to “N” processing clusters (e.g., cluster 2014A, cluster 2014B, through cluster 2014N). In at least one embodiment, each cluster 2014A-2014N of processing array 2012 can execute a large number of concurrent threads. In at least one embodiment, scheduler 2010 can use various scheduling and / or work distribution algorithms to assign work to clusters 2014A-2014N of processing array 2012, which can vary depending on workload produced by each type of program or computation. In at least one embodiment, scheduling can be handled dynamically by scheduler 2010, or can be assisted in part by compiler logic during compilation of program logic configured for execution by processing array 2012. In at least one embodiment, different clusters 2014A-2014N of processing array 2012 can be allocated for processing different types of programs or for performing different types of computations.

[0215] In at least one embodiment, processing array 2012 can be configured to perform various types of parallel processing operations. In at least one embodiment, processing array 2012 is configured to perform general purpose parallel compute operations. For example, in at least one embodiment, processing array 2012 can include logic to perform processing tasks including filtering of video and / or audio data, performing modeling operations including physical operations, and performing data transformations.

[0216] In at least one embodiment, processing array 2012 is configured to perform parallel graphics processing operations. In at least one embodiment, processing array 2012 can include additional logic to support performance of such graphics processing operations, including but not limited to texture sampling logic to perform texture operations, tessellation logic, and other vertex processing logic. In at least one embodiment, processing array 2012 can be configured to execute shader programs associated with graphics processing, such as but not limited to vertex shaders, tessellation shaders, geometry shaders, and pixel shaders. In at least one embodiment, parallel processing unit 2002 can transfer data to be processed from a system memory over I / O unit 2004. In at least one embodiment, data transferred can be stored to on-chip memory (e.g., parallel processor memory 2022) during processing, and written back to system memory upon completion of processing.

[0217] In at least one embodiment, when parallel processing unit 2002 is used to perform graphics processing, scheduler 2010 can be configured to divide the processing workload into approximately equal sized tasks, to better enable distribution of graphics processing operations across multiple clusters 2014A-2014N of processing array 2012. In at least one embodiment, portions of processing array 2012 can be configured to perform different types of processing. For example, in at least one embodiment, a first portion can be configured to perform vertex shading and topology generation, a second portion can be configured to perform tessellation and geometry shading, and a third portion can be configured to perform pixel shading or other screen space operations, to produce a rendered image for display on a display device. In at least one embodiment, intermediate data produced by one or more of clusters 2014A-2014N can be stored in buffers to allow transmission of the intermediate data between clusters 2014A-2014N for further processing.

[0218] In at least one embodiment, processing array 2012 can receive processing tasks to be executed via scheduler 2010, which receives commands defining the processing tasks from front end 2008. In at least one embodiment, a processing task can include an index to data to be processed, which can include surface (patch) data, raw data, vertex data, and / or pixel data, for example, as well as state parameters and commands defining how the data is to be processed (e.g., what program is to be executed). In at least one embodiment, scheduler 2010 can be configured to fetch the index corresponding to a task, or can receive the index from front end 2008. In at least one embodiment, front end 2008 can be configured to ensure that processing array 2012 is configured in an effective state before launching a workload specified by an incoming command buffer (e.g., a batch-buffer, a push buffer, etc.).

[0219] In at least one embodiment, each of one or more instances of parallel processing unit 2002 can be coupled to a parallel processor memory 2022. In at least one embodiment, parallel processor memory 2022 can be accessed by the processing array 2012, as well as the I / O unit 2004, via a memory crossbar 2016. In at least one embodiment, memory crossbar 2016 can be used to access memory banks 2024A-2024N of parallel processor memory 2022 by the processing array 2012, I / O unit 2004, and possibly other processing units or processing arrays in at least one embodiment. In at least one embodiment, central hub 2010 can be used to distribute commands and data to various units of the parallel processing unit 2002. In at least one embodiment, one or more instances of central hub 2010 can be coupled to a memory controller 2018 that can control access to memory crossbar 2016, and by extension, parallel processor memory 2022, via one or more memory ports. In at least one embodiment, memory crossbar 2016 can be configured to provide parallel memory access to multiple memory ports using separate bus channels.

[0220] In at least one embodiment, memory units 2024A-2024N can include various types of memory devices including dynamic random access memory (DRAM) or graphics random access memory, such as synchronous graphics random access memory (SGRAM). In at least one embodiment, memory units 2024A-2024N can also include 3D stacked memory including but not limited to high bandwidth memory (HBM). In at least one embodiment, rendering targets such as frame buffers or texture maps can be stored across memory units 2024A-2024N allowing partition units 2020A-2020N to write portions of each rendering target in parallel to effectively use available bandwidth of parallel processor memory 2022. In at least one embodiment, local instances of parallel processor memory 2022 can be excluded in favor of a unified memory design that utilizes system memory in combination with local cache memory.

[0221] In at least one embodiment, any of clusters 2014A-2014N of processing array 2012 can process data that is to be written into any of memory units 2024A-2024N within parallel processor memory 2022. In at least one embodiment, memory crossbar 2016 can be configured to transmit outputs of each cluster 2014A-2014N to any partition unit 2020A-2020N or another cluster 2014A-2014N, which can perform further processing operations on the outputs. In at least one embodiment, each cluster 2014A-2014N can communicate with memory interface 2018 through memory crossbar 2016 to read from or write to various external memory devices. In at least one embodiment, memory crossbar 2016 has a connection to memory interface 2018 to communicate with I / O unit 2004, as well as a local instance of parallel processor memory 2022, to enable processing elements within different processing clusters 2014A-2014N to communicate with system memory or other memories not local to the parallel processing units 2002. In at least one embodiment, memory crossbar 2016 can use virtual channels to separate traffic flows among clusters 2014A-2014N and partition units 2020A-2020N.

[0222] In at least one embodiment, multiple instances of parallel processing unit 2002 can be provided on a single add-in card, or multiple add-in cards can be interconnected. In at least one embodiment, different instances of parallel processing unit 2002 can be configured to operate in cooperation with one another to enable parallel processing operations. Although

[0223] Figure 20B Processing cluster 2094 is shown according to at least one embodiment. In at least one embodiment, processing cluster 2094 is included in a parallel processing unit. In at least one embodiment, processing cluster 2094 is a Figure 20AAn example of one of the processing clusters 2014A-2014N. In at least one embodiment, the processing cluster 2094 can be configured to execute a number of threads in parallel, wherein the term "thread" refers to an instance of a specific program executing on a particular set of input data. In at least one embodiment, Single Instruction Multiple Data (SIMD) instruction issuing technology is used to support the parallel execution of a large number of threads without providing multiple independent instruction units. In at least one embodiment, Single Instruction Multiple Threading (SIMT) technology is used to support the parallel execution of a large number of generally synchronous threads, which uses a common instruction unit configured to issue instructions to a set of processing engines within each processing cluster 2094.

[0224] In at least one embodiment, regarding Figure 20B At least one component shown or described is used to implement the combination. Figures 1-8 The described technologies and / or functions. In at least one embodiment, the at least one embodiment includes one or more security devices having information for detecting whether the security device has been restarted, thereby indicating that the one or more security devices have been tampered with.

[0225] In at least one embodiment, the operation of the processing cluster 2094 can be controlled by a pipeline manager 2032 that assigns processing tasks to the SIMT parallel processors. In at least one embodiment, the pipeline manager 2032... Figure 20A The scheduler 2010 receives instructions and manages the execution of these instructions via the graphics multiprocessor 2034 and / or texture unit 2036. In at least one embodiment, the graphics multiprocessor 2034 is an exemplary instance of a SIMT parallel processor. However, in at least one embodiment, the processing cluster 2094 may include various types of SIMT parallel processors with different architectures. In at least one embodiment, the processing cluster 2094 may include one or more instances of the graphics multiprocessor 2034. In at least one embodiment, the graphics multiprocessor 2034 can process data, and the data crossover switch 2040 can be used to distribute the processed data to one of a number of possible destinations (including other shader units). In at least one embodiment, the pipeline manager 2032 can facilitate the distribution of processed data by specifying the destination of the processed data to be distributed via the data crossover switch 2040.

[0226] In at least one embodiment, each graphics multiprocessor 2034 within processing cluster 2094 can include an identical set of functional execution logic (e.g., arithmetic logic units, load store units (LSUs), etc.). In at least one embodiment, functional execution logic can be configured in a pipelined manner in which new instructions can be issued before previous instructions are complete. In at least one embodiment, functional

[0227] In at least one embodiment, instructions transmitted to processing cluster 2094 form a thread for execution. In at least one embodiment, a set of threads executing across a set of parallel processing engines forms a warp. In at least one embodiment, threads within a warp are aligned by size and execute the same program, but can be in different states. In at least one embodiment, all threads within a warp are executed coherently by one or more of processing engines within graphics multiprocessor 2034. In at least one embodiment, all threads of a warp are executed coherently, meaning that the state of each thread is managed and governed in a manner that takes into consideration the state of each of the other threads within the same warp.

[0228] In at least one embodiment, graphics multiprocessor 2034 includes internal cache memory to perform load and store operations. In at least one embodiment, graphics multiprocessor 2034 can bypass internal cache and use register file 2036 within processing cluster 2094 for load and store operations. In at least one embodiment, register file 2036 can be on- chip memory, shared with graphics processing cluster 2094. Figure 20AL2 cache within partition units 2020A-2020N that are shared among all processing clusters 2094 and can be used to transmit data between threads. In at least one embodiment, graphics multiprocessor 2034 can also have access to off-chip global memory, which can include one or more of a local parallel processor memory and / or a system memory. In at least one embodiment, any memory external to parallel processor unit 2002 can be accessed by graphics multiprocessor 2034. In at least one embodiment, processing clusters 2094 include multiple instances of graphics multiprocessor 2034 that share common instructions and data, which can be stored in Ll cache 2048.

[0229] In at least one embodiment, each processing cluster 2094 can include an MMU 2045 that is configured to translate virtual addresses into physical addresses, in at least one embodiment, one or more instances of MMU 2045 can reside within memory interface 2018. In at least one embodiment, MMU 2045 includes a set of page table entries (PTEs) that are used to map Figure 20A virtual addresses to physical addresses at a granularity of one or more cache lines. In at least one embodiment, MMU 2045 can translate virtual addresses to physical addresses as part of page walk requests that are initiated by software. In at least one embodiment, a page walk is a process by which a system maps a virtual address to a physical address. In at least one embodiment, software may

[0230] In at least one embodiment, processing cluster 2094 can be configured such that each graphics multiprocessor 2034 is coupled to a texture unit 2036 for performing texture mapping operations, e.g., determining texture sample positions, reading texture data, and filtering texture data. In at least one embodiment, texture data is read from an internal texture Ll cache (not shown) or from an Ll cache within graphics multiprocessor 2034 and stored in L2 cache, local parallel processor memory, or system memory, as needed, for processing operations. In at least one embodiment, each graphics multiprocessor 2034 outputs processed tasks to data crossbar 2040 to provide processed tasks to another processing cluster 2094 for further processing or to store processed task data in an L2 cache, local parallel processor memory, or system memory via memory crossbar 2016. In at least one embodiment, pre-ROP 2042 is configured to receive data from graphics multiprocessor 2034, direct data to a ROP unit which can be located formally within graphics multiprocessor 2034 or which can be part of on-chip system memory, and to perform address translations upon write operations. In at least one embodiment, pre-ROP 2042 can handle vertex data, execution of vertex shaders and geometry processing. Figure 20A

[0231] Figure 20C A graphics multiprocessor 2096, according to at least one embodiment, is shown. In at least one embodiment, graphics multiprocessor 2096 is a graphics multiprocessor 2034 of Figure 20B FIG. 20A-20N. In at least one embodiment, graphics multiprocessor 2096 is coupled with a pipeline manager 2032 of a processing cluster 2094. In at least one embodiment, graphics multiprocessor 2096 has a graphics processing pipeline that includes, without limitation, an instruction cache 2052, an instruction unit 2054, an address mapping unit 2056, a register file 2058, one or more GPGPU cores 2062, and one or more LSUs 2066. GPGPU cores 2062 and LSUs 2066 are coupled with cache memory 2072 and shared memory 2070 via a memory and cache interconnect 2068.

[0232] ​In at least one embodiment, instruction cache 2052 receives a stream of instructions 2050 to be executed by graphics processing engine 2010 from pipeline manager 2032. In at least one embodiment, instructions 2050 are cached in instruction cache 2052 and dispatched for execution by instruction unit 2054. In one embodiment, instruction unit 2054 can dispatch instructions to threads assigned to different ones of GPGPU cores 2062 as thread groups, where each thread can be assigned to a different one of GPGPU cores 2062. In at least one embodiment, instructions can be accessed from an unified address space within a single program by each processor and thread, for example, within a single chip.

[0233] In at least one embodiment, register file 2058 provides a set of registers for functional units of graphics processing engine 2096. In at least one embodiment, register file 2058 provides temporary storage for operands of the data

[0234] In at least one embodiment, GPGPU cores 2062 can each include FPUs and / or ALUs for performing instructions of graphics processing engine 2096. GPGPU cores 2062 can be similar in architecture or can differ in architecture. In at least one embodiment, a first portion of GPGPU cores 2062 includes single precision FPUs and integer ALUs, while a second portion of GPGPU cores 2062 includes double precision FPUs. In at least one embodiment, FPUs can implement IEEE 754-2008 standard for floating point arithmetic or enable variable precision floating point arithmetic. In at least one embodiment, graphics processing engine 2096 can additionally include one or more fixed function or special function units to perform specific functions or

[0235] In at least one embodiment, GPGPU cores 2062 include SIMD logic capable of performing a single instruction on multiple sets of data. In at least one embodiment, GPGPU cores 2062 can physically execute SIMD 4, SIMD 8, and SIMD 9 instructions and logically execute a SIMD 1, SIMD 2, and SIMD 32 instructions. In at least one embodiment, SIMD instructions for GPGPU cores can be generated by a shader compiler during compilation of code to be executed by GPGPU cores 2062 or automatically generated by a driver program or execution system. In at least one embodiment, multiple threads can be executed via a single SIMD instruction by GPGPU cores 2062 configured for a single instruction multiple thread (SIMT) execution model. For example, in at least one embodiment, a first eight threads can be executed via a first SIMD 8 instruction and a second eight threads can be executed via a second SIMD 8 instruction.

[0236] In at least one embodiment, memory and cache interconnect 2068 is an interconnect network that connects each functional unit of graphics multiprocessor 2096 to register file 2058 and shared memory 2070. In at least one embodiment, memory and cache interconnect 2068 is a crossbar interconnect that allows LSUs 2066 to perform load and store operations between shared memory 2070 and register file 2058. In at least one embodiment, register file 2058 can operate at same frequency as GPGPU cores 2062 so that there is very little latency in data transfers between GPGPU cores 2062 and register file 2058. In at least one embodiment, shared memory 2070 can be used to enable

[0237] In at least one embodiment, parallel processor or GPGPU as described herein is communicatively coupled to host / processor cores to accelerate graphics operations, machine learning operations, pattern analysis operations, and various general purpose GPU (GPGPU) functions. In at least one embodiment, GPU can be communicatively coupled to host processor / cores over a bus or other interconnect (e.g., a high-speed

[0238] Figure 21 A graphics processor 2100 according to at least one embodiment is shown. In at least one embodiment, graphics processor 2100 includes a ring interconnect 2102, a front-end 2104, a media engine 2137, and graphics cores 2180A-2180N. In at least one embodiment, ring interconnect 2102 couples graphics processor 2100 to other processing units, including other graphics processors or one or more general-purpose processor cores. In at least one embodiment, graphics processor 2100 is one of a number of processors integrated within a multi-core processing system.

[0239] In at least one embodiment, with respect to Figure 21 At least one component shown or described is used to implement a technique and / or functionality described Figures 1-8 In at least one embodiment, the at least one embodiment includes one or more secure devices having information that will be used to detect whether the secure device has been re-booted, indicating that the one or more secure devices have been tampered with.

[0240] In at least one embodiment, graphics processor 2100 receives batches of commands via ring interconnect 2102. In at least one embodiment, incoming commands are interpreted by a command streamer 2103 in pipeline front-end 2104. In at least one embodiment, graphics processor 2100 includes scalable execution logic to perform 3D geometry processing and media processing via graphics core(s) 2180A-2180N. In at least one embodiment, for 3D geometry processing commands, command streamer 2103 supplies commands to geometry pipeline 2136. In at least one embodiment, for at least some media processing commands, command streamer 2103 supplies commands to video front end 2134, which couples with a media engine 2137. In at least one embodiment, media engine 2137 includes a video quality engine (VQE) 2130 for video and image post-processing, and a multi-format encode / decode (MFX) 2133 engine to

[0241] In at least one embodiment, graphics processor 2100 includes a scalable thread execution resource characterized by modular graphics cores 2180A-2180N (sometimes referred to as core slices), each with multiple sub-cores 2150A-2150N, 2160A-2160N (sometimes referred to as core sub-slices). In at least one embodiment, graphics processor 2100 can have any number of graphics cores 2180A-2180N. In at least one embodiment, graphics processor 2100 includes graphics core 2180A having at least a first sub-core 2150A and a second sub-core 2160A. In at least one embodiment, graphics processor 2100 is a low power processor with a single sub-core (e.g., 2150A). In at least one embodiment, graphics processor 2100 includes multiple graphics cores 2180A-2180N, each including a set of first sub-cores 2150A-2150N and a set of second sub-cores 2160A-2160N. In at least one embodiment, each sub-core in first sub-cores 2150A-2150N includes at least a first set of execution units (EUs) 2152A-2152N and a media / texture sampler 2154A-2154N. In at least one embodiment, each sub-core in second sub-cores 2160A-2160N includes at least a second set of execution units 2162A-2162N and a sampler 2164A-2164N. In at least one embodiment, each sub-core 2150A-2150N, 2160A-2160N shares a set of shared resources 2170A-2170N. In at least one embodiment, shared resources include shared cache memory and pixel operation logic.

[0242] Figure 22 A processor 2200 is shown in accordance with at least one embodiment. In at least one embodiment, processor 2200 can include, without limitation, a logic circuit that executes instructions. In at least one embodiment, processor 2200 can execute instructions including x86 instructions, ARM instructions, specialized instructions for ASICs, and the like. In at least one embodiment, processor 2210 can include registers to store packed data, for example 64-bit wide MMX™ registers as enabled in microprocessors employing MMX technology by Intel Corporation of Santa Clara, California. In at least one embodiment, MMX registers available in integer and floating point form can operate with packed data elements that accompany SIMD and Streaming SIMD Extensions (“SSE”) instructions. In at least one embodiment, 128-bit wide XMM registers related to SSE2, SSE3, SSE4, AVX, or higher (“SSEx” in general) technology can hold such packed data operands. In at least one embodiment, processor 2210 can execute instructions to accelerate CUDA programs.

[0243] In at least one embodiment, with respect to Figure 22 At least one component shown or described is used to implement a technique and / or functionality described Figures 1-8 In at least one embodiment, the at least one embodiment includes one or more secure devices having information that will be used to detect whether the secure device has been restarted indicating that the one or more secure devices have been tampered with.

[0244] In at least one embodiment, processor 2200 includes an in-order front-end (“front-end”) 2201 to fetch instructions to execute and prepare instructions to use in pipeline stages later in processor. In at least one embodiment, front-end 2201 can include several units. In at least one embodiment, instruction prefetcher 2226 fetches instructions from memory and provides instructions to instruction decoder 2228 which in turn decodes or interprets instructions. For example, in at least one embodiment, instruction decoder 2228 decodes a received instruction into one or more operations called “micro-instructions” or “micro-operations” (also called “micro ops” or “uops”) for execution. In at least one embodiment, instruction decoder 2228 parses the instruction into an operation code that it uses to identify an operation (e.g., a load, a store, or a jump) and corresponding data and control fields that are used to execute the instruction. In at least one embodiment, a trace cache 2230 can assemble decoded micro ops into program ordered sequences or traces in a micro op queue 2234 for execution. In at least one embodiment, when trace cache 2230 encounters a complex instruction, a microcode ROM 2232 provides micro ops needed to complete the operation.

[0245] In at least one embodiment, some instructions can be converted into a single micro op, while others can need several micro ops to complete the full operation. In at least one embodiment, if more than four micro ops are needed to complete an instruction, then instruction decoder 2228 can access microcode ROM 2232 to execute the instruction. In at least one embodiment, instructions can be decoded into a small number of micro ops for processing at instruction decoder 2228. In at least one embodiment, if multiple micro ops are needed to complete an operation, then the instruction can be stored in microcode ROM 2232. In at least one embodiment, trace cache 2230 references a entry point programmable logic array (“PLA”) to determine the correct micro op pointer for reading microcode sequences from microcode ROM 2232 to complete one or more instructions, in at least one embodiment. In at least one embodiment, after microcode ROM 2232 completes sequencing of micro ops for an instruction, front-end 2201 of the machine can resume fetching micro ops from trace cache 2230.

[0246] In at least one embodiment, out-of-order execution engine (“out-of-order engine”) 2203 can prepare instructions for execution. In at least one embodiment, out-of-order execution logic has multiple buffers to smooth and reorder instruction flow to optimize performance as instructions are pipelined down and dispatched for execution. Out-of-order execution engine 2203 includes, without limitation, an allocator / register renamer 2240, a memory micro instruction queue 2242, an integer / floating point micro instruction queue 2244, a memory scheduler 2246, a fast scheduler 2202, a slow / general floating point scheduler (“slow / general FP scheduler”) 2204, and a simple floating point scheduler (“simple FP scheduler”) 2206. In at least one embodiment, fast scheduler 2202, slow / general floating point scheduler 2204, and simple floating point scheduler 2206 are also collectively referred to as “micro instruction schedulers 2202, 2204, 2206.” Allocator / register renamer 2240 allocates machine buffers and resources needed for each micro instruction to execute in order. In at least one embodiment, allocator / register renamer 2240 renames logical registers to entries in a register file. In at least one embodiment, allocator / register renamer 2240 also allocates entries for each micro instruction in one of two micro instruction queues, memory micro instruction queue 2242 for memory operations and integer / floating point micro instruction queue 2244 for non-memory operations, in front of memory scheduler 2246 and micro instruction schedulers 2202, 2204, 2206. In at least one embodiment, micro instruction schedulers 2202, 2204, 2206 determine when micro instructions are ready to execute based on readiness of their dependent input register operand sources and availability of execution resource micro instructions needed to complete. In at least one embodiment, fast scheduler 2202 of at least one embodiment can schedule on every half of a main clock cycle, while slow / general floating point scheduler 2204 and simple floating point scheduler 2206 can schedule once per main processor clock cycle. In at least one embodiment, micro instruction schedulers 2202, 2204, 2206 arbitrate for a dispatch port to dispatch micro instructions for execution.

[0247] In at least one embodiment, execution block 2211 includes, without limitation, an integer register file / bypass network 2208, a floating point register file / bypass network (“FP register file / bypass network”) 2210, address generation units (“AGUs”) 2212 and 2214, fast ALUs 2216 and 2218, a slow ALU 2220, a floating point ALU (“FP”) 2222, and a floating point move unit (“FP move”) 2224. In at least one embodiment, integer register file / bypass network 2208 and floating point register file / bypass network 2210 are also referred to herein as “register files 2208, 2210.” In at least one embodiment, AGUs 2212 and 2214, fast ALUs 2216 and 2218, slow ALU 2220, floating point ALU 2222, and floating point move unit 2224 are also referred to herein as “execution units 2212, 2214, 2216, 2218, 2220, 2222, and 2224.” In at least one embodiment, execution block can include, without limitation, any number (including zero) and type of register files, bypass networks, address generation units, and execution units (in any combination).

[0248] In at least one embodiment, register files 2208, 2210 can be arranged between microinstruction schedulers 2202, 2204, 2206 and execution units 2212, 2214, 2216, 2218, 2220, 2222, and 2224. In at least one embodiment, integer register file / bypass network 2208 performs integer operations. In at least one embodiment, floating point register file / bypass network 2210 performs floating point operations. In at least one embodiment, each of register files 2208, 2210 can include, without limitation, a bypass network that can bypass or forward a just-completed result that has not yet been written into a register file to a new dependee. In at least one embodiment, register files 2208, 2210 can communicate data with each other. In at least one embodiment, integer register file / bypass network 2208 can include, without limitation, two separate register files, one for lower 32 bits of data and a second for upper 32 bits of data. In at least one embodiment, floating point register file / bypass network 2210 can include, without limitation, 128 bit wide entries, as floating point instructions typically have operands that are 64 to 128 bits wide.

[0249] In at least one embodiment, execution units 2212, 2214, 2216, 2218, 2220, 2222, 2224 can execute instructions. In at least one embodiment, register files 2208, 2210 store integer and floating point data operand values upon which microinstructions require execution. In at least one embodiment, processor 2200 can include, without limitation, any number of execution units 2212, 2214, 2216, 2218, 2220, 2222, 2224 and combinations thereof. In at least one embodiment, floating point ALU 2222 and floating point move unit 2224 can execute floating point, MMX, SIMD, AVX and SSE, or other operations, including specialized machine learning instructions. In at least one embodiment, floating point ALU 2222 can include, without limitation, a 64 bit by 64 bit floating point divider to execute divide, square root, and remainder micro-ops. In at least one embodiment, instructions for dealing with floating point values can be handled with floating point hardware. In at least one embodiment, ALU operations can be passed to fast ALUs 2216, 2218. In at least one embodiment, fast ALUS 2216, 2218 can execute fast operations with an effective latency of half a clock cycle. In at least one embodiment, most complex integer operations enter slow ALU 2220 as slow ALU 2220 can include, without limitation, integer execution hardware for long latency type operations such as multiplies, shifts, flag logic, and branch processing. In at least one embodiment, memory load / store operations can be executed by AGUs 2212, 2214. In at least one embodiment, fast ALU 2216, fast ALU 2218, and slow ALU 2220 can execute integer operations on 64 bit data operands. In at least one embodiment, fast ALU 2216, fast ALU 2218, and slow ALU 2220 can be implemented to support a variety of data bit sizes including 16, 32, 128, 256, etc. In at least one embodiment, floating point ALU 2222 and floating point move unit 2224 can be implemented to support a range of operands having bits of various widths. In at least one embodiment, floating point ALU 2222 and floating point move unit 2224 can operate on 128 bits wide packed data operands in conjunction with SIMD and multimedia instructions.

[0250] In at least one embodiment, micro-instruction schedulers 2202, 2204, 2206 schedule dependent operations before parent load completion execution. In at least one embodiment, because micro-instructions can be speculatively scheduled and executed in processor 2200, processor 2200 can also include logic to handle memory misses. In at least one embodiment, if a data load in a data cache misses, there can be a dependent operation running in a pipeline that causes a scheduler to temporarily have incorrect data. In at least one embodiment, a replay mechanism tracks and re-executes instructions that use incorrect data. In at least one embodiment, dependent operations can need to be replayed and independent operations can be allowed to complete. In at least one embodiment, a scheduler and replay mechanism of at least one embodiment of a processor can also be designed to capture instruction sequences for text string compare operations.

[0251] In at least one embodiment, the term “register” can refer to an on-board processor storage location that can be used as part of an instruction that identifies an operand. In at least one embodiment, a register can be one that can be used from outside of a processor (from a programmer’s perspective). In at least one embodiment, a register can not be limited to a particular type of circuit. Rather, in at least one embodiment, a register can store data, provide data, and perform functions described herein. In at least one embodiment, registers described herein can be implemented by circuitry within a processor using a variety of different techniques, such as dedicated physical registers, physical registers allocated dynamically with register renaming, a combination of dedicated and dynamically allocated physical registers, etc. In at least one embodiment, an integer register stores 32 bits of integer data. A register file of at least one embodiment also contains eight multimedia SIMD registers for packing data.

[0252] Figure 23 Processor 2300 is shown in accordance with at least one embodiment. In at least one embodiment, processor 2300 includes, without limitation, one or more processor cores (cores)

[0253] 2302A-2302N, an integrated memory controller 2314, and an integrated graphics processor 2308. In at least one embodiment, processor 2300 can include additional cores up to and including an additional core 2302N represented by a dashed box. In at least one embodiment, each processor core 2302A-2302N includes one or more internal cache units 2304A-2304N. In at least one embodiment, each processor core can also access one or more shared cache units 2306. In at least one embodiment, one or more processor cores 2302A-2302N are referred to as one or more compute units or execution units.

[0254] In at least one embodiment, in connection with Figure 23 At least one component shown or described is used in implementing the techniques and / or functionality described Figures 1-8 described. In at least one embodiment, the at least one embodiment includes one or more secure devices having information that will be used to detect whether the secure device has been re-booted indicating that the one or more secure devices have been tampered with.

[0255] In at least one embodiment, internal cache units 2304A-2304N and shared cache unit 2306 represent a cache memory hierarchy within processor 2300. In at least one embodiment, cache memory units 2304A-2304N can include one or more levels of cache, such as L2, L3, 4-level (L4), or other levels of cache, within each processor core and a shared on-chip cache unit, e.g., L2, L3, L4, or other level(s) of cache, where the highest level of cache is classified as an LLC before off-chip memory. In at least one embodiment, cache coherence logic maintains coherence between various cache units 2306 and 2304A-2304N.

[0256] In at least one embodiment, processor 2300 can also include a set of one or more bus controller units 2316 and a system agent core 2310. In at least one embodiment, one or more bus controller units 2316 manage a set of peripheral buses, such as one or more PCI or PCI Express buses. In at least one embodiment, system agent core 2310 provides management functionality for various processor components. In at least one embodiment, system agent core 2310 includes one or more integrated memory controllers 2314 to manage access to various external memory devices (not shown), such as one or more dynamic random access memory DRAM or static RAM (SRAM) devices.

[0257] In at least one embodiment, one or more processor cores 2302A-2302N include support for simultaneous multi-threading. In at least one embodiment, system agent core 2310 includes components for coordinating and operating processor cores 2302A-2302N during multi-threaded processing. In at least one embodiment, system agent core 2310 can additionally include a power control unit (PCU), including logic and components to regulate one or more power states of processor cores 2302A-2302N and graphics processor 2308.

[0258] In at least one embodiment, processor 2300 additionally includes a graphics processor 2308 to perform graphics processing operations. In at least one embodiment, graphics processor 2308 couples with shared cache unit 2306 and system agent core 2310 including one or more integrated memory controllers 2314. In at least one embodiment, system agent core 2310 also includes a display controller 2311 to drive one or more coupled displays to display graphics processor output. In at least one embodiment, display controller 2311 can also be a separate module coupled with graphics processor 2308 via at least one interconnect, or can be integrated within graphics processor 2308.

[0259] In at least one embodiment, ring based interconnect unit 2312 is used to couple internal components of processor 2300. In at least one embodiment, alternative interconnect units can be used, such as a point-to-point interconnect, a switched interconnect, or other technology. In at least one embodiment, graphics processor 2308 couples with ring interconnect 2312 via I / O link 2313.

[0260] In at least one embodiment, I / O link 2313 represents at least one of a variety of I / O interconnects, including a package I / O interconnect facilitating communication between various processor components and a high performance embedded memory module 2318, such as an eDRAM module. In at least one embodiment, each of processor cores 2302A-2302N and graphics processor 2308 uses embedded memory module 2318 as a shared LLC.

[0261] In at least one embodiment, processor cores 2302A-2302N are homogenous cores executing a common instruction set architecture. In at least one embodiment, processor cores 2302A-2302N are heterogeneous in terms of ISA, with one or more processor cores 2302A-2302N executing a common instruction set while one or more other processor cores 2302A-2302N execute a subset of the common instruction set or a different instruction set. In at least one embodiment, processor cores 2302A-2302N are heterogeneous in terms of microarchitecture, with one or more cores having a relatively high power consumption coupled with one or more power cores having a lower power consumption. In at least one embodiment, processor 2300 can be implemented on one or more chips or as a SoC integrated circuit.

[0262] Figure 24A graphics processor core 2400 is shown in accordance with at least one embodiment described. In at least one embodiment, graphics processor core 2400 is included within a graphics core array. In at least one embodiment, graphics processor core 2400 (sometimes called a core slice) can be one or more graphics cores within a modular graphics processor. In at least one embodiment, graphics processor core 2400 is an example of one graphics core slice, and graphics processors described herein can include multiple graphics core slices based on target power and performance envelopes. In at least one embodiment, each graphics core 2400 can include fixed function block 2430, also known as a sub slice, that includes modules of general- purpose and fixed function logic for handling various compute and sampling tasks.

[0263] In at least one embodiment, with respect to Figure 24 At least one component shown or described is used in implementing techniques and / or functionality described Figures 1-8 In at least one embodiment, the at least one embodiment includes one or more secure devices having information that will be used to detect whether the secure device has been restarted indicating that the one or more secure devices have been tampered with.

[0264] In at least one embodiment, fixed function block 2430 includes geometry / fixed function pipeline 2436 that can be shared by all the sub-cores in graphics processor 2400, for example, in lower performance and / or lower power graphics processor implementations. In at least one embodiment, geometry / fixed function pipeline 2436 includes a 3D fixed function pipeline, a video front-end unit, a thread generator and thread dispatcher, and a unified return buffer manager that manages a unified return buffer.

[0265] In at least one embodiment, fixed function block 2430 also includes a graphics SoC interface 2437, a graphics microcontroller 2438, and a media pipeline 2439. Graphics SoC interface 2437 provides an interface between graphics core 2400 and other processor cores within a SoC integrated circuit. In at least one embodiment, graphics microcontroller 2438 is a programmable sub-processor that is configurable to manage various functions of graphics processor 2400, including thread dispatch, scheduling, and pre-emption. In at least one embodiment, media pipeline 2439 includes logic to facilitate decoding, encoding, pre-processing, and / or post-processing of multimedia data, including image and video data. In at least one embodiment, media pipeline 2439 implements media operations via requests to compute or sample logic within sub-cores 2401-2401F.

[0266] In at least one embodiment, SoC interface 2437 enables graphics core 2400 to communicate with general application processor cores (e.g., CPUs) and / or other components within the SoC, including memory hierarchy elements such as shared LLC cache, system RAM, and / or embedded on-chip or package DRAM. In at least one embodiment, SoC interface 2437 can also enable communication with fixed function devices (e.g., camera imaging pipeline) within the SoC, and to use and / or implement global memory atoms that can be shared between graphics core 2400 and the CPU(s) within the SoC. In at least one embodiment, SoC interface 2437 can also implement power management controls for graphics core 2400 and enable an interface between a clock domain of graphics core 2400 with other clock domains within the SoC. In at least one embodiment, SoC interface 2437 enables receiving command buffers from command streamer and global thread dispatcher, which are configured to provide commands and instructions to each of one or more graphics cores within a graphics processor. In at least one embodiment, commands and instructions can be dispatched to a media pipeline 2439 when media operations are to be performed, or to a geometry and fixed function pipeline (e.g., geometry and fixed function pipeline 2436, geometry and fixed function pipeline 2414) when graphics processing operations are to be performed.

[0267] In at least one embodiment, graphics microcontroller 2438 can be configured to perform various scheduling and management tasks for graphics core 2400. In at least one embodiment, graphics microcontroller 2438 can perform graphics and / or compute workload scheduling on various graphics processing engines within execution unit (EU) arrays 2402A-2402F, 2404A-2404F in sub-cores 2401A-2401F. In at least one embodiment, host software executing on CPU cores of an SoC including graphics core 2400 can submit workloads for one of graphics processor doorbell, which invokes scheduling operations on appropriate graphics engines. In at least one embodiment, scheduling operations include determining which workload to run next, submitting the workload to a command streamer, pre-empting existing workloads running on an engine, monitoring progress of the workload, and notifying host software when the workload is complete. In at least one embodiment, graphics microcontroller 2438 can also facilitate low power or idle state for graphics core 2400, providing graphics core 2400 with the ability to save and restore registers across low power state transitions independently of operating system and / or graphics driver software on the system.

[0268] In at least one embodiment, graphics core 2400 can have more or less than the illustrated sub-cores 2401 A-2401F, up to N modular sub-cores. For each set of N sub-cores, graphics core 2400 can also include, in at least one embodiment, shared function logic 2410, shared and / or cache memory 2412, geometry / fixed function pipeline 2414, and additional fixed function logic 2416 to accelerate various graphics and compute processing operations. In at least one embodiment, shared function logic 2410 can include logic units (e.g., samplers, math, and / or inter-thread communication logic) that are shareable by each N sub-core within graphics core 2400. Shared and / or cache memory 2412 can be an LLC for N sub-cores 2401A-2401F within graphics core 2400, and can also serve as shared memory accessible by multiple sub-cores. In at least one embodiment, geometry / fixed function pipeline 2414 can be included in place of geometry / fixed function pipeline 2436 within fixed function block 2430, and can include the same or similar logic units.

[0269] In at least one embodiment, graphics core 2400 includes additional fixed function logic 2416 that can include various fixed function acceleration logic used by graphics core 2400. In at least one embodiment, additional fixed function logic 2416 includes an additional geometry pipeline used for position-only shading. In position-only shading, there are at least two geometry pipelines, while in a full geometry pipeline and cull pipeline within geometry / fixed function pipeline 2416, 2436, which is an additional geometry pipeline that can be included in additional fixed function logic 2416. In at least one embodiment, the cull pipeline is a trimmed down version of the full geometry pipeline. In at least one embodiment, the full pipeline and the cull pipeline can execute different instances of an application, each with a separate environment. In at least one embodiment, position-only shading can hide long cull runs of triangles that are discarded, which can complete shading earlier in some cases. For example, in at least one embodiment, cull pipeline logic in additional fixed function logic 2416 can execute position shaders in parallel with a main application, and often generate critical results faster than the full pipeline because the cull pipeline takes and shades position attributes of vertices without performing rasterization and rendering pixels to a frame buffer. In at least one embodiment, the cull pipeline can use generated critical results to compute visibility information for all triangles, regardless of whether those triangles are culled or not. In at least one embodiment, the full pipeline, which can be referred to as a replay pipeline in this case, can consume the visibility information to skip culled triangles to only shade visible triangles that are ultimately passed to a rasterization stage.

[0270] In at least one embodiment, additional fixed function logic 2416 can also include general purpose processing acceleration logic, such as fixed function matrix multiplication logic, for implementing a reduced CUAD program.

[0271] In at least one embodiment, within each graphics sub-core 2401A-2401F includes a set of execution resources that can be used to perform graphics, media, and compute operations in response to requests by graphics pipeline, media pipeline, or shader programs. In at least one embodiment, graphics sub-cores 2401A-2401F include multiple arrays of execution units 2402A-2402F, 2404A-2404F, thread dispatch and inter-thread communication (TD / IC) logic 2403A-2403F, 3D (e.g., texture) samplers 2405A-2405F, media samplers 2406A-2406F, shader processors 2407A-2407F, and shared local memory (SLM) 2408A-2408F. Arrays of execution units 2402A-2402F, 2404A-2404F each include multiple execution units, which are GPGUs capable of performing floating point and integer / fixed point logic operations in support of graphics, media, and compute operations, including graphics, media, or compute shader programs. In at least one embodiment, TD / IC logic 2403A-2403F performs local thread dispatch and thread control operations for execution units within a sub-core and facilitate

[0272] Figure 25A parallel processing unit (“PPU”) 2500, in accordance with at least one embodiment, is shown. In at least one embodiment, PPU 2500 is configured with machine-readable code that, if executed by PPU 2500, causes PPU 2500 to perform some or all of the processes and techniques described throughout this document. In at least one embodiment, PPU 2500 is a multi-threaded processor implemented on one or more integrated circuit devices and utilizes multi-threading as a latency-hiding technique designed to process computer-readable instructions (also referred to as machine-readable instructions or simply instructions) that are executed in parallel across multiple threads. In at least one embodiment, a thread refers to an execution thread and is an instance of a set of instructions configured to be executed by PPU 2500. In at least one embodiment, PPU 2500 is a graphics processing unit (“GPU”) configured to implement a graphics rendering pipeline for processing three-dimensional (“3D”) graphics data in order to generate two-dimensional (“2D”) image data for display on a display device, such as an LCD device. In at least one embodiment, PPU 2500 is used to perform computations, such as linear algebraic operations and machine learning operations. Figure 25 The example parallel processor is shown for illustrative purposes only and should be construed as a non-limiting example of a processor architecture implemented in at least one embodiment.

[0273] In at least one embodiment, with respect to Figure 25 At least one component shown or described is used to implement techniques and / or functionality described Figures 1-8 In at least one embodiment, the at least one embodiment includes one or more secure devices having information that will be used to detect whether the secure device has been restarted indicating that the one or more secure devices have been tampered with.

[0274] In at least one embodiment, PPU(s) 2500 are configured to accelerate high- performance computing (“HPC”), datacenter, and machine learning applications. In at least one embodiment, PPU(s) 2500 are configured to accelerate CUDA programs. In at least one embodiment, PPU 2500 includes, without limitation, I / O unit 2506, front-end unit 2510, scheduler unit 2512, work distribution unit 2514, hub 2516, crossbar (“Xbar”) 2520, one or more general processing clusters (“GPCs”) 2518, and one or more partition units (“memory partition units”) 2522. In at least one embodiment, PPU(s) 2500 connect to a host processor or other PPU(s) 2500 by one or more high-speed GPU interconnects (“GPU interconnects”) 2508. In at least one embodiment, PPU(s) 2500 connect to host processor or other peripherals by system bus or interconnect 2502. In at least one embodiment, PPU(s) 2500 connect to a local memory comprising one or more memory devices (“memory”) 2504. In at least one embodiment, memory devices 2504 include, without limitation, one or more dynamic random access memory (“DRAM”) devices. In at least one embodiment, one or more DRAM devices are configured and / or configurable as high-bandwidth memory (“HBM”) subsystems with multiple DRAM dies stacked

[0275] In at least one embodiment, high-speed GPU interconnect 2508 can refer to a link-based parallel computer bus that systems use to scale and includes one or more PPUs 2500 in conjunction with one or more CPUs (“CPUs”) that support cache coherency between PPUs 2500 and CPUs and CPU mastering. In at least one embodiment, high-speed GPU interconnect 2508 transports data and / or commands through hub 2516 to other units of PPU(s) 2500, such as one or more copy engines, video encoders, video decoders, power management units, and / or other components not explicitly shown in Figure 25

[0276] In at least one embodiment, I / O unit 2506 is configured to facilitate communication between PPU(s) 2500 and a host processor (not shown), other PPU(s) 2500, and / or one or more Figure 25 ​The I / O units 2506 send and receive communications (e.g., commands, data) to and from the system bus 2502. In at least one embodiment, the I / O units 2506 communicate directly with the host processor(s) via the system bus 2502 or through one or more intermediate devices such as a memory hub. In at least one embodiment, the I / O units 2506 can communicate with one or more other processors, such as one or more PPUs 2500, via the system bus 2502. In at least one embodiment, the I / O units 2506 implement a PCIe interface for communications over a PCIe bus. In at least one embodiment, the I / O units 2506 implement interfaces for communicating with external devices.

[0277] In at least one embodiment, the I / O units 2506 decode packets received via the system bus 2502. In at least one embodiment, at least some packets represent commands configured to cause the PPU 2500 to perform various operations. In at least one embodiment, the I / O units 2506 send decoded commands to various other units of the PPU 2500 as designated by the commands. In at least one embodiment, commands are sent to the front-end unit 2510 and / or to the hub 2516 or other units of the PPU 2500 such as one or more copy engines, a video encoder, a video decoder, a power management unit, etc. Figure 25 In at least one embodiment, the I / O units 2506 are not explicitly shown in FIG. 25. In at least one embodiment, the I / O units 2506 are configured to route communications between various logical units of the PPU 2500.

[0278] In at least one embodiment, a program executed by the host processor encodes a command stream in a buffer that provides a workload to the PPU 2500 for processing. In at least one embodiment, the workload includes instructions and data to be processed by those instructions. In at least one embodiment, the buffer is a region in memory that is accessible (e.g., read / write) by both the host processor and the PPU 2500 - the host interface unit can be configured to access memory requests transmitted by the I / O units 2506 over the system bus 2502 to connect to a buffer in system memory. In at least one embodiment, the host processor writes the command stream to the buffer and then sends a pointer to the beginning of the command stream to the PPU 2500 so that the front-end unit 2510 receives the one or more command stream pointers and manages the one or more command streams, reading commands from the command stream and forwarding the commands to various units of the PPU 2500.

[0279] In at least one embodiment, front-end unit 2510 is coupled to a scheduler unit 2512 which configures various GPCs 2518 to process tasks defined by one or more command streams. In at least one embodiment, scheduler unit 2512 is configured to track state information related to various tasks managed by scheduler unit 2512, where state information can indicate which task is assigned to which GPC 2518, whether a task is active or inactive, a priority associated with a task, etc. In at least one embodiment, scheduler unit 2512 manages multiple tasks that are executed on one or more GPCs 2518.

[0280] In at least one embodiment, scheduler unit 2512 is coupled to a work distribution unit 2514, which is configured to dispatch tasks for execution on GPCs 2518. In at least one embodiment, work distribution unit 2514 tracks a number of scheduled tasks received from scheduler unit 2512 and work distribution unit 2514 manages a pending task pool and an active task pool for each GPC 2518. In at least one embodiment, the pending task pool includes a number of slots (e.g., 32 slots) that hold tasks assigned to be processed by a particular GPC 2518; the active task pool can include a number of slots (e.g., 4 slots) for tasks that are actively being processed by a GPC 2518, such that as one task completes execution on a GPC 2518, that task is evicted from the active task pool for GPC 2518 and one of the other tasks from the pending task pool is selected and scheduled for execution on the GPC 2518. In at least one embodiment, if an active task is idle, for example, while waiting for a data dependency to resolve, the active task is evicted from the GPC 2518 and returned to the pending task pool while another task is selected from the pending task pool and scheduled for execution on the GPC 2518.

[0281] In at least one embodiment, work distribution unit 2514 communicates with one or more GPCs 2518 via an Xbar 2520. In at least one embodiment, Xbar 2520 is an interconnect network that couples many units of PPU 2500 to other units of PPU 2500 and can be configured to couple work distribution unit 2514 to a particular GPC 2518. In at least one embodiment, other units of one or more PPUs 2500 can also be connected to Xbar 2520 via hub 2516.

[0282] In at least one embodiment, tasks are managed by a scheduler unit 2512 and dispatched to one of GPCs 2518 by a work distribution unit 2514. GPCs 2518 are configured to process tasks and generate results. In at least one embodiment, results can be consumed by other tasks within GPC 2518, routed to a different GPC 2518 via Xbar 2520, or stored in memory 2504. In at least one embodiment, results can be written to memory 2504 via a partition unit 2522, which implements a memory interface for reading and writing data to / from memory 2504. In at least one embodiment, results can be transmitted to another PPU 2500 or CPU via a high-speed GPU interconnect 2508. In at least one embodiment, PPU 2500 includes, without limitation, U partition units 2522 equal to the number of separate and distinct memory devices 2504 coupled to PPU 2500.

[0283] In at least one embodiment, a host processor executes a driver core that implements an application programming interface (API) that enables one or more applications executing on host processor to schedule operations to be performed on PPU 2500. In one embodiment, multiple compute applications are executed simultaneously by PPU 2500 and PPU 2500 provides isolation, quality-of-service (“QoS”), and independent address spaces for multiple compute applications. In at least one embodiment, an application generates instructions (e.g., in the form of API calls) that cause driver core to generate one or more tasks for execution by PPU 2500 and driver core outputs tasks to one or more streams that are processed by PPU 2500. In at least one embodiment, each task includes one or more related thread groups, which can be referred to as warps. In at least one embodiment, a warp includes multiple related threads (e.g., 32 threads) that can be executed in parallel. In at least one embodiment, a cooperative thread can refer to a plurality of threads including instructions for performing a task and exchanging data via shared memory.

[0284] Figure 26 A GPC 2600 according to at least one embodiment is shown. In at least one embodiment, GPC 2600 is a Figure 25GPC 2518. In at least one embodiment, each GPC 2600 includes, without limitation, a plurality of hardware units for handling graphics processing tasks and each GPC 2600 includes, without limitation, a pipeline manager 2602, a pre-raster operations unit (“PROP”) 2604, a raster engine 2608, a work distribution crossbar (“WDX”) 2616, a memory management unit (“MMU”) 2618, one or more data processing clusters (“DPCs”) 2606, and any suitable combination of such components.

[0285] In at least one embodiment, with respect to Figure 26 At least one component shown or described is used to implement a technique and / or functionality described Figures 1-8 In at least one embodiment, the at least one embodiment includes one or more secure devices having information that will be used to detect whether the secure device has been restarted indicating that the one or more secure devices have been tampered with.

[0286] In at least one embodiment, operation of GPC 2600 is controlled by pipeline manager 2602. In at least one embodiment, pipeline manager 2602 manages configuration of one or more DPCs 2606 to process tasks assigned to GPC 2600. In at least one embodiment, pipeline manager 2602 configures at least one of one or more DPCs 2606 to implement at least a portion of a graphics rendering pipeline. In at least one embodiment, DPC 2606 is configured to execute vertex shader programs on a programmable streaming multi-processor (“SM”) 2614. In at least one embodiment, pipeline manager 2602 is configured to route packets received from a work distribution unit to appropriate logical units within GPC 2600, and in at least one embodiment, some packets can be routed to fixed function hardware units in PROP 2604 and / or raster engine 2608 while other packets can be routed to DPCs 2606 for processing by a primitive engine 2612 or SM 2614. In at least one embodiment, pipeline manager 2602 configures at least one of DPCs 2606 to implement a neural network model and / or compute pipeline. In at least one embodiment, pipeline manager 2602 configures at least one of DPCs 2606 to execute at least a portion of a CUDA program.

[0287] In at least one embodiment, PROP unit 2604 is configured to route data generated by raster engine 2608 and DPCs 2606 to a raster operations (“ROP”) unit in a partition unit, such as described above in connection with Figure 25The memory partition unit 2522, among other things, is described in greater detail. In at least one embodiment, the PROP unit 2604 is configured to perform optimizations for color blending, organize pixel data, perform address translations, and the like. In at least one embodiment, the raster engine 2608 includes, without limitation, a number of fixed function hardware units configured to perform various raster operations, and in at least one embodiment, the raster engine 2608 includes, without limitation, a setup engine, a coarse raster engine, a cull engine, a clip engine, a fine raster engine, a tile aggregation engine, and any suitable combinations thereof. In at least one embodiment, the setup engine receives transformed vertices and generates plane equations associated with geometric primitives defined by the vertices; the plane equations are passed to the coarse raster engine to generate coverage information (e.g., x, y coverage masks for tiles) for the primitive; output from the coarse raster engine is passed to the cull engine, where fragments associated with primitives that fail a z-test are culled, and to the clip engine, where fragments that are outside a viewing frustum are clipped. In at least one embodiment, the culled and clipped fragments are passed to the fine raster engine to generate attributes of pixel fragments based on the plane equations generated by the setup engine. In at least one embodiment, output from the raster engine 2608 includes fragments that are to be processed by any suitable entity, such as by a fragment shader implemented within the DPC 2606.

[0288] In at least one embodiment, each DPC 2606 included in GPC 2600 includes, without limitation, an M-Pipe Controller (“MPC”) 2610; a primitive engine 2612; one or more SMs 2614; and any suitable combination thereof. In at least one embodiment, the MPC 2610 controls operation of the DPC 2606, routing received packets from the pipeline manager 2602 to appropriate units in the DPC 2606. In at least one embodiment, packets associated with vertices are routed to the primitive engine 2612, which is configured to fetch vertex attributes associated with the vertices from memory; in contrast, data packets associated with a shader program can be transmitted to the SM 2614.

[0289] In at least one embodiment, the SM 2614 includes, but is not limited to, a programmable streaming processor configured to process tasks represented by multiple threads. In at least one embodiment, the SM 2614 is multithreaded and configured to execute multiple threads (e.g., 32 threads) from a specific thread group concurrently, and implements a Single Instruction, Multiple Data (“SIMD”) architecture, wherein each thread in a group of threads (e.g., a thread bundle) is configured to process a different dataset based on the same instruction set. In at least one embodiment, all threads in the thread group execute the same instructions. In at least one embodiment, the SM 2614 implements a Single Instruction, Multiple Thread (“SIMT”) architecture, wherein each thread in a group of threads is configured to process a different dataset based on the same instruction set, but wherein individual threads in the thread group are allowed to diverge during execution. In at least one embodiment, a program counter, call stack, and execution state are maintained for each thread bundle, thereby achieving concurrency between the thread bundle and serial execution within the thread bundle when threads in the thread bundle diverge. In another embodiment, a program counter, call stack, and execution state are maintained for each individual thread, thereby ensuring equal concurrency among all threads within and between thread bundles. In at least one embodiment, an execution state is maintained for each individual thread, and threads executing the same instructions can be converged and executed in parallel to improve efficiency. The following is in conjunction with... Figure 27 At least one embodiment of SM 2614 is described in more detail.

[0290] In at least one embodiment, the MMU 2618 is integrated with the GPC 2600 and memory partitioning unit (e.g., Figure 25 The MMU 2618 provides an interface between partition units 2522 and provides virtual address to physical address translation, memory protection, and memory request arbitration. In at least one embodiment, the MMU 2618 provides one or more translation back buffers (“TLBs”) for performing virtual address to physical address translation in memory.

[0291] Figure 27 A streaming multiprocessor (“SM”) 2700 according to at least one embodiment is illustrated. In at least one embodiment, the SM 2700 is Figure 26SM 2614. In at least one embodiment, SM 2700 includes, without limitation, an instruction cache 2702; one or more scheduler units 2704; a register file 2708; one or more processing cores (“cores”) 2710; one or more special-function units (“SFUs”) 2712; one or more load / store units (“LSUs”) 2714; an interconnect network 2716; shared memory / level-one (“LI”) cache 2718; and any suitable combination thereof. In at least one embodiment, a work distribution unit dispatches tasks for execution on general processing clusters (“GPCs”) of parallel processing units (“PPUs”) and each task is assigned a specific data processing cluster (“DPC”) within a GPC and, if task is associated with a shader program, to one of SMs 2700. In at least one embodiment, scheduler units 2704 receive tasks from a work distribution unit and manage scheduling of instructions (e.g., thread blocks) to be executed by one or more of SMs 2700. In at least one embodiment, scheduler units 2704 schedule thread blocks for execution as thread warps of parallel threads, with each thread block allocated at least one thread warp. In at least one embodiment, each thread warp executes a thread. In at least one embodiment, scheduler units 2704 manage a plurality of different thread blocks, allocating thread warps to different thread blocks, and then dispatching instructions from different ones of the plurality of different thread blocks to various functional units (e.g., processing cores 2710, SFUs 2712, and LSUs 2714) during each clock cycle. In at least one embodiment, SM 2700 includes one or more clusters of thread blocks, where a cluster of thread blocks enables programmatic control over locality at a greater granularity than a single thread block of a streaming multi-processor (SM). In at least one embodiment, a cluster of thread blocks (also referred to as a “cluster”) enables multiple thread blocks running concurrently across streaming multi-processors to synchronize and cooperatively fetch, exchange, or otherwise use data.

[0292] In at least one embodiment, with respect to Figure 27 At least one component shown or described as being implemented within Figures 1-8 the techniques and / or functionality described. In at least one embodiment, the at least one embodiment includes one or more secure devices having information that will be used to detect whether the secure device has been restarted indicating that the one or more secure devices have been tampered with.

[0293] In at least one embodiment, a “cooperative group” can refer to a programming model for organizing groups of communication threads that allows developers to express the granularity at which threads are communicating, enabling richer, more efficient parallel decomposition. In at least one embodiment, a cooperative launch API supports synchronization between thread blocks to execute parallel algorithms. In at least one embodiment, an API of a conventional programming model provides a single, simple construct for synchronizing cooperative threads: a barrier across all threads of a thread block (e.g., a syncthreads() function). However, in at least one embodiment, a programmer can define thread groups at less than a thread block granularity and synchronize within defined groups to achieve higher performance, design flexibility, and software reuse in the form of collective group-wide function interfaces. In at least one embodiment, cooperative groups enable programmers to explicitly define thread groups at sub-block and multi-block granularity and perform collective operations, such as synchronizing threads in a cooperative group. In at least one embodiment, sub-block granularity is as small as a single thread. In at least one embodiment, a programming model supports clean composition across software boundaries, so that library and utility functions can safely synchronize in their local environment without having to make assumptions about convergence. In at least one embodiment, cooperative group primitives enable new patterns of cooperative parallelism, including but not limited to producer-consumer parallelism, opportunistic parallelism, and global synchronization across a grid of thread blocks.

[0294] In at least one embodiment, dispatch unit 2706 is configured to send instructions to one or more of the functional units, and a scheduler unit 2704 includes, without limitation, two dispatch units 2706 that enable two different instructions from the same thread to be dispatched in each clock cycle. In at least one embodiment, each scheduler unit 2704 includes a single dispatch unit 2706 or an additional dispatch unit 2706.

[0295] In at least one embodiment, each SM 2700 includes, without limitation, a register file 2708 that provides a set of registers for functional units of SM 2700. In at least one embodiment, register file 2708 is partitioned between functional units, with each functional unit being allocated a dedicated portion of register file 2708. In at least one embodiment, register file 2708 is partitioned between different warps executed by SM 2700 and register file 2708 provides temporary storage for operands of the operations calculated and performed by functional units. In at least one embodiment, each SM 2700 includes, without limitation, a plurality L of processing cores 2710. In at least one embodiment, SM 2700 includes, without limitation, a large number (e.g., 128 or more) of different processing cores 2710. In at least one embodiment, each processing core 2710 includes, without limitation, a full-pipe, single precision, double precision, and / or mixed precision processing unit including, without limitation, a floating point arithmetic logic unit and an integer arithmetic logic unit. In at least one embodiment, floating point arithmetic logic units implement IEEE 754-2008 standard for floating point arithmetic. In at least one embodiment, processing cores 2710 include, without limitation, 64 single-precision (32-bit) floating point cores, 64 integer cores, 32 double-precision (64-bit) floating point cores, and 8 tensor cores.

[0296] In at least one embodiment, tensor cores are configured to perform matrix operations. In at least one embodiment, one or more tensor cores are included in processing cores 2710. In at least one embodiment, tensor cores are configured to perform deep learning matrix arithmetic, such as convolution operations for neural network training and inferencing. In at least one embodiment, each tensor core operates on 4x4 matrices and performs matrix multiplication and accumulation operations D = A x B + C, where A, B, C, and D are 4x4 matrices.

[0297] In at least one embodiment, matrix multiplication inputs A and B are 16-bit floating point matrices, and accumulation matrices C and D are 16-bit floating point or 32-bit floating point matrices. In at least one embodiment, a tensor core performs 32-bit floating point accumulation operations on 16-bit floating point input data. In at least one embodiment, 16-bit floating point multiplication uses 64 operations and results in a full precision product, which is then accumulated with other intermediate products using 32-bit floating point addition for 4x4x4 matrix multiplication. In at least one embodiment, tensor cores are used to perform larger two-dimensional or higher dimensional matrix operations composed of these smaller elements. In at least one embodiment, an API such as a CUDA-C++ API exposes specialized matrix load, matrix multiply and accumulate, and matrix store operations to efficiently use tensor cores from a CUDA-C++ program. In at least one embodiment, at a CUDA level, a warp level interface assumes a 16x16 size matrix across all 32 warp threads.

[0298] In at least one embodiment, each SM 2700 includes, without limitation, M SFUs 2712 to perform special functions (e.g., certain math functions, exponentials, logarithms, etc.). In at least one embodiment, SFUs 2712 include, without limitation, tree traversal units configured to traverse a hierarchical tree data structure. In at least one embodiment, SFUs 2712 include, without limitation, texture units configured to perform texture mapping operations. In at least one embodiment, texture units are configured to load a texture map (e.g., a 2D array of texture pixels) from memory and sample the texture map to produce sampled texture values for use by a shader program executed by SM 2700. In at least one embodiment, texture maps are stored in shared memory / L1 cache 2718. In at least one embodiment, texture units use mip-maps (e.g., different levels of detail for a texture map) to perform texture operations such as filtering operations. In at least one embodiment, each SM 2700 includes, without limitation, two texture units.

[0299] In at least one embodiment, each SM 2700 includes, without limitation, N LSUs 2714 that implement load and store operations between shared memory / L1 cache 2718 and register file 2708. In at least one embodiment, each SM 2700 includes, without limitation, interconnect network 2716 that connects each of the functional units to register file 2708 and LSUs 2714 to register file 2708 and shared memory / L1 cache 2718. In at least one embodiment, interconnect network 2716 is a cross-bar switch that can be configured to connect any of the functional units to any of the registers in register file 2708 and connect LSUs 2714 to registers in register file 2708 and memory locations in shared memory / L1 cache 2718.

[0300] In at least one embodiment, shared memory / L1 cache 2718 is an array of on-chip memory that, in at least one embodiment, allows data storage and communication between SMs 2700 and the graphics processing engine and threads within SM 2700. In at least one embodiment, shared memory / L1 cache 2718 includes, without limitation, 128 KB of storage and is located on a path from SM 2700 to partition units. In at least one embodiment, shared memory / L1 cache 2718 is used for caching reads and writes, in at least one embodiment. In at least one embodiment, one or more of shared memory / L1 cache 2718, L2 cache, and memory are backing stores.

[0301] In at least one embodiment, combining data cache and shared memory functionality into a single memory block provides improved performance for both types of memory accesses. In at least one embodiment, capacity is used by programs that do not use shared memory or use it as a cache, for example if shared memory is configured to use half the capacity, then textures and load / store operations can use remaining capacity. According to at least one embodiment, integration within shared memory / L1 cache 2718 enables shared memory / L1 cache 2718 to be used as a high-throughput pipe for streaming data while also providing high bandwidth and low latency access to frequently reused data. In at least one embodiment, when configured for general purpose parallel computation, a simpler configuration can be used compared to graphics processing. In at least one embodiment, fixed function GPU is bypassed, creating a more straightforward programming model. In at least one embodiment, in a general purpose parallel computation configuration, work distribution unit allocates and distributes blocks of threads directly to DPCs. In at least one embodiment, threads in a block execute the same program, use unique thread IDs during computation to ensure that each thread generates unique results, use SM 2700 to execute the program and perform computations, use shared memory / L1 cache 2718 to communicate between threads, and use LSUs 2714 to read and write to global memory through shared memory / L1 cache 2718 and memory partition unit. In at least one embodiment, when configured for general purpose parallel computation, SM 2700 write scheduler unit 2704 can be used to launch new work on DPCs. In at least one embodiment, SM 2700 includes one or more distributed shared memories (or distributed shared memories) that enable direct SM to SM operations, such as cross-SM shared memory block loads, stores, and atomics.

[0302] In at least one embodiment, SM 2700 includes one or more asynchronous execution functions including a tensor memory accelerator (TMA) unit that can transfer data blocks between global memory and shared memory. In at least one embodiment, one or more processors use or access one or more TMAs to perform bi-directional copy operations, e.g., from global memory to shared memory, and vice versa. In at least one embodiment, SM 2700 includes one or more TMAs for asynchronously copying between thread blocks in a cluster. In at least one embodiment, SM 2700 includes one or more asynchronous transactional barriers for performing atomic data movement and synchronization. In at least one embodiment, SM 2700 includes a tensor core converter engine including software and one or more cores for accelerating converter model training and inference. In at least one embodiment, a converter executing one or more of the processor cores of one or more tensor core converter engines manages FP8 and 16-bit computations and dynamically selects between FP8 and 16-bit computations by re-casting and scaling between FP8 and 16-bit in each layer of one or more neural networks.

[0303] In at least one embodiment, PPU is included in or coupled to a desktop computer, a laptop computer, a tablet computer, servers, supercomputers, a smart-phone (e.g., a wireless, hand-held device), a PDA, a digital camera, a vehicle, a head mounted display, a hand-held electronic device, and the like. In at least one embodiment, PPU is implemented on a single semiconductor

[0304] In at least one embodiment, PPU can be included on a graphics card that includes one or more memory devices. In at least one embodiment, graphics card can be configured to interface with a PCIe slot on a motherboard of a desktop computer. In at least one embodiment, PPU can be an integrated GPU (“iGPU”) included in a chipset of a motherboard.

[0305] Software constructs for general purpose computing

[0306] The following figures illustrate, without limitation, example software constructs for implementing at least one embodiment.

[0307] Figure 28A software stack of a programming platform according to at least one embodiment is illustrated. In at least one embodiment, the programming platform is a platform for accelerating computational tasks by utilizing hardware on a computing system. In at least one embodiment, software developers can access the programming platform through libraries, compiler instructions, and / or extensions to programming languages. In at least one embodiment, the programming platform may be, but is not limited to, CUDA, Radeon Open Computing Platform (“ROCm”), OpenCL (OpenCL developed by Khronosgroup). TM ), SYCL or Intel One API.

[0308] In at least one embodiment, regarding Figure 28 At least one component shown or described is used to implement the combination. Figures 1-8 The described technologies and / or functions. In at least one embodiment, the at least one embodiment includes one or more security devices having information for detecting whether the security device has been restarted, thereby indicating that the one or more security devices have been tampered with.

[0309] In at least one embodiment, the software stack 2800 of the programming platform provides an execution environment for the application 2801. In at least one embodiment, the application 2801 may include any computer software capable of being launched on the software stack 2800. In at least one embodiment, the application 2801 may include, but is not limited to, artificial intelligence (“AI”) / machine learning (“ML”) applications, high-performance computing (“HPC”) applications, virtual desktop infrastructure (“VDI”) or data center workloads.

[0310] In at least one embodiment, application 2801 and software stack 2800 run on hardware 2807. In at least one embodiment, hardware 2807 may include one or more GPUs, CPUs, FPGAs, AI engines, and / or other types of computing devices supporting a programming platform. In at least one embodiment, such as using CUDA, software stack 2800 may be vendor-specific and compatible only with devices from a specific vendor. In at least one embodiment, such as using OpenCL, software stack 2800 may be used with devices from different vendors. In at least one embodiment, hardware 2807 includes a host connected to one or more devices that can be accessed via application programming interface (API) calls to perform computational tasks. In at least one embodiment, compared to the host within hardware 2807, which may include, but is not limited to, a CPU (but may also include computing devices) and its memory, devices within hardware 2807 may include, but are not limited to, GPUs, FPGAs, AI engines, or other computing devices (but may also include CPUs) and their memory.

[0311] In at least one embodiment, software stack 2800 of a programming platform includes, without limitation, a plurality of libraries 2803, a runtime 2805, and an in-device kernel driver 2806. In at least one embodiment, each of libraries 2803 can include data and programming code that can be used by computer programs and utilized during software development. In at least one embodiment, libraries 2803 can include, without limitation, pre-written code and subroutines, classes, values, type specifications, configuration data, documentation, help data, and / or message templates. In at least one embodiment, libraries 2803 include functions optimized for execution on one or more types of devices. In at least one embodiment, libraries 2803 can include, without limitation, functions for performing mathematical, deep learning, and / or other types of operations on a device. In at least one embodiment, libraries 2803 are associated with a corresponding API 2802, which can include one or more APIs that expose functions implemented in libraries 2803. In at least one embodiment, a processor (e.g., CPU, GPU) executes, calls, or otherwise uses one or more APIs to prioritize kernels. For example, a first kernel (e.g., a parent kernel) can launch a second kernel (e.g., a child kernel), and the second kernel can be used by a processor to launch an additional kernel (e.g., a grandchild kernel) independent of the first kernel. In at least one embodiment, a processor executes an API or calls an API to be executed from memory to support dynamic stream prioritization (e.g., updating priorities as streams are being used to perform operations). For example, when a processor executes the API, it allows a programmer to copy stream priorities from one stream to one or more other streams.

[0312] In at least one embodiment, software stack 2800 includes an API to support dynamic stream priority (e.g., priority is updated when a stream is being used to perform an operation), which allows a programmer to set a priority of a stream at any time after creation. In at least one embodiment, software stack 2800 includes an API to support dynamic stream priority (e.g., priority is updated when a stream is being used to perform an operation), which allows a programmer to get a current priority of a stream, where priority is one of a plurality of attributes of a stream. In at least one embodiment, software stack 2800 includes an API to support dynamic stream priority (e.g., priority is updated when a stream is being used to perform an operation), which allows a programmer to get a current priority of a stream as a single attribute. In at least one embodiment, software stack 2800 includes an API to support dynamic stream priority (e.g., priority is updated when a stream is being used to perform an operation), which allows a programmer to launch a kernel to perform an operation on a stream at a set priority (which can be different than a stream priority). In at least one embodiment, software stack 2800 includes an API to indicate whether an object (e.g., a thread synchronization object such as a barrier) tracks whether all data movement operations for a set of threads operating on a GPU are complete and have a specified state after a specified period of time, where a specified state can be a state that indicates data has been moved and is ready for use, and uses an expected parity value as an input to the API to specify.

[0313] In at least one embodiment, software stack 2800 includes one or more APIs to update a kernel. In at least one embodiment, a processor executes an API or calls an API to be executed from memory to update to an existing API to support a contextless kernel, which allows a programmer to add a kernel node to a graph without a graphics context, such that a graphics context can be dynamically associated with a kernel at runtime. In at least one embodiment, software stack 2800 includes one or more APIs that allow a programmer to get a kernel identifier and a graphics context from a kernel node as separate parameters, thereby obtaining parameters from a kernel and a contextless kernel. In at least one embodiment, software stack 2800 includes one or more APIs that use a parallel processor (such as one or more graphics processing units) to launch a task graph (e.g., a task graph) and execute one or more task graphs (e.g., including one or more programs).

[0314] In at least one embodiment, the software stack 2800 includes one or more APIs for associating one or more instructions with one or more memory sorting operations (e.g., fence or memory bar operations). In at least one embodiment, instructions are associated with one or more domains such that memory sorting operations are performed in association with one or more specific domains without interfering with instructions in other domains. The software stack 2800 includes one or more APIs for indicating that a thread has reached (e.g., at a thread synchronization barrier) or completed a work phase associated with an asynchronous data movement operation on the GPU. In at least one embodiment, the software stack 2800 includes one or more APIs for allowing a programmer to manually indicate an expected transaction count when a thread completes a work phase; this count is used to update an object that tracks whether all data movement operations of a set of threads have completed.

[0315] In at least one embodiment, application 2801 is written as source code, which is compiled into executable code, as follows: Figures 28-31 This will be discussed in more detail. In at least one embodiment, the executable code of application 2801 may run at least partially on an execution environment provided by software stack 2800. In at least one embodiment, during the execution of application 2801, code that needs to run on the device (compared to the host) may be obtained. In this case, in at least one embodiment, runtime 2805 may be invoked to load and start the necessary code on the device. In at least one embodiment, runtime 2805 may include any technically feasible runtime system capable of supporting the execution of application 2801.

[0316] In at least one embodiment, runtime 2805 is implemented as one or more runtime libraries associated with a corresponding API (shown as API 2804). In at least one embodiment, one or more such runtime libraries may include, but are not limited to, functions for memory management, execution control, device management, error handling, and / or synchronization, etc. In at least one embodiment, memory management functions may include, but are not limited to, functions for allocating, dealing with, and copying device memory, and for transferring data between host memory and device memory. In at least one embodiment, execution control functions may include, but are not limited to, functions for launching functions on the device (sometimes referred to as "kernels" when the function is a global function callable from the host), and functions for setting attribute values ​​in buffers maintained by the runtime library for a given function to be executed on the device.

[0317] In at least one embodiment, runtime libraries and corresponding APIs 2804 can be implemented in any technically feasible manner. In at least one embodiment, one (or any number) of APIs can expose a low-level set of functions for fine-grained control of a device, while another (or any number) of APIs can expose a higher-level set of functions. In at least one embodiment, high-level runtime APIs can be built on top of low-level APIs. In at least one embodiment, one or more runtime APIs can be language-specific APIs layered on top of language-independent runtime APIs.

[0318] In at least one embodiment, one or more processors disclosed in “processing systems” can execute, access, or otherwise utilize a software stack 2800. For example, APU 1500, CPU 1600, Figures 18A-18B In at least one embodiment, exemplary graphics processor, general purpose graphics processing unit (“GPGPU”) 1930, parallel processor 2000, processing cluster 2094, graphics multiprocessor 2034, graphics multiprocessor 2096, graphics processor 2100, processor 2200, processor 2300, parallel processing unit (“PPU”) 2500, GPC 2600, and / or streaming multiprocessor (“SM”) 2700 can execute, use, call or otherwise implement one or more APIs included in software stack 2800, for example, by accessing memory.

[0319] In at least one embodiment, a device kernel driver 2806 is configured to facilitate communication with underlying devices. In at least one embodiment, device kernel driver 2806 can provide low-level functions relied upon by APIs such as APIs 2804 and / or other software. In at least one embodiment, device kernel driver 2806 can be configured to compile intermediate representation (“IR”) code into binary code at runtime. In at least one embodiment, for CUDA, device kernel driver 2806 can compile non-hardware-specific parallel thread execution (“PTX”) IR code into binary code for a particular target device at runtime (caching compiled binary code), which is sometimes also referred to as “final” code. In at least one embodiment, doing so can allow final code to run on a target device that can not have existed when source code was initially compiled into PTX code. Alternatively, in at least one embodiment, device source code can be compiled into binary code offline without requiring device kernel driver 2806 to compile IR code at runtime.

[0320] Figure 29 FIG. 28 illustrates a software stack 2800, according to at least one embodiment Figure 28CUDA implementation of the software stack 2800. In at least one embodiment, the CUDA software stack 2900 on which the application 2901 can be launched includes CUDA libraries 2903, a CUDA runtime 2905, a CUDA driver 2907, and a device kernel driver 2908. In at least one embodiment, the CUDA software stack 2900 executes on hardware 2909, which can include a CUDA-enabled GPU, developed by NVIDIA Corporation of Santa Clara, California.

[0321] In at least one embodiment, with respect to Figure 29 At least one component shown or described as being implemented with software only is also implemented with one or more hardware components. In at least one embodiment, at least one component shown or described as being implemented with hardware only is also implemented with one or more software components. Figures 1-8 In at least one embodiment, the technology and / or functionality described herein is implemented with one or more software programs running on one or more computers or servers, which programs and computers are not shown or described herein. In at least one embodiment, at least one component shown or described as being implemented with software only is also implemented with one or more hardware components. In at least one embodiment, at least one component shown or described as being implemented with hardware only is also implemented with one or more software components.

[0322] In at least one embodiment, the application 2901, the CUDA runtime 2905, and the device kernel driver 2908 can perform similar functions to the application 2801, the runtime 2805, and the device kernel driver 2806, respectively, described above in connection with Figure 28CUDA Driver 2907, which includes a library (libcuda.so) that implements the CUDA Driver API 2906. In at least one embodiment, similar to the CUDA Runtime API 2904 implemented by the CUDA Runtime Library (cudart), the CUDA Driver API 2906 can expose, without limitation, functions for memory management, execution control, device management, error handling, synchronization, and / or graphics interoperability, among others. In at least one embodiment, the CUDA Driver API 2906 differs from the CUDA Runtime API 2904 in that the CUDA Runtime API 2904 simplifies device code management by providing implicit initialization, context (similar to a process) management, and module (similar to a dynamically loaded library) management. In contrast to the high-level CUDA Runtime API 2904, in at least one embodiment, the CUDA Driver API 2906 is a low-level API that provides more granular control over a device, particularly with respect to context and module loading. In at least one embodiment, the CUDA Driver API 2906 can expose functions for context management that are not exposed by the CUDA Runtime API 2904. In at least one embodiment, the CUDA Driver API 2906 is also language agnostic and supports, for example, OpenCL in addition to the CUDA Runtime API 2904. Further, in at least one embodiment, development libraries including the CUDA Runtime 2905 can be considered separate from driver components, including the user-mode CUDA Driver 2907 and the kernel-mode device driver 2908 (sometimes also referred to as a “display” driver).

[0323] In at least one embodiment, CUDA Libraries 2903 can include, without limitation, mathematical libraries, deep learning libraries, parallel algorithm libraries, and / or signal / image / video processing libraries that can be utilized by parallel computing applications, such as application 2901. In at least one embodiment, CUDA Libraries 2903 can include mathematical libraries such as a cuBLAS library, which is an implementation of basic linear algebra subprograms (“BLAS”) for performing linear algebra operations; a cuFFT library for computing fast Fourier transforms (“FFTs”), and a cuRAND library for generating random numbers, among others. In at least one embodiment, CUDA Libraries 2903 can include deep learning libraries such as a cuDNN library for primitives of deep neural networks and a TensorRT platform for high-performance deep learning inference, among others.

[0324] Figure 30 FIG. 29 illustrates a computing system 2900 in accordance with at least one embodiment, in which a CUDA runtime library 2905 and CUDA driver library 2906 are used to implement CUDA runtime API 2904 and CUDA driver API 2906, respectively, in accordance with at least one embodiment. Figure 28ROCm implementation of software stack 2800. In at least one embodiment, ROCm software stack 3000 on which application 3001 can be launched includes a language runtime 3003, a system runtime 3005, a thunk 3007, and a ROCm kernel driver 3008. In at least one embodiment, ROCm software stack 3000 executes on hardware 3009, which can include a GPU that supports ROCm, which is developed by AMD Corporation of Santa Clara, California.

[0325] In at least one embodiment, with respect to Figure 30 At least one component shown or described is used to implement the techniques and / or functionality described herein. In at least one embodiment, the at least one embodiment includes one or more secure devices having information that will be used to detect whether the secure device has been restarted, indicating that the one or more secure devices have been tampered with. Figures 1-8 At least one component shown or described is used to implement the techniques and / or functionality described herein. In at least one embodiment, the at least one embodiment includes one or more secure devices having information that will be used to detect whether the secure device has been restarted, indicating that the one or more secure devices have been tampered with.

[0326] In at least one embodiment, application 3001 can perform similar functionality to application 2801 discussed above in connection with Figure 28 In at least one embodiment, language runtime 3003 and system runtime 3005 can perform similar functionality to runtime 2805 discussed above in connection with Figure 28 In at least one embodiment, language runtime 3003 and system runtime 3005 can perform similar functionality to runtime 2805 discussed above in connection with Figure 29 In at least one embodiment, language runtime 3003 and system runtime 3005 can perform similar functionality to runtime 2805 discussed above in connection with

[0327] In at least one embodiment, the thunk (ROCt) 3007 is an interface 3006 that can be used to interact with the underlying ROCm driver 3008. In at least one embodiment, the ROCm driver 3008 is a ROCk driver, which is a combination of an AMD GPU driver and an HSA core driver (amdkfd). In at least one embodiment, the AMD GPU driver is a device core driver for GPUs developed by AMD, which performs the above-described combination. Figure 28 The device kernel driver 2806 discussed has similar functionality. In at least one embodiment, the HSA kernel driver is a driver that allows different types of processors to share system resources more efficiently via hardware features.

[0328] In at least one embodiment, various libraries (not shown) may be included in the ROCm software stack 3000 above the language runtime 3003, and provide integration with the above. Figure 29 The discussed CUDA library 2903 has similar functionality. In at least one embodiment, various libraries may include, but are not limited to, mathematical, deep learning, and / or other libraries, such as the hipBLAS library which implements functions similar to CUDA cuBLAS, the rocFFT library which is similar to CUDA cuFFT for computing FFT, etc.

[0329] Figure 31 The illustration shows an embodiment according to at least one of the embodiments. Figure 28 The software stack 2800 is an OpenCL implementation. In at least one embodiment, the OpenCL software stack 3100 on which the application 3101 can be launched includes an OpenCL framework 3110, an OpenCL runtime 3106, and a driver 3107. In at least one embodiment, the OpenCL software stack 3100 executes on non-vendor-specific hardware 2909. In at least one embodiment, because devices developed by different vendors support OpenCL, specific OpenCL drivers may be required for interoperability with hardware from such vendors.

[0330] In at least one embodiment, regarding Figure 31 At least one component shown or described is used to implement the combination. Figures 1-8 The described technologies and / or functions. In at least one embodiment, the at least one embodiment includes one or more security devices having information for detecting whether the security device has been restarted, thereby indicating that the one or more security devices have been tampered with.

[0331] In at least one embodiment, application 3101, OpenCL runtime 3106, device kernel driver 3107, and hardware 3108 can perform similar functions to application 2801, runtime 2805, device kernel driver 2806, and hardware 2807, respectively, described above in conjunction with FIG. 28. In at least one embodiment, application 3101 also includes OpenCL kernels 3102 having code to be executed on a device. Figure 28

[0332] In at least one embodiment, OpenCL defines a “platform” that allows a host to control devices connected to that host. In at least one embodiment, OpenCL framework provides a platform layer API and a runtime API, shown as platform API 3103 and runtime API 3105. In at least one embodiment, runtime API 3105 uses a context to manage execution of kernels on a device. In at least one embodiment, each identified device can be associated with a respective context, which runtime API 3105 can use to manage a command queue, program objects and kernel objects, shared memory objects, etc. for that device. In at least one embodiment, platform API 3103 exposes functions that allow a device context to be used to select and initialize a device, submit work to a device via a command queue, and enable data transfers to and from a device, among other things. Additionally, in at least one embodiment, OpenCL framework provides various built-in functions (not shown), including mathematical functions, relational functions, and image processing functions, among others.

[0333] In at least one embodiment, compiler 3104 is also included in OpenCL framework 3110. In at least one embodiment, source code can be compiled offline before an application is executed or online during execution of an application. In contrast to CUDA and ROCm, OpenCL applications in at least one embodiment can be compiled online by compiler 3104, which is included to represent any number of compilers that can be used to compile source code and / or IR code (e.g., Standard Portable Intermediate Representation (“SPIR-V”) code) into binary code. Alternatively, in at least one embodiment, OpenCL applications can be compiled offline before such applications are executed.

[0334] Figure 32 ​Software supported by a programming platform is shown. In at least one embodiment, programming platform 3204 is configured to support various programming models 3203, middleware and / or libraries 3202, and frameworks 3201 that an application 3200 can rely on. In at least one embodiment, application 3200 can be an AI / ML application implemented using, for example, a deep learning framework (e.g., MXNet, PyTorch, or TensorFlow) that can rely on libraries such as cuDNN, NVIDIA Collective Communications Library (“NCCL”), and / or NVIDIA Developer Data Loading Library (“DALI”) CUDA libraries to provide accelerated computing on underlying hardware.

[0335] In at least one embodiment, with respect to Figure 32 At least one component shown or described is used to implement a technique and / or functionality described. Figures 1-8 In at least one embodiment, the at least one embodiment includes one or more secure devices having information that will be used to detect whether the secure device has been restarted indicating that the one or more secure devices have been tampered with.

[0336] In at least one embodiment, programming platform 3204 can be one of the CUDA, ROCm, or OpenCL platforms described above in connection with Figure 29 , Figure 30 and Figure 31 respectively. In at least one embodiment, programming platform 3204 supports multiple programming models 3203, which are abstractions of underlying computing systems that allow expression of algorithms and data structures. In at least one embodiment, programming models 3203 can expose features of underlying hardware in order to improve performance. In at least one embodiment, programming models 3203 can include, but are not limited to, CUDA, HIP, OpenCL, C++ Accelerated Massive Parallelism (“C++ AMP”), Open Multi-Processing (“OpenMP”), Open Accelerators (“OpenACC”), and / or Vulcan Compute.

[0337] In at least one embodiment, libraries and / or middleware 3202 provide implementations of abstractions of programming models 3204. In at least one embodiment, such libraries include data and programming code that can be used by computer programs and utilized during software development. In at least one embodiment, such middleware includes software that provides services to applications in addition to those that can be obtained from a programming platform 3204. In at least one embodiment, libraries and / or middleware 3202 can include, but are not limited to, cuBLAS, cuFFT, cuRAND, and other CUDA libraries, or rocBLAS, rocFFT, rocRAND, and other ROCm libraries. Additionally, in at least one embodiment, libraries and / or middleware 3202 can include NCCL and ROCm Communication Collectives Library (“RCCL”) libraries, which provide communication routines for GPUs, MIOpen libraries for deep learning acceleration, and Eigen libraries for linear algebra, matrix and vector operations, geometric transformations, numerical solvers, and related algorithms.

[0338] In at least one embodiment, application frameworks 3201 rely on libraries and / or middleware 3202. In at least one embodiment, each application framework 3201 is a software framework used to implement a standard structure for application software. Returning to AI / ML example discussed above, in at least one embodiment, a framework such as Caffe, Caffe2, TensorFlow, Keras, PyTorch, or MxNet deep learning framework can be used to implement an AI / ML application.

[0339] Figure 33 FIG. 33 illustrates compiled code to execute on one of Figures 28-31 programming platforms, in accordance with at least one embodiment. In at least one embodiment, compiler 3301 receives source code 3300, which includes both host code as well as device code. In at least one embodiment, compiler 3301 is configured to convert source code 3300 into host executable code 3302 for execution on a host and device executable code 3303 for execution on a device. In at least one embodiment, source code 3300 can be compiled offline, prior to execution of an application, or online during execution of an application. In at least one embodiment, compiler 3301 includes or has access to one or more libraries for identifying sequences of API calls to perform a single fused API, where the single fused API is a combined API of two or more APIs.

[0340] In at least one embodiment, with respect to Figure 33 At least one component shown or described is used to implement a combination of Figures 1-8technology and / or functionality described. In at least one embodiment, the at least one embodiment includes one or more security devices having information that will be used to detect whether the security device has been re-started indicating that the one or more security devices have been tampered with.

[0341] In at least one embodiment, source code 3300 can include code in any programming language supported by compiler 3301, such as C++, C, Fortran, etc. In at least one embodiment, source code 3300 can include a single-source file having a mix of host code and device code, with locations of device code indicated therein. In at least one embodiment, the single-source file can be a.cu file including CUDA code or a.hip.cpp file including HIP code. Alternatively, in at least one embodiment, source code 3300 can include multiple source code files, rather than a single-source file, with host code and device code separated.

[0342] In at least one embodiment, compiler 3301 is configured to compile source code 3300 into host executable code 3302 for execution on a host and device executable code 3303 for execution on a device. In at least one embodiment, compiler 3301 performs operations including parsing source code 3300 into an abstract syntax tree (AST), performing optimizations, and generating executable code. In at least one embodiment in which source code 3300 includes a single-source file, compiler 3301 can separate device code from host code in such single-source file, compile device code and host code separately into device executable code 3303 and host executable code 3302, and link device executable code 3303 and host executable code 3302 together in a single file, as discussed below with respect to FIG. 4. Figure 34 discussed in more detail.

[0343] In at least one embodiment, host executable code 3302 and device executable code 3303 can be in any suitable format, such as binary code and / or IR code. In the case of CUDA, in at least one embodiment, host executable code 3302 can include native object code, while device executable code 3303 can include PTX intermediate representation code. In the case of ROCm, in at least one embodiment, both host executable code 3302 and device executable code 3303 can include object binary code.

[0344] Figure 34 is compiled code to be executed on a host in accordance with at least one embodiment. Figures 28-31more detail on one of the programming platforms described above. In at least one embodiment, a compiler 3401 is configured to receive source code 3400, compile source code 3400, and output executable 3410. In at least one embodiment, source code 3400 is a single source file, such as a.cu file, a.hip.cpp file, or other format of file, that includes both host code and device code. In at least one embodiment, compiler 3401 can be, without limitation, an NVIDIA CUDA compiler (“NVCC”) for compiling CUDA code in a.cu file, or an HCC compiler for compiling HIP code in a.hip.cpp file.

[0345] In at least one embodiment, with respect to Figure 34 At least one of the components shown or described is used in implementing the techniques and / or functionality described Figures 1-8 described. In at least one embodiment, the at least one embodiment includes one or more secure devices having information that will be used to detect whether the secure device has been re-started, indicating that the one or more secure devices have been tampered with.

[0346] In at least one embodiment, compiler 3401 includes a compiler front end 3402, a host compiler 3405, a device compiler 3406, and a linker 3409. In at least one embodiment, compiler front end 3402 is configured to separate device code 3404 from host code 3403 in source code 3400. In at least one embodiment, device code 3404 is compiled by device compiler 3406 into device executable code 3408, which can include binary code or IR code, as described. In at least one embodiment, host code 3403 is separately compiled by host compiler 3405 into host executable code 3407. In at least one embodiment, for NVCC, host compiler 3405 can be, without limitation, a general C / C++ compiler that outputs native object code, while device compiler 3406 can be, without limitation, a low-level virtual machine (“LLVM”) based compiler that forks the LLVM compiler infrastructure and outputs PTX code or binary code. In at least one embodiment, for HCC, both host compiler 3405 and device compiler 3406 can be, without limitation, LLVM based compilers that output target binary code.

[0347] In at least one embodiment, after source code 3400 is compiled into host executable code 3407 and device executable code 3408, linker 3409 links host and device executable code 3407 and 3408 together in executable file 3410. In at least one embodiment, host and PTX native object code or device binary code can be linked together in an Executable and Linkable Format (“ELF”) file, which is a container format for storing object code.

[0348] Figure 35 Compiling source code prior to translation is shown, in accordance with at least one embodiment. In at least one embodiment, source code 3500 is passed through translation tool 3501, which translates source code 3500 into translated source code 3502. In at least one embodiment, compiler 3503 is used to compile translated source code 3502 into host executable code 3504 and device executable code 3405, a process similar to that of compiler 3301 compiling source code 3300 into host executable code 3302 and device executable code 3303, as described above in connection with FIG. 3. Figure 33

[0349] In at least one embodiment, at least one component shown or described is used to implement techniques and / or functionality described in connection with Figure 35 In at least one embodiment, at least one component shown or described is used to implement techniques and / or functionality described in connection with Figures 1-8 In at least one embodiment, the at least one embodiment includes one or more secure devices having information that will be used to detect whether the secure device has been restarted, indicating that the one or more secure devices have been tampered with.

[0350] In at least one embodiment, translation performed by translation tool 3501 is used to port source code 3500 to execute in a different environment than originally intended. In at least one embodiment, translation tool 3501 can include, without limitation, a HIP translator used to “hipify” CUDA code for a CUDA platform into HIP code that can be compiled and executed on a ROCm platform. In at least one embodiment, translation of source code 3500 can include parsing source code 3500 and translating calls to APIs provided by one programming model (e.g., CUDA) to corresponding calls to APIs provided by another programming model (e.g., HIP), as described below in connection with FIG. 4. Figure 36A Figure 37 ​​discussed in more detail. Returning to the example of porting CUDA code, in at least one embodiment, calls to CUDA runtime APIs, CUDA driver APIs, and / or CUDA libraries can be converted to corresponding HIP API calls. In at least one embodiment, the automatic conversion performed by the conversion tool 3501 can sometimes be incomplete, requiring additional human effort to fully port the source code 3500.

[0351] Configuring GPUs for general purpose computing

[0352] The following figures set forth, without limitation, exemplary architectures for compiling and executing compute source code, in accordance with at least one embodiment.

[0353] Figure 36A A system 3600 configured to compile and execute CUDA source code 3610 using different types of processing units is shown, in accordance with at least one embodiment. In at least one embodiment, system 3600 includes, without limitation, CUDA source code 3610, a CUDA compiler 3650, host executable code 3670(1), host executable code 3670(2), CUDA device executable code 3684, a CPU 3690, a CUDA-enabled GPU 3694, a GPU 3692, a CUDA to HIP conversion tool 3620, HIP source code 3630, a HIP compiler driver 3640, an HCC 3660, and HCC device executable code 3682.

[0354] In at least one embodiment, with respect to Figure 36A At least one component shown or described is used to implement a technique and / or functionality described Figures 1-8 described. In at least one embodiment, the at least one embodiment includes one or more secure devices having information that will be used to detect whether the secure device has been restarted, indicating that the one or more secure devices have been tampered with.

[0355] In at least one embodiment, CUDA source code 3610 is a set of human-readable code in a CUDA programming language. In at least one embodiment, CUDA code is human-readable code in a CUDA programming language. In at least one embodiment, a CUDA programming language is an extension of a C++ programming language that includes, without limitation, mechanisms to define device code and to distinguish between device code and host code. In at least one embodiment, device code is source code that, upon compilation, is executable in parallel on a device. In at least one embodiment, a device can be a processor optimized for parallel instruction processing, such as CUDA-enabled GPU 3690, GPU 36192, or another GPGPU, etc. In at least one embodiment, host code is source code that, upon compilation, is executable on a host. In at least one embodiment, a host is a processor optimized for sequential instruction processing, such as CPU 3690.

[0356] In at least one embodiment, CUDA source code 3610 includes, without limitation, any number (including zero) of global functions 3612, any number (including zero) of device functions 3614, any number (including zero) of host functions 3616, and any number (including zero) of host / device functions 3618. In at least one embodiment, global functions 3612, device functions 3614, host functions 3616, and host / device functions 3618 can be intermingled in CUDA source code 3610. In at least one embodiment, each global function 3612 is executable on a device and is invocable from a host. Thus, in at least one embodiment, one or more of global functions 3612 can serve as an entry point to a device. In at least one embodiment, each global function 3612 is a kernel. In at least one embodiment, and in a technique known as dynamic parallelism, one or more global functions 3612 define a kernel that is executable on a device and is invocable from such a device. In at least one embodiment, a kernel is executed N times in parallel by N different threads on a device during execution (where N is any positive integer).

[0357] In at least one embodiment, each device function 3614 is executable on a device and is invocable only from such a device. In at least one embodiment, each host function 3616 is executable on a host and is invocable only from such a host. In at least one embodiment, each host / device function 3616 defines both a host version of a function that is executable on a host and is invocable only from such a host, and a device version of a function that is executable on a device and is invocable only from such a device.

[0358] In at least one embodiment, CUDA source code 3610 can also include, without limitation, any number of calls to any number of functions defined by CUDA runtime API 3602. In at least one embodiment, CUDA runtime API 3602 can include, without limitation, any number of functions executing on a host that are used to allocate and de-allocate device memory, transfer data between host memory and device memory, manage systems with multiple devices, etc. In at least one embodiment, CUDA source code 3610 can also include any number of calls to any number of functions specified in any number of other CUDA APIs. In at least one embodiment, a CUDA API can be any API designed to be used by CUDA code. In at least one embodiment, CUDA APIs include, without limitation, CUDA runtime API 3602, a CUDA driver API, APIs for any number of CUDA libraries, etc. In at least one embodiment and relative to CUDA runtime API 3602, a CUDA driver API is a lower-level API but can provide more fine-grained control over a device. In at least one embodiment, examples of CUDA libraries include, without limitation, cuBLAS, cuFFT, cuRAND, cuDNN, etc.

[0359] In at least one embodiment, CUDA compiler 3650 compiles input CUDA code (e.g., CUDA source code 3610) to generate host executable code 3670(1) and CUDA device executable code 3684. In at least one embodiment, CUDA compiler 3650 is NVCC. In at least one embodiment, host executable code 3670(1) is a compiled version of host code included in input source code that is executable on CPU 3690. In at least one embodiment, CPU 3690 can be any processor optimized for sequential instruction processing.

[0360] In at least one embodiment, CUDA device executable code 3684 is a compiled version of device code included in input source code executable on a CUDA-enabled GPU 3694. In at least one embodiment, CUDA device executable code 3684 includes, without limitation, binary code. In at least one embodiment, CUDA device executable code 3684 includes, without limitation, IR code, such as PTX code, which is further compiled into binary code for a particular target device (e.g., CUDA-enabled GPU 3694) at runtime by a device driver. In at least one embodiment, CUDA-enabled GPU 3694 can be any processor optimized for parallel instruction processing and supporting CUDA. In at least one embodiment, CUDA-enabled GPU 3694 is developed by NVIDIA Corporation of Santa Clara, CA.

[0361] In at least one embodiment, CUDA to HIP translation tool 3620 is configured to translate CUDA source code 3610 into functionally similar HIP source code 3630. In at least one embodiment, HIP source code 3630 is a collection of human-readable code in the HIP programming language. In at least one embodiment, HIP code is human-readable code in the HIP programming language. In at least one embodiment, the HIP programming language is an extension of the C++ programming language that includes, without limitation, functionally similar versions of CUDA mechanisms for defining device code and distinguishing device code from host code. In at least one embodiment, the HIP programming language can include a subset of functionality of the CUDA programming language. In at least one embodiment, for example, the HIP programming language includes, without limitation, mechanisms to define global functions 3612, but such a HIP programming language can lack support for dynamic parallelism, and thus global functions 3612 defined in HIP code are only callable from a host.

[0362] In at least one embodiment, HIP source code 3630 includes, without limitation, any number (including zero) of global functions 3612, any number (including zero) of device functions 3614, any number (including zero) of host functions 3616, and any number (including zero) of host / device functions 3618. In at least one embodiment, HIP source code 3630 can also include any number of calls to any number of functions specified in a HIP runtime API 3632. In one embodiment, HIP runtime API 3632 includes, without limitation, functionally similar versions of a subset of functions included in CUDA runtime API 3602. In at least one embodiment, HIP source code 3630 can also include any number of calls to any number of functions specified in any number of other HIP APIs. In at least one embodiment, a HIP API can be any API designed for use with HIP code and / or ROCm. In at least one embodiment, HIP APIs include, without limitation, HIP runtime API 3632, a HIP driver API, APIs for any number of HIP libraries, APIs for any number of ROCm libraries, and the like.

[0363] In at least one embodiment, CUDA to HIP translation tool 3620 translates each kernel call in CUDA code from CUDA syntax to HIP syntax, and translates any number of other CUDA calls in CUDA code to any number of other functionally similar HIP calls. In at least one embodiment, a CUDA call is a call to a function specified in a CUDA API, and a HIP call is a call to a function specified in a HIP API. In at least one embodiment, CUDA to HIP translation tool 3620 translates any number of calls to functions specified in CUDA runtime API 3602 to any number of calls to functions specified in HIP runtime API 3632.

[0364] In at least one embodiment, CUDA to HIP translation tool 3620 is a tool known as hipify-perl, which performs a text-based translation process. In at least one embodiment, CUDA to HIP translation tool 3620 is a tool known as hipify-clang, which performs a more complex and robust translation process relative to hipify-perl, involving parsing of CUDA code using clang (a compiler frontend), followed by translation of resulting symbols. In at least one embodiment, in addition to modifications performed by CUDA to HIP translation tool 3620, proper translation of CUDA code to HIP code can also require modifications (e.g., manual edits).

[0365] In at least one embodiment, HIP compiler driver 3640 is a front end that determines target device 3646 and then configures a compiler compatible with target device 3646 to compile HIP source code 3630. In at least one embodiment, target device 3646 is a processor optimized for parallel instruction processing. In at least one embodiment, HIP compiler driver 3640 can determine target device 3646 in any technically feasible manner.

[0366] In at least one embodiment, if target device 3646 is compatible with CUDA (e.g., a CUDA-enabled GPU 3694), then HIP compiler driver 3640 generates HIP / NVCC compilation commands 3642. In at least one embodiment and in conjunction with Figure 36B In more detail, HIP / NVCC compilation commands 3642 configure CUDA compiler 3650 to compile HIP source code 3630 using, without limitation, a HIP-to-CUDA translation header and a CUDA runtime library. In at least one embodiment and in response to HIP / NVCC compilation commands 3642, CUDA compiler 3650 generates host executable code 3670(1) and CUDA device executable code 3684.

[0367] In at least one embodiment, if target device 3646 is not compatible with CUDA, then HIP compiler driver 3640 generates HIP / HCC compilation commands 3644. In at least one embodiment and in conjunction with Figure 36C In more detail, HIP / HCC compilation commands 3644 configure HCC 3660 to compile HIP source code 3630 using a HCC header and a HIP / HCC runtime library. In at least one embodiment and in response to HIP / HCC compilation commands 3644, HCC 3660 generates host executable code 3670(2) and HCC device executable code 3682. In at least one embodiment, HCC device executable code 3682 is a compiled version of device code contained in HIP source code 3630 that is executable on GPU 3692. In at least one embodiment, GPU 3692 can be any processor optimized for parallel instruction processing that is not compatible with CUDA and is compatible with HCC. In at least one embodiment, GPU 3692 is developed by AMD Corporation of Santa Clara, California. In at least one embodiment, GPU 3692 is a GPU 3692 that is not CUDA-enabled.

[0368] For illustrative purposes only, in Figure 36ACUDA source code 3610 to execute on CPU 3690 and a different device in at least one embodiment. In at least one embodiment, a direct CUDA flow compiles CUDA source code 3610 to execute on CPU 3690 and a CUDA-enabled GPU 3694 without converting CUDA source code 3610 to HIP source code 3630. In at least one embodiment, an indirect CUDA flow converts CUDA source code 3610 to HIP source code 3630 and then compiles HIP source code 3630 to execute on CPU 3690 and a CUDA-enabled GPU 3694. In at least one embodiment, a CUDA / HCC flow converts CUDA source code 3610 to HIP source code 3630 and then compiles HIP source code 3630 to execute on CPU 3690 and a GPU 3692.

[0369] A direct CUDA flow that can be implemented in at least one embodiment can be depicted by the dashed line and series of bubble annotations Al-A3. In at least one embodiment, and as shown by bubble annotation Al, CUDA compiler 3650 receives CUDA source code 3610 and a CUDA compile command 3648 that configures CUDA compiler 3650 to compile CUDA source code 3610. In at least one embodiment, CUDA source code 3610 used in the direct CUDA flow is written in the CUDA programming language that is based on a programming language other than C++ (e.g., C, Fortran, Python, Java, etc.). In at least one embodiment, and in response to CUDA compile command 3648, CUDA compiler 3650 generates host executable code 3670(1) and CUDA device executable code 3684 (denoted by bubble annotation A2). In at least one embodiment and as shown by bubble annotation A3, host executable code 3670(1) and CUDA device executable code 3684 can be executed on CPU 3690 and CUDA-enabled GPU 3694, respectively. In at least one embodiment, CUDA device executable code 3684 includes, without limitation, binary code. In at least one embodiment, CUDA device executable code 3684 includes, without limitation, PTX code and is further compiled into binary code for a particular target device at runtime.

[0370] The indirect CUDA flow that can be implemented in at least one embodiment can be described by the dashed line and series of bubble annotations Bl- B6. In at least one embodiment and as shown by bubble annotation Bl, the CUDA to HIP translation tool 3620 receives CUDA source code 3610. In at least one embodiment and as shown by bubble annotation B2, the CUDA to HIP translation tool 3620 translates the CUDA source code 3610 to HIP source code 3630. In at least one embodiment and as shown by bubble annotation B3, the HIP compiler driver 3640 receives the HIP source code 3630 and determines whether the target device 3646 has CUDA enabled.

[0371] In at least one embodiment and as shown by bubble annotation B4, the HIP compiler driver 3640 generates HIP / NVCC compilation commands 3642 and sends both the HIP / NVCC compilation commands 3642 and the HIP source code 3630 to the CUDA compiler 3650. In at least one embodiment and as described in greater detail below in connection with FIG. 36B, the HIP / NVCC compilation commands 3642 configure the CUDA compiler 3650 to compile the HIP source code 3630 using, without limitation, a HIP to CUDA translation header and a CUDA runtime library. Figure 36B In at least one embodiment and in response to the HIP / NVCC compilation commands 3642, the CUDA compiler 3650 generates host executable code 3670(1) and CUDA device executable code 3684 (represented by bubble annotation B5). In at least one embodiment and as shown by bubble annotation B6, the host executable code 3670(1) and the CUDA device executable code 3684 can be executed on a CPU 3690 and a CUDA-enabled GPU 3694, respectively. In at least one embodiment, the CUDA device executable code 3684 includes, without limitation, binary code. In at least one embodiment, the CUDA device executable code 3684 includes, without limitation, PTX code and is further compiled into binary code for a particular target device at runtime.

[0372] The CUDA / HCC flow that can be implemented in at least one embodiment can be described by the solid line and series of bubble annotations Cl- C6. In at least one embodiment and as shown by bubble annotation Cl, the CUDA to HIP translation tool 3620 receives CUDA source code 3610. In at least one embodiment and as shown by bubble annotation C2, the CUDA to HIP translation tool 3620 translates the CUDA source code 3610 to HIP source code 3630. In at least one embodiment and as shown by bubble annotation C3, the HIP compiler driver 3640 receives the HIP source code 3630 and determines that the target device 3646 does not have CUDA enabled. In at least one embodiment and as shown by bubble annotation C4, the HIP compiler driver 3640 generates HCC compilation commands 3644 and sends both the HCC compilation commands 3644 and the HIP source code 3630 to the HCC compiler 3660. In at least one embodiment and as described in greater detail below in connection with FIG. 36C, the HCC compilation commands 3644 configure the HCC compiler 3660 to compile the HIP source code 3630 using, without limitation, a HIP to HCC translation header and a HCC runtime library.

[0373] In at least one embodiment, HIP compiler driver 3640 generates HIP / HCC compilation commands 3664 and sends both HIP / HCC compilation commands 3664 and HIP source code 3630 to HCC 3660 (represented with bubble annotation C4). In at least one embodiment and as described in greater detail below, HIP / HCC compilation commands 3664 configure HCC 3660 to compile HIP source code 3630 using, without limitation, HCC headers and a HIP / HCC runtime library. In at least one embodiment and in response to HIP / HCC compilation commands 3664, HCC 3660 generates host executable code 3670(2) and HCC device executable code 3682 (represented with bubble annotation C5). In at least one embodiment and as represented by bubble annotation C6, host executable code 3670(2) and HCC device executable code 3682 can be executed on CPU 3690 and GPU 3692, respectively. Figure 36C In at least one embodiment and as described in greater detail below, HIP / HCC compilation commands 3664 configure HCC 3660 to compile HIP source code 3630 using, without limitation, HCC headers and a HIP / HCC runtime library. In at least one embodiment and in response to HIP / HCC compilation commands 3664, HCC 3660 generates host executable code 3670(2) and HCC device executable code 3682 (represented with bubble annotation C5). In at least one embodiment and as represented by bubble annotation C6, host executable code 3670(2) and HCC device executable code 3682 can be executed on CPU 3690 and GPU 3692, respectively.

[0374] In at least one embodiment, after CUDA source code 3610 is converted to HIP source code 3630, HIP compiler driver 3640 can then be used to generate executable code for CUDA-enabled GPU 3694 or GPU 3692 without CUDA being re-implemented as HIP conversion tool 3620. In at least one embodiment, CUDA-to-HIP conversion tool 3620 converts CUDA source code 3610 to HIP source code 3630, which is then stored in memory. In at least one embodiment, HIP compiler driver 3640 then configures HCC 3660 to generate host executable code 3670(2) and HCC device executable code 3682 based on HIP source code 3630. In at least one embodiment, HIP compiler driver 3640 then configures CUDA compiler 3650 to generate host executable code 3670(1) and CUDA device executable code 3684 based on stored HIP source code 3630.

[0375] Figure 36B System 3604 is shown configured to compile and execute CUDA source code 3610 using CPU 3690 and CUDA-enabled GPU 3694, in accordance with at least one embodiment. In at least one embodiment, system 3604 includes, without limitation, CUDA source code 3610, CUDA-to-HIP conversion tool 3620, HIP source code 3630, HIP compiler driver 3640, CUDA compiler 3650, host executable code 3670(1), CUDA device executable code 3684, CPU 3690, and CUDA-enabled GPU 3694. Figure 36A System 3604 is shown configured to compile and execute CUDA source code 3610 using CPU 3690 and CUDA-enabled GPU 3694, in accordance with at least one embodiment. In at least one embodiment, system 3604 includes, without limitation, CUDA source code 3610, CUDA-to-HIP conversion tool 3620, HIP source code 3630, HIP compiler driver 3640, CUDA compiler 3650, host executable code 3670(1), CUDA device executable code 3684, CPU 3690, and CUDA-enabled GPU 3694.

[0376] In at least one embodiment and as previously mentioned herein Figure 36A As described, the CUDA source code 3610 includes, but is not limited to, any number (including zero) of global functions 3612, any number (including zero) of device functions 3614, any number (including zero) of host functions 3616, and any number (including zero) of host / device functions 3618. In at least one embodiment, the CUDA source code 3610 also includes, but is not limited to, any number of calls to any number of functions specified in any number of CUDA APIs.

[0377] In at least one embodiment, the CUDA to HIP conversion tool 3620 converts CUDA source code 3610 into HIP source code 3630. In at least one embodiment, the CUDA to HIP conversion tool 3620 converts each kernel call in the CUDA source code 3610 from CUDA syntax to HIP syntax, and converts any number of other CUDA calls in the CUDA source code 3610 into any number of other functionally similar HIP calls.

[0378] In at least one embodiment, the HIP compiler driver 3640 determines that the target device 3646 is CUDA enabled and generates HIP / NVCC compilation commands 3642. In at least one embodiment, the HIP compiler driver 3640 then configures the CUDA compiler 3650 via the HIP / NVCC compilation commands 3642 to compile HIP source code 3630. In at least one embodiment, as part of configuring the CUDA compiler 3650, the HIP compiler driver 3640 provides access to a HIP-to-CUDA translation header 3652. In at least one embodiment, the HIP-to-CUDA translation header 3652 translates an arbitrary number of mechanisms (e.g., functions) specified in an arbitrary number of HIP APIs into an arbitrary number of mechanisms specified in an arbitrary number of CUDA APIs. In at least one embodiment, the CUDA compiler 3650 uses the HIP-to-CUDA translation header 3652 in conjunction with a CUDA runtime library 3654 corresponding to the CUDA runtime API 3602 to generate host executable code 3670(1) and CUDA device executable code 3684. In at least one embodiment, host executable code 3670(1) and CUDA device executable code 3684 can then be executed on CPU 3690 and CUDA-enabled GPU 3694, respectively. In at least one embodiment, CUDA device executable code 3684 includes, but is not limited to, binary code. In at least one embodiment, CUDA device executable code 3684 includes, but is not limited to, PTX code, and is further compiled at runtime into binary code for a specific target device.

[0379] Figure 36C A system 3606 according to at least one embodiment is shown, the system 3606 being configured to compile and execute using a CPU 3690 and a GPU 3692 with CUDA disabled. Figure 36A The CUDA source code 3610. In at least one embodiment, the system 3606 includes, but is not limited to, the CUDA source code 3610, the CUDA to HIP conversion tool 3620, the HIP source code 3630, the HIP compiler driver 3640, the HCC 3660, the host executable code 3670(2), the HCC device executable code 3682, the CPU 3690, and the GPU 3692.

[0380] In at least one embodiment, and as previously mentioned herein Figure 36AAs described, the CUDA source code 3610 includes, but is not limited to, any number (including zero) of global functions 3612, any number (including zero) of device functions 3614, any number (including zero) of host functions 3616, and any number (including zero) of host / device functions 3618. In at least one embodiment, the CUDA source code 3610 also includes, but is not limited to, any number of calls to any number of functions specified in any number of CUDA APIs.

[0381] In at least one embodiment, the CUDA to HIP conversion tool 3620 converts CUDA source code 3610 into HIP source code 3630. In at least one embodiment, the CUDA to HIP conversion tool 3620 converts each kernel call in the CUDA source code 3610 from CUDA syntax to HIP syntax, and converts any number of other CUDA calls in the source code 3610 into any number of other functionally similar HIP calls.

[0382] In at least one embodiment, the HIP compiler driver 3640 then determines that the target device 3646 is not CUDA enabled and generates a HIP / HCC compilation command 3644. In at least one embodiment, the HIP compiler driver 3640 then configures the HCC 3660 to execute the HIP / HCC compilation command 3644, thereby compiling the HIP source code 3630. In at least one embodiment, the HIP / HCC compilation command 3644 configures the HCC 3660 to use, but not limited to, the HIP / HCC runtime library 3658 and the HCC header 3656 to generate host executable code 3670(2) and HCC device executable code 3682. In at least one embodiment, the HIP / HCC runtime library 3658 corresponds to the HIP runtime API 3632. In at least one embodiment, the HCC header 3656 includes, but is not limited to, any number and type of interoperability mechanisms for HIP and HCC. In at least one embodiment, host executable code 3670(2) and HCC device executable code 3682 can be executed on CPU 3690 and GPU 3692, respectively.

[0383] Figure 37 The diagram illustrates a method according to at least one embodiment. Figure 36CAn exemplary kernel is converted by the CUDA to HIP conversion tool 3620. In at least one embodiment, the CUDA source code 3610 divides the overall problem that a given kernel is designed to solve into relatively coarse subproblems that can be solved independently using thread blocks. In at least one embodiment, each thread block includes, but is not limited to, any number of threads. In at least one embodiment, each subproblem is divided into relatively small pieces that can be solved in parallel by the threads within the thread block. In at least one embodiment, threads within a thread block can cooperate by sharing data through shared memory and by coordinating memory accesses through synchronized execution.

[0384] In at least one embodiment, regarding Figure 37 At least one component shown or described is used to implement the combination. Figures 1-8 The described technologies and / or functions. In at least one embodiment, the at least one embodiment includes one or more security devices having information for detecting whether the security device has been restarted, thereby indicating that the one or more security devices have been tampered with.

[0385] In at least one embodiment, CUDA source code 3610 organizes thread blocks associated with a given kernel into a one-dimensional, two-dimensional, or three-dimensional grid of thread blocks. In at least one embodiment, each thread block includes, but is not limited to, any number of threads, and the grid includes, but is not limited to, any number of thread blocks.

[0386] In at least one embodiment, the kernel is a function in the device code defined using the "__global__" declaration specifier. In at least one embodiment, CUDA kernel startup syntax 3710 is used to specify the size of the mesh for executing the kernel for a given kernel call and the associated flow. In at least one embodiment, CUDA kernel startup syntax 3710 is specified as "KernelName <<<GridSize,BlockSize,SharedMemorySize,Stream> >>

[0387] (KernelArguments);". In at least one embodiment, the execution configuration syntax is a "<<<...>>>" construct, which is inserted between the kernel name ("KernelName") and the bracketed list of kernel parameters ("KernelArguments"). In at least one embodiment, the CUDA kernel startup syntax 3710 includes, but is not limited to, the CUDA startup function syntax instead of the execution configuration syntax.

[0388] In at least one embodiment, "GridSize" is of type dim3 and specifies the size and dimensions of the grid. In at least one embodiment, type dim3 is a CUDA-defined structure that includes, but is not limited to, unsigned integers x, y, and z. In at least one embodiment, if z is not specified, z defaults to 1. In at least one embodiment, if y is not specified, y defaults to 1.

[0389] In at least one embodiment, the number of thread blocks in the grid is equal to the product of GridSize.x, GridSize.y, and GridSize.z. In at least one embodiment, "BlockSize" is of type dim3 and specifies the size and dimensions of each thread block. In at least one embodiment, the number of threads per thread block is equal to the product of BlockSize.x, BlockSize.y, and BlockSize.z. In at least one embodiment, each thread executing the kernel has a unique thread ID, which can be accessed within the kernel via a built-in variable (e.g., "threadIdx").

[0390] In at least one embodiment, regarding CUDA kernel boot syntax 3710,

[0391] "SharedMemorySize" is an optional parameter that specifies the number of bytes dynamically allocated for each thread block in shared memory for a given kernel call, excluding statically allocated memory. In at least one embodiment and with respect to CUDA kernel startup syntax 3710, SharedMemorySize defaults to zero. In at least one embodiment and with respect to CUDA kernel startup syntax 3710, "stream" is an optional parameter that specifies an associated stream and defaults to zero to specify a default stream. In at least one embodiment, a stream is a sequence of commands executed sequentially (which may be issued by different host threads). In at least one embodiment, different streams may execute commands out of order or simultaneously relative to each other.

[0392] In at least one embodiment, CUDA source code 3610 includes, but is not limited to, kernel definitions and a main function for the exemplary kernel "MatAdd". In at least one embodiment, the main function is host code executed on the host and includes, but is not limited to, kernel calls that cause the kernel MatAdd to execute on the device. In at least one embodiment, as shown, the kernel MatAdd adds two matrices A and B of size NxN, where N is a positive integer, and stores the result in matrix C.

[0393] In at least one embodiment, the main function defines the `threadsPerBlock` variable as 16 x 16 and the `numBlocks` variable as N / 16 x N / 16. In at least one embodiment, the main function then specifies that the kernel calls `MatAdd<<`.<numBlocks,threadsPerBlock> >>(A, B, C);”. In at least one embodiment, and in accordance with CUDA kernel startup syntax 3710, a grid of thread blocks of size N / 16 × N / 16 is used to execute the kernel MatAdd, where each thread block is 16 × 16. In at least one embodiment, each thread block comprises 256 threads, creating a grid with enough blocks to allow one thread per matrix element, and each thread in the grid executes the kernel MatAdd to perform a pairwise addition.

[0394] In at least one embodiment, while converting CUDA source code 3610 to HIP source code 3630, the CUDA-to-HIP conversion tool 3620 converts each kernel call in the CUDA source code 3610 from CUDA kernel startup syntax 3710 to HIP kernel startup syntax 3720, and converts any number of other CUDA calls in the source code 3610 into any number of other functionally similar HIP calls. In at least one embodiment, the HIP kernel startup syntax 3720 is specified as “hipLaunchKernelGGL(KernelName,GridSize,BlockSize,SharedMemorySize,Stream,KernelArguments);”. In at least one embodiment, each of KernelName, GridSize, BlockSize, ShareMemorySize, Stream, and KernelArguments has the same meaning in the HIP kernel startup syntax 3720 as it does in the CUDA kernel startup syntax 3710 (as previously described herein). In at least one embodiment, the parameters SharedMemorySize and Stream are required in HIP kernel startup syntax 3720, but optional in CUDA kernel startup syntax 3710.

[0395] In at least one embodiment, in addition to the kernel call that causes the kernel MatAdd to execute on the device, Figure 37 The part of HIP source code 3630 described in the text is related to Figure 37The kernel MatAdd is identical to a portion of the CUDA source code 3610 depicted in the HIP source code 3630. In at least one embodiment, the kernel MatAdd is defined in the HIP source code 3630, having the same "__global__" declaration specifier as the kernel MatAdd defined in the CUDA source code 3610. In at least one embodiment, the kernel call in the HIP source code 3630 is "hipLaunchKernelGGL(MatAdd, numBlocks, threadsPerBlock, 0, 0, A, B, C);", while the corresponding kernel call in the CUDA source code 3610 is "MatAdd <<<numBlocks,threadsPerBlock> >>(A, B, C);”.

[0396] Figure 38 A more detailed description is provided according to at least one embodiment. Figure 36C The GPU 3692 is a CUDA-unenabled GPU. In at least one embodiment, the GPU 3692 was developed by AMD Inc. of Santa Clara City. In at least one embodiment, the GPU 3692 can be configured to perform computational operations in a highly parallel manner. In at least one embodiment, the GPU 3692 is configured to perform graphics pipeline operations, such as drawing commands, pixel operations, geometric calculations, and other operations associated with rendering images to a display. In at least one embodiment, the GPU 3692 is configured to perform graphics-independent operations. In at least one embodiment, the GPU 3692 is configured to perform both graphics-related and graphics-independent operations. In at least one embodiment, the GPU 3692 can be configured to execute device code included in HIP source code 3630.

[0397] In at least one embodiment, regarding Figure 38 At least one component shown or described is used to implement the combination. Figures 1-8 The described technologies and / or functions. In at least one embodiment, the at least one embodiment includes one or more security devices having information for detecting whether the security device has been restarted, thereby indicating that the one or more security devices have been tampered with.

[0398] In at least one embodiment, the GPU 3692 includes, but is not limited to, any number of programmable processing units 3820, an command processor 3810, an L2 cache 3822, a memory controller 3870, a DMA engine 3880(1), a system memory controller 3882, a DMA engine 3880(2), and a GPU controller 3884. In at least one embodiment, each programmable processing unit 3820 includes, but is not limited to, a workload manager 3830 and any number of compute units 3840. In at least one embodiment, the command processor 3810 reads commands from one or more command queues (not shown) and distributes the commands to the workload manager 3830. In at least one embodiment, for each programmable processing unit 3820, the associated workload manager 3830 distributes work to the compute units 3840 included in the programmable processing unit 3820. In at least one embodiment, each compute unit 3840 can execute any number of thread blocks, but each thread block executes on a single compute unit 3840. In at least one embodiment, the workgroup is a thread block.

[0399] In at least one embodiment, each computing unit 3840 includes, but is not limited to, any number of SIMD units 3850 and shared memory 3860. In at least one embodiment, each SIMD unit 3850 implements a SIMD architecture and is configured to perform operations in parallel. In at least one embodiment, each SIMD unit 3850 includes, but is not limited to, a vector ALU 3852 and a vector register file 3854. In at least one embodiment, each SIMD unit 3850 executes a different thread bundle. In at least one embodiment, a thread bundle is a group of threads (e.g., 16 threads), where each thread in the thread bundle belongs to a single thread block and is configured to process different datasets based on a single instruction set. In at least one embodiment, prediction can be used to disable one or more threads in a thread bundle. In at least one embodiment, a channel is a thread. In at least one embodiment, a work item is a thread. In at least one embodiment, a wavefront is a thread bundle. In at least one embodiment, different wavefronts in a thread block can be synchronized together and communicate via shared memory 3860. In at least one embodiment, computing unit 3840 includes one or more distributed shared memories (or distributed shared memory blocks) that enable direct streaming multiprocessor (SM) to streaming multiprocessor (SMs) operations related to loading, storing, and performing atomic operations across multiple SM shared memory blocks. Computing unit 3840 includes one or more clustered distributed shared memories (DSMEMs), which are memory blocks within a cluster that are enabled to directly access each other's shared memories.

[0400] In at least one embodiment, the programmable processing unit 3820 is referred to as a "shading engine". In at least one embodiment, in addition to the computing unit 3840, each programmable processing unit 3820 also includes, but is not limited to, any number of dedicated graphics hardware. In at least one embodiment, each programmable processing unit 3820 includes, but is not limited to, any number (including zero) of geometry processors, any number (including zero) of rasterizers, any number (including zero) of rendering backends, a workload manager 3830, and any number of computing units 3840.

[0401] In at least one embodiment, compute units 3840 share an L2 cache 3822. In at least one embodiment, the L2 cache 3822 is partitioned. In at least one embodiment, all compute units 3840 in the GPU 3692 have access to the GPU memory 3890. In at least one embodiment, a memory controller 3870 and a system memory controller 3882 facilitate data transfer between the GPU 3692 and the host, and a DMA engine 3880(1) enables asynchronous memory transfers between the GPU 3692 and the host. In at least one embodiment, a memory controller 3870 and a GPU controller 3884 facilitate data transfers between the GPU 3692 and other GPUs 3692, and a DMA engine 3880(2) enables asynchronous memory transfers between the GPU 3692 and other GPUs 3692.

[0402] In at least one embodiment, GPU 3692 includes, but is not limited to, any number and type of system interconnects that facilitate data and control transfers between any number and type of directly or indirectly linked components, either internally or externally to GPU 3692. In at least one embodiment, GPU 3692 includes, but is not limited to, any number and type of I / O interfaces (e.g., PCIe) coupled to any number and type of peripheral devices. In at least one embodiment, GPU 3692 may include, but is not limited to, any number (including zero) of display engines and any number (including zero) of multimedia engines. In at least one embodiment, GPU 3692 implements a memory subsystem that includes, but is not limited to, any number and type of memory controllers (e.g., memory controller 3870 and system memory controller 3882) and memory devices dedicated to a component or shared among multiple components (e.g., shared memory 3860). In at least one embodiment, GPU3692 implements a cache subsystem that includes, but is not limited to, one or more cache memories (e.g., L2 cache 3822), each cache memory being either private or shared among any number of components (e.g., SIMD unit 3850, compute unit 3840, and programmable processing unit 3820).

[0403] Figure 39 This illustrates how threads of an exemplary CUDA grid 3920, according to at least one embodiment, are mapped to... Figure 38 Different computational units 3840. In at least one embodiment, and for illustrative purposes only, grid 3920 has a GridSize of BX multiplied by BY multiplied by 1 and a BlockSize of TX multiplied by TY multiplied by 1. Therefore, in at least one embodiment, grid 3920 includes, but is not limited to, (BX*BY) thread blocks 3930, and each thread block 3930 includes, but is not limited to, (TX*TY) threads 3940. Threads 3940 in Figure 39 It is depicted as a curved arrow.

[0404] In at least one embodiment, regarding Figure 39 At least one component shown or described is used to implement the combination. Figures 1-8 The described technologies and / or functions. In at least one embodiment, the at least one embodiment includes one or more security devices having information for detecting whether the security device has been restarted, thereby indicating that the one or more security devices have been tampered with.

[0405] In at least one embodiment, grid 3920 is mapped to programmable processing unit 3820(1), which includes, but is not limited to, computing units 3840(1)-3840(C). In at least one embodiment, and as shown, (BJ*BY) thread block 3930 is mapped to computing unit 3840(1), and the remaining thread blocks 3930 are mapped to computing unit 3840(2). In at least one embodiment, each thread block 3930 may include, but is not limited to, any number of thread bundles, and each thread bundle is mapped to... Figure 38 Different SIMD units 3850.

[0406] In at least one embodiment, the thread bundles in a given thread block 3930 can be synchronized together and communicate via shared memory 3860 included in the associated computing unit 3840. For example, and in at least one embodiment, the thread bundles in thread block 3930(BJ, 1) can be synchronized together and communicate via shared memory 3860(1). For example, and in at least one embodiment, the thread bundles in thread block 3930(BJ+1, 1) can be synchronized together and communicate via shared memory 3860(2).

[0407] Figure 40 This document illustrates how to migrate existing CUDA code to data-parallel C++ code according to at least one embodiment. Data-parallel C++ (DPC++) can refer to an open, standards-based alternative to a single-architecture proprietary language that allows developers to reuse code across hardware targets (CPUs and accelerators, such as GPUs and FPGAs) and also perform custom tweaks for specific accelerators. DPC++ uses similar and / or the same C and C++ constructs as ISOC++, which developers may be familiar with. DPC++ incorporates the Khronos Group's standard SYCL to support data parallelism and heterogeneous programming. SYCL stands for Cross-Platform Abstraction Layer, which builds upon the underlying concepts, portability, and efficiency of OpenCL, enabling code for heterogeneous processors to be written in a "single-source" style using standard C++. SYCL enables single-source development, where C++ template functions can contain both host code and device code to build complex algorithms accelerated using OpenCL, and then reuse them throughout the source code for different types of data.

[0408] In at least one embodiment, regarding Figure 40 At least one component shown or described is used to implement the combination. Figures 1-8The described technologies and / or functions. In at least one embodiment, the at least one embodiment includes one or more security devices having information for detecting whether the security device has been restarted, thereby indicating that the one or more security devices have been tampered with.

[0409] In at least one embodiment, a DPC++ compiler is used to compile DPC++ source code that can be deployed across various hardware targets. In at least one embodiment, the DPC++ compiler is used to generate DPC++ applications that can be deployed across various hardware targets, and DPC++ compatibility tools are used to migrate CUDA applications to multi-platform programs in DPC++. In at least one embodiment, the DPC++ basic toolkit includes: a DPC++ compiler for deploying applications across various hardware targets; DPC++ libraries for improving productivity and performance on CPUs, GPUs, and FPGAs; DPC++ compatibility tools for migrating CUDA applications to multi-platform applications; and any suitable combination thereof.

[0410] In at least one embodiment, the DPC++ programming model simplifies one or more aspects related to programming CPUs and accelerators by using modern C++ features to express parallelism with a programming language called Data Parallel C++. The DPC++ programming language can be used for code reuse against hosts (e.g., CPUs) and accelerators (e.g., GPUs or FPGAs) using a single-source language, and clearly communicates execution and memory dependencies. Mappings within the DPC++ code can be used to translate applications to run on the hardware or set of hardware devices that best accelerate workloads. Even on platforms without available accelerators, the host can be used to simplify the development and debugging of device code.

[0411] In at least one embodiment, CUDA source code 4000 is provided as input to the DPC++ compatibility tool 4002 to generate human-readable DPC++ 4004. In at least one embodiment, the human-readable DPC++ 4004 includes inline comments generated by the DPC++ compatibility tool 4002, which guide developers on how and / or where to modify the DPC++ code to complete the coding and tune it to the desired performance 4006, thereby generating DPC++ source code 4008.

[0412] In at least one embodiment, CUDA source code 4000 is or includes a collection of human-readable source code in the CUDA programming language. In at least one embodiment, CUDA source code 4000 is human-readable source code using the CUDA programming language. In at least one embodiment, the CUDA programming language is an extension of the C++ programming language, including but not limited to mechanisms for defining device code and distinguishing between device code and host code. In at least one embodiment, the device code is source code that, after compilation, can be executed on a device (e.g., a GPU or FPGA) and may include one or more parallelizable workflows that can be executed on one or more processor cores of the device. In at least one embodiment, the device may be a processor optimized for parallel instruction processing, such as a CUDA-enabled GPU, GPU, or another GPGPU. In at least one embodiment, the host code is source code that, after compilation, can be executed on a host machine. In at least one embodiment, some or all of the host code and device code can be executed in parallel across CPU and GPU / FPGA. In at least one embodiment, the host is a processor optimized for sequential instruction processing, such as a CPU. Figure 40 The CUDA source code 4000 described is consistent with what is discussed elsewhere in this document.

[0413] In at least one embodiment, DPC++ compatibility tool 4002 refers to an executable tool, program, application, or any other suitable type of tool for facilitating the migration of CUDA source code 4000 to DPC++ source code 4008. In at least one embodiment, DPC++ compatibility tool 4002 is a command-line based code migration tool that can be used as part of the DPC++ toolkit for porting existing CUDA sources to DPC++. In at least one embodiment, DPC++ compatibility tool 4002 converts some or all of the source code of a CUDA application from CUDA to DPC++ and generates a result file, at least partially written in DPC++, referred to as human-readable DPC++ 4004. In at least one embodiment, human-readable DPC++ 4004 includes comments generated by DPC++ compatibility tool 4002 to indicate where user intervention may be required. In at least one embodiment, user intervention is necessary when CUDA source code 4000 calls a CUDA API that does not have a DPC++ API; other examples requiring user intervention will be discussed in more detail later.

[0414] In at least one embodiment, the workflow for migrating CUDA source code 4000 (e.g., an application or a portion thereof) includes creating one or more build database files; migrating CUDA to DPC++ using DPC++ compatibility tool 4002; completing the migration and verifying its correctness to generate DPC++ source code 4008; and compiling DPC++ source code 4008 using the DPC++ compiler to generate a DPC++ application. In at least one embodiment, the compatibility tool provides a utility that intercepts commands used during Makefile execution and stores them in the build database files. In at least one embodiment, the files are stored in JSON format. In at least one embodiment, intercepted build commands translate Makefile commands into DPC compatibility commands.

[0415] In at least one embodiment, intercept-build is a utility script that intercepts the build process to capture build options, macro definitions, and include paths, and writes this data to a build database file. In at least one embodiment, the build database file is a JSON file. In at least one embodiment, the DPC++ compatibility tool 4002 parses the build database and applies options when migrating input sources. In at least one embodiment, the use of intercept-build is optional but strongly recommended for Make or CMake-based environments. In at least one embodiment, the migration database includes commands, directories, and files: commands may include necessary build flags; directories may include paths to header files; and files may include paths to CUDA files.

[0416] In at least one embodiment, the DPC++ compatibility tool 4002 migrates CUDA code (e.g., applications) written in CUDA to DPC++ by generating DPC++ as much as possible. In at least one embodiment, the DPC++ compatibility tool 4002 is available as part of a toolkit. In at least one embodiment, the DPC++ toolkit includes an intercept-build tool. In at least one embodiment, the intercept-build tool creates a build database that captures build commands to migrate CUDA files. In at least one embodiment, the DPC++ compatibility tool 4002 uses the build database generated by the intercept-build tool to migrate CUDA code to DPC++. In at least one embodiment, non-CUDA C++ code and files are migrated as is. In at least one embodiment, the DPC++ compatibility tool 4002 generates a human-readable DPC++ 4004, which may be DPC++ code, such as that generated by the DPC++ compatibility tool 4002, that cannot be compiled by the DPC++ compiler and requires additional pipelines to verify incorrectly migrated code portions, and may involve manual intervention, such as intervention by a developer. In at least one embodiment, the DPC++ compatibility tool 4002 provides hints or tools embedded in the code to help developers manually migrate additional code that cannot be migrated automatically. In at least one embodiment, the migration is a one-time activity for a source file, project, or application.

[0417] In at least one embodiment, the DPC++ compatibility tool 40002 is capable of successfully migrating all portions of CUDA code to DPC++, and can simply include optional steps for manually verifying and tuning the performance of the generated DPC++ source code. In at least one embodiment, the DPC++ compatibility tool 4002 directly generates DPC++ source code 4008 that is compiled by the DPC++ compiler, without requiring or utilizing manual intervention to modify the DPC++ code generated by the DPC++ compatibility tool 4002. In at least one embodiment, the DPC++ compatibility tool generates compilable DPC++ code that developers can selectively tune based on performance, readability, maintainability, and various other considerations, or any combination thereof.

[0418] In at least one embodiment, one or more CUDA source files are migrated to DPC++ source files, at least in part, using DPC++ compatibility tool 4002. In at least one embodiment, the CUDA source code includes one or more header files, which may include CUDA header files. In at least one embodiment, the CUDA source files include text that can be printed.<cuda.h> header files and<stdio.h> Header file. In at least one embodiment, a portion of the CUDA source file for the vector addition kernel may be written as or related to:

[0419]

[0420]

[0421]

[0422] In at least one embodiment, and in conjunction with the CUDA source files presented above, the DPC++ compatibility tool 4002 parses the CUDA source code and replaces the header files with appropriate DPC++ and SYCL header files. In at least one embodiment, the DPC++ header files include helper declarations. In CUDA, there is a concept of thread IDs; correspondingly, in DPC++ or SYCL, there is a local identifier for each element.

[0423] In at least one embodiment, and in relation to the CUDA source file presented above, there are two vectors A and B, which are initialized and the result of vector addition is placed into vector C as part of VectorAddKernel(). In at least one embodiment, as part of migrating CUDA code to DPC++ code, the DPC++ compatibility tool 4002 converts the CUDA thread ID used to index worker elements to the SYCL standard addressing of the worker elements via the local ID. In at least one embodiment, the DPC++ code generated by the DPC++ compatibility tool 4002 can be optimized—for example, by reducing the dimension of nd_item, thereby increasing memory and / or processor utilization.

[0424] In at least one embodiment, and in conjunction with the CUDA source files presented above, memory allocation is migrated. In at least one embodiment, relying on SYCL concepts such as platform, device, context, and queue, cudaMalloc() is migrated to a unified shared memory SYCL call malloc_device() to which the device and context are passed. In at least one embodiment, the SYCL platform may have multiple devices (e.g., host and GPU devices); a device may have multiple queues to which jobs can be submitted; each device may have a context; and a context may have multiple devices and manage shared memory objects.

[0425] In at least one embodiment, and in conjunction with the CUDA source files presented above, the `main()` function invokes or calls `VectorAddKernel()` to add two vectors A and B and store the result in vector C. In at least one embodiment, the CUDA code calling `VectorAddKernel()` is replaced by DPC++ code to submit the kernel to the command queue for execution. In at least one embodiment, the command group handler `cgh` passes the data submitted to the queue, synchronization, and computation, and `parallel_for` is called for multiple global elements and multiple work items in the workgroup that call `VectorAddKernel()`.

[0426] In at least one embodiment, and in conjunction with the CUDA source files presented above, CUDA calls to copy device memory and then free memory for vectors A, B, and C are migrated to corresponding DPC++ calls. In at least one embodiment, C++ code (e.g., standard ISOC++ code for printing vectors of floating-point variables) is migrated as is without modification by the DPC++ compatibility tool 4002. In at least one embodiment, the DPC++ compatibility tool 4002 modifies the CUDA API used for memory setup and / or host calls to execute the kernel on an accelerated device. In at least one embodiment, and in conjunction with the CUDA source files presented above, the corresponding human-readable DPC++4004 (e.g., compilable) is written as or related to:

[0427]

[0428]

[0429] In at least one embodiment, human-readable DPC++4004 refers to the output generated by the DPC++ compatibility tool 4002 and can be optimized in one or another. In at least one embodiment, the human-readable DPC++4004 generated by the DPC++ compatibility tool 4002 can be manually edited by developers after migration to make it more maintainable, performant, or for other considerations. In at least one embodiment, the DPC++ code generated by the DPC++ compatibility tool 40002 (e.g., publicly available DPC++) can be optimized by removing duplicate calls to get_current_device() and / or get_default_context() for each malloc_device() call. In at least one embodiment, the DPC++ code generated above uses a 3D nd_range, which can be refactored to use only a single dimension, thereby reducing memory usage. In at least one embodiment, developers can manually edit the DPC++ code generated by the DPC++ compatibility tool 4002 to replace the use of unified shared memory with accessors. In at least one embodiment, the DPC++ compatibility tool 4002 has the option to change how it migrates CUDA code to DPC++ code. In at least one embodiment, the DPC++ compatibility tool 4002 is verbose because it uses a generic template to migrate CUDA code to DPC++ code, which is suitable for a wide range of situations.

[0430] In at least one embodiment, the CUDA to DPC++ migration workflow includes the following steps: preparing the migration using an intercept-build script; performing the migration of the CUDA project to DPC++ using the DPC++ compatibility tool 4002; reviewing and editing the source files of the migration to ensure their integrity and correctness; and compiling the final DPC++ code to generate the DPC++ application. In at least one embodiment, manual review of the DPC++ source code may be required in one or more scenarios, including but not limited to: the migrated API not returning error codes (CUDA code can return error codes that can subsequently be used by the application, but SYCL uses exceptions to report errors, therefore error codes are not used to expose errors); DPC++ does not support CUDA compute capability-related logic; statements cannot be deleted. In at least one embodiment, scenarios requiring manual intervention in the DPC++ code may include, but are not limited to: replacing error code logic with (*,0) code or commenting it out; equivalent DPC++ APIs being unavailable; CUDA compute capability-related logic; hardware-related APIs (clock()); APIs lacking unsupported features; performing time measurement logic; handling built-in vector type conflicts; migrating the cuBLAS API; and more.

[0431] In at least one embodiment, one or more techniques described herein utilize the oneAPI programming model. In at least one embodiment, the oneAPI programming model refers to a programming model for interacting with various computing accelerator architectures. In at least one embodiment, oneAPI refers to an application programming interface (API) designed to interact with various computing accelerator architectures. In at least one embodiment, the oneAPI programming model utilizes the DPC++ programming language. In at least one embodiment, the DPC++ programming language refers to a high-level language for data-parallel programming productivity. In at least one embodiment, the DPC++ programming language is at least partially based on the C and / or C++ programming languages. In at least one embodiment, the oneAPI programming model is a programming model such as those developed by Intel Corporation of Santa Clara, California.

[0432] In at least one embodiment, oneAPI and / or the oneAPI programming model are used to interact with various accelerators, GPUs, processors, and / or their variants and architectures. In at least one embodiment, oneAPI includes a set of libraries that implement various functions. In at least one embodiment, oneAPI includes at least the oneAPIDPC++ library, the oneAPI math kernel library, the oneAPI data analysis library, the oneAPI deep neural network library, the oneAPI collection communication library, the oneAPI thread building block library, the oneAPI video processing library, and / or their variants.

[0433] In at least one embodiment, the oneAPIDPC++ library (also known as oneDPL) is a library that implements algorithms and functions to accelerate DPC++ kernel programming. In at least one embodiment, oneDPL implements one or more Standard Template Library (STL) functions. In at least one embodiment, oneDPL implements one or more parallel STL functions. In at least one embodiment, oneDPL provides a set of library classes and functions, such as parallel algorithms, iterators, function object classes, range-based APIs, and / or their variants. In at least one embodiment, oneDPL implements one or more classes and / or functions from the C++ Standard Library. In at least one embodiment, oneDPL implements one or more random number generator functions.

[0434] In at least one embodiment, the oneAPI math kernel library (also known as oneMKL) is a library that implements various optimized and parallelized routines for the various mathematical functions and / or operations. In at least one embodiment, oneMKL implements one or more Basic Linear Algebra Subroutines (BLAS) and / or Linear Algebra Encapsulations (LAPACK) dense linear algebra routines. In at least one embodiment, oneMKL implements one or more sparse BLAS linear algebra routines. In at least one embodiment, oneMKL implements one or more random number generators (RNGs). In at least one embodiment, oneMKL implements one or more vector mathematics (VM) routines for performing mathematical operations on vectors. In at least one embodiment, oneMKL implements one or more Fast Fourier Transform (FFT) functions.

[0435] In at least one embodiment, the oneAPI data analytics library (also known as oneDAL) is a library that implements various data analytics applications and distributed computing. In at least one embodiment, oneDAL implements various algorithms for preprocessing, transformation, analysis, modeling, validation, and decision-making in batch, online, and distributed computing processing modes. In at least one embodiment, oneDAL implements various C++ and / or Java APIs and connectors to one or more data sources. In at least one embodiment, oneDAL implements the DPC++ API extension to the traditional C++ interface and enables GPUs to be used for the various algorithms.

[0436] In at least one embodiment, the oneAPI deep neural network library (also referred to as oneDNN) is a library that implements various deep learning functions. In at least one embodiment, oneDNN implements various neural networks, machine learning and deep learning functions, algorithms and / or variations thereof.

[0437] In at least one embodiment, the oneAPI centralized communication library (also known as oneCCL) is a library for implementing various applications of deep learning and machine learning workloads. In at least one embodiment, oneCCL is built on top of lower-level communication middleware such as Message Passing Interface (MPI) and libfabrics. In at least one embodiment, oneCCL implements a set of deep learning-specific optimizations, such as prioritization, persistent operations, out-of-order execution, and / or variations thereof. In at least one embodiment, oneCCL implements individual CPU and GPU functionalities.

[0438] In at least one embodiment, the oneAPI thread building block library (also referred to as oneTBB) is a library that implements various parallelization processes for different applications. In at least one embodiment, oneTTB is used for task-based shared parallel programming on a host machine. In at least one embodiment, oneTTB implements general-purpose parallel algorithms. In at least one embodiment, oneTTB implements concurrent containers. In at least one embodiment, oneTTB implements a scalable memory allocator. In at least one embodiment, oneTTB implements a work-stealing task scheduler. In at least one embodiment, oneTTB implements low-level synchronization primitives. In at least one embodiment, oneTTB is compiler-independent and can be used on various processors, such as GPUs, PPUs, CPUs, and / or variants thereof.

[0439] In at least one embodiment, the oneAPI video processing library (also known as oneVPL) is a library for accelerating video processing in one or more applications. In at least one embodiment, oneVPL implements various video decoding, encoding, and processing functions. In at least one embodiment, oneVPL implements various functions for media pipelines on CPUs, GPUs, and other accelerators. In at least one embodiment, oneVPL implements device discovery and selection for media-centric and video analytics workloads. In at least one embodiment, oneVPL implements API primitives for zero-copy buffer sharing.

[0440] In at least one embodiment, the oneAPI programming model utilizes the DPC++ programming language. In at least one embodiment, the DPC++ programming language is a programming language that is, but is not limited to, a functionally similar version of the CUDA mechanism that defines device code and distinguishes between device code and host code. In at least one embodiment, the DPC++ programming language may include a subset of the functionality of a CUDA programming language. In at least one embodiment, the oneAPI programming model is used with the DPC++ programming language to perform one or more CUDA programming model operations.

[0441] Figure 41This is a system diagram illustrating a system 4100 for interfacing with an application 4102 to process data, according to at least one embodiment. In at least one embodiment, the application 4102 uses a Large Language Model (LLM) 4112 to generate output data 4120 based at least in part on input data 4110. In at least one embodiment, the input data 4110 is a text prompt. In at least one embodiment, the input data 4110 includes unstructured text. In at least one embodiment, the input data 4110 includes a series of tokens. In at least one embodiment, the tokens are part of the input data. In at least one embodiment, the tokens are words. In at least one embodiment, the tokens are characters. In at least one embodiment, the tokens are subwords. In at least one embodiment, the input data 4110 is formatted in Chat Markup Language (ChatML). In at least one embodiment, the input data 4110 is an image. In at least one embodiment, the input data 4110 is one or more video frames. In at least one embodiment, the input data 4110 is any other expressive medium.

[0442] In at least one embodiment, regarding Figure 41 At least one component shown or described is used to implement the combination. Figures 1-8 The described technologies and / or functions. In at least one embodiment, the at least one embodiment includes one or more security devices having information for detecting whether the security device has been restarted, thereby indicating that the one or more security devices have been tampered with.

[0443] In at least one embodiment, the large language model 4112 includes a deep neural network. In at least one embodiment, the deep neural network is a neural network having two or more layers. In at least one embodiment, the large language model 4112 includes a converter model. In at least one embodiment, the large language model 4112 includes a neural network configured to perform natural language processing. In at least one embodiment, the large language model 4112 is configured to process one or more data sequences. In at least one embodiment, the large language model 4112 is configured to process text. In at least one embodiment, the weights and biases of the large language model 4112 are configured to process text. In at least one embodiment, the large language model 4112 is configured to determine patterns in data to perform one or more natural language processing tasks. In at least one embodiment, the natural language processing task includes text generation. In at least one embodiment, the natural language processing task includes question answering. In at least one embodiment, performing the natural language processing task produces output data 4120.

[0444] In at least one embodiment, the processor uses input data 4110 to query a retrieval database 4114. In at least one embodiment, the retrieval database 4114 is a key-value store. In at least one embodiment, the retrieval database 4114 is a corpus used to train a large language model 4112. In at least one embodiment, the processor uses the retrieval database 4114 to provide updated information to the large language model 4112. In at least one embodiment, the retrieval database 4114 includes data from Internet sources. In at least one embodiment, the large language model 4112 does not use the retrieval database 4114 to perform inference.

[0445] In at least one embodiment, the encoder encodes input data 4110 into one or more feature vectors. In at least one embodiment, the encoder encodes input data 4110 into sentence embedding vectors. In at least one embodiment, the processor uses the sentence embedding vectors to perform nearest neighbor search to generate one or more neighbors 4116. In at least one embodiment, one or more neighbors 4116 are values ​​retrieved from database 4114 corresponding to a key including input data 4110. In at least one embodiment, one or more neighbors 4116 include text data. In at least one embodiment, the encoder 4118 encodes one or more neighbors 4116. In at least one embodiment, the encoder 4118 encodes one or more neighbors 4116 into text embedding vectors. In at least one embodiment, the encoder 4118 encodes one or more neighbors 4116 into sentence embedding vectors. In at least one embodiment, the large language model...

Claims

1. A processor, comprising: One or more circuits are used to detect whether the security device has been restarted using information stored in one or more storage locations of the security device.

2. The processor of claim 1, wherein the absence of the information in one or more storage locations indicates that the security device has been tampered with.

3. The processor of claim 1, wherein the information includes at least one of a random byte sequence, a string, a random number, or a counter.

4. The processor of claim 1, wherein the information in the one or more storage locations remains intact until the one or more security devices are restarted.

5. The processor of claim 1, wherein the one or more security devices include a Trusted Platform Module (TPM).

6. The processor of claim 1, wherein the one or more circuits are configured to store a copy of the information.

7. The processor of claim 1, wherein the one or more circuits are configured to cause the information to be stored in the one or more storage locations of the one or more security devices via a secure session.

8. The processor of claim 1, wherein the information is erased upon restarting the one or more security devices.

9. A method comprising: Information stored in one or more storage locations of one or more security devices is used to detect whether the security device has been restarted.

10. The method of claim 9, further comprising: An indication that the security device has been restarted is generated, at least in part, based on the absence of the information in one or more of the storage locations.

11. The method of claim 9, wherein the information includes a unique identifier.

12. The method according to claim 9, further comprising: When restarting one or more security devices, the information is replaced with a different message.

13. The method of claim 9, wherein each of the one or more security devices includes one or more registers that store hash values ​​to record the state of the computer.

14. The method of claim 9, wherein the central processing unit (CPU) creates a copy of the information and stores the copy in a secure storage device on the CPU.

15. The method of claim 9, wherein one or more circuits are configured to store a public key, the public key being used to compare with a public key stored on the one or more security devices to verify whether the one or more security devices are the intended security devices with which the CPU will communicate via a secure channel.

16. A system comprising: One or more processors are used to detect whether the security device has been restarted using information stored in one or more storage locations of the security device.

17. The system of claim 16, wherein the absence of the information in one or more storage locations indicates that the security device has been tampered with.

18. The system of claim 16, wherein the one or more processors are configured to create the information and store the information on the one or more processors and in the one or more storage locations of the one or more security devices.

19. The system of claim 16, wherein each of the one or more security devices includes one or more registers that store hash values ​​to record the state of the computer, wherein only a processor among the one or more processors having a public key that matches a public key stored on the one or more security devices is authorized to extend the hash value.

20. The system of claim 16, wherein the one or more security devices are restarted by a command issued through a session different from that of the processor in the one or more processors through which the information is stored in the one or more storage locations of the one or more security devices.

Images