CUDA Runtime

運行時在cudart庫中實現，該庫通過cudart靜態地鏈接到應用程序。

所有入口都有cuda的前綴。

正如在異構編程中提到的，CUDA編程模型假設一個由主機和設備組成的系統，每個設備都有自己的獨立內存。

Initialization

運行時沒有顯式的初始化函數。 它在第一次調用運行時函數 (更確切地說，是參考手冊中錯誤處理和版本管理部分的函數以外的任何函數) 時初始化 。

運行時為系統中的每個設備創建一個CUDA Context 。該上下文是該設備的primary context， 在該設備上需要活動上下文的第一個運行時函數時初始化 。 它在應用程序的所有主機線程之間共享 。 作為創建上下文的一部分，如果需要的話，設備代碼將被實時編譯并加載到設備內存中 。這一切都是透明的。如果需要，例如，為了驅動API的互操作性，可以從驅動API訪問設備的主上下文。

當主機線程調用cudaDeviceReset()時，這將銷毀主機線程當前操作的設備的 primary context (即在device Selection中定義的當前設備)。當前擁有該設備的任何主機線程的下一個運行時函數調用將為該設備創建一個新的 primary context。

注意:CUDA接口使用全局狀態，該狀態在主機程序啟動時初始化，在主機程序終止時銷毀。CUDA運行時和驅動程序無法檢測此狀態是否無效，因此在程序啟動或main后終止期間使用任何這些接口(隱式或顯式)將導致未定義的行為。

Device Memory

正如在異構編程中提到的，CUDA編程模型假設一個由主機和設備組成的系統，每個設備都有自己的獨立內存。內核在設備內存之外運行，因此運行時提供了分配、釋放和復制設備內存的函數，以及在主機內存和設備內存之間傳輸數據。

設備內存可以分配作為linear memory 或 CUDA arrays。

• CUDA arrays 是為 texture fetching 優化的不透明內存布局。
• Linear memory 在單一的統一地址空間中分配，這意味著分別分配的實體可以通過指針相互引用，例如，在二叉樹或鏈表中。地址空間的大小取決于主機系統(CPU)和使用的GPU的計算能力.

Graphics Interoperability 介紹了運行時提供的與兩個主要圖形API，OpenGL和 Direct3D互操作的各種功能。

Texture and Surface Memory 提供了紋理和表面存儲器空間，提供了訪問設備內存的另一種方式;它們還公開了GPU紋理硬件的一個子集。

Call Stack 提到了用于管理CUDA c++調用棧的運行時函數。

Error Checking 描述如何正確檢查運行時生成的錯誤。

Multi-Device System 展示了編程模型如何擴展到具有多個設備連接到同一主機的系統。

Asynchronous Concurrent Execution 描述了用于在系統的各個級別上支持異步并發執行的概念和API。

Page-Locked Host Memory 引入了頁鎖定主機內存，它需要在內核執行與主機和設備內存之間的數據傳輸重疊。

Shared Memory演示了如何使用線程層次結構中引入的共享內存來最大化性能。

Linear memory 通常使用 cudaMalloc() 分配，使用cudaFree()釋放，主機內存和設備內存之間的數據傳輸通常使用cudaMemcpy()完成。在kernel的vector加法代碼示例中，需要將vector從主機內存復制到設備內存:

// Device code

__global__

void

VecAdd

(

float

* A,

float

* B,

float

* C,

int

N)

{

int

i = blockDim.x * blockIdx.x + threadIdx.x;

if

(i < N)

C[i] = A[i] + B[i];

}

// Host code

int

main

()

{

int

N = ...;

size_t

size = N *

sizeof

(

float

);

// Allocate input vectors h_A and h_B in host memory

float

* h_A = (

float

*)

malloc

(size);

float

* h_B = (

float

*)

malloc

(size);

float

* h_C = (

float

*)

malloc

(size);

// Initialize input vectors

...

// Allocate vectors in device memory

float

* d_A;

cudaMalloc(&d_A, size);

float

* d_B;

cudaMalloc(&d_B, size);

float

* d_C;

cudaMalloc(&d_C, size);

// Copy vectors from host memory to device memory

cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);

cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

// Invoke kernel

int

threadsPerBlock =

256

;

int

blocksPerGrid =

(N + threadsPerBlock -

1

) / threadsPerBlock;

VecAdd<<>>(d_A, d_B, d_C, N);

// Copy result from device memory to host memory

// h_C contains the result in host memory

cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

// Free device memory

cudaFree(d_A);

cudaFree(d_B);

cudaFree(d_C);

// Free host memory

...

}

Linear memory 也可以通過cudaMallocPitch()和cudaMalloc3D()來分配。這些函數被推薦用于2D或3D數組的分配，因為它確保分配被適當填充，以滿足設備內存訪問中描述的對齊要求，因此在訪問行地址或在2D數組和設備內存的其他區域之間執行復制時(使用cudaMemcpy2D()和cudaMemcpy3D()函數)確保最佳性能。返回的pitch(或stride)必須用于訪問數組元素。

• 下面的代碼示例分配了一個 width x height 的二維浮點值數組，并展示了如何在設備代碼中循環遍歷數組元素:

// Host code

int

width =

64

, height =

64

;

float

* devPtr;

size_t

pitch;

cudaMallocPitch(&devPtr, &pitch,

width *

sizeof

(

float

), height);

MyKernel<<<

100

,

512

>>>(devPtr, pitch, width, height);

// Device code

__global__

void

MyKernel

(

float

* devPtr,

size_t

pitch,

int

width,

int

height)

{

for

(

int

r =

0

; r < height; ++r) {

float

* row = (

float

*)((

char

*)devPtr + r * pitch);

for

(

int

c =

0

; c < width; ++c) {

float

element = row[c];

}

}

}

• 下面的代碼示例分配了一個 width x height x depth 的浮點值3D數組，并展示了如何在設備代碼中循環遍歷數組元素:

// Host code

int

width =

64

, height =

64

, depth =

64

;

cudaExtent extent = make_cudaExtent(width *

sizeof

(

float

),

height, depth);

cudaPitchedPtr devPitchedPtr;

cudaMalloc3D(&devPitchedPtr, extent);

MyKernel<<<

100

,

512

>>>(devPitchedPtr, width, height, depth);

// Device code

__global__

void

MyKernel

(cudaPitchedPtr devPitchedPtr,

int

width,

int

height,

int

depth)

{

char

* devPtr = devPitchedPtr.ptr;

size_t

pitch = devPitchedPtr.pitch;

size_t

slicePitch = pitch * height;

for

(

int

z =

0

; z < depth; ++z) {

char

* slice = devPtr + z * slicePitch;

for

(

int

y =

0

; y < height; ++y) {

float

* row = (

float

*)(slice + y * pitch);

for

(

int

x =

0

; x < width; ++x) {

float

element = row[x];

}

}

}

}

下面的代碼示例演示了通過運行時API訪問全局變量的各種方法:

__constant__

float

constData[

256

];

float

data[

256

];

cudaMemcpyToSymbol(constData, data,

sizeof

(data));

cudaMemcpyFromSymbol(data, constData,

sizeof

(data));

__device__

float

devData;

float

value =

3.14f

;

cudaMemcpyToSymbol(devData, &value,

sizeof

(

float

));

__device__

float

* devPointer;

float

* ptr;

cudaMalloc(&ptr,

256

*

sizeof

(

float

));

cudaMemcpyToSymbol(devPointer, &ptr,

sizeof

(ptr));

cudaGetSymbolAddress()用于檢索指向分配給在全局內存空間中聲明的變量的內存的地址。所分配內存的大小通過cudaGetSymbolSize()獲得。

Device Memory L2 Access Management

當CUDA內核重復訪問全局內存中的數據區域時，可以認為這種數據訪問是 persisting 。

另一方面，如果數據只被訪問一次，則可以將這種數據訪問視為 streaming 。

從CUDA 11.0開始，具有8.0及以上計算能力的設備能夠影響L2緩存中的數據持久性，從而可能提供更高的帶寬和更低的全局內存訪問延遲。

L2 cache Set-Aside for Persisting Accesses

L2緩存的一部分可以被預留出來，用于持久化對全局內存的數據訪問 。持久化訪問優先使用L2緩存的預留部分，而正常的或流的全局內存訪問只能在持久化訪問未使用時使用L2的這部分。

用于持久化訪問的L2緩存預留大小可以在限制范圍內進行調整:

cudaGetDeviceProperties(&prop, device_id);

size_t

size = min(

int

(prop.l2CacheSize *

0.75

), prop.persistingL2CacheMaxSize);

cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, size);

/* set-aside 3/4 of L2 cache for persisting accesses or the max allowed*/

當GPU配置為MIG (Multi-Instance GPU)模式時，L2緩存預留功能不可用。

當使用多進程服務(MPS)時，L2緩存預留大小不能通過cudaDeviceSetLimit來改變。相反，只能在啟動MPS服務器時通過環境變量CUDA_DEVICE_DEFAULT_PERSISTING_L2_CACHE_PERCENTAGE_LIMIT指定預留大小。

L2 Policy for Persisting Accesses

訪問策略窗口指定全局內存的連續區域和L2緩存中的持久性屬性，以便在該區域內進行訪問。

下面的代碼示例展示了如何使用CUDA流設置L2持久化訪問窗口。

• CUDA Stream Example

cudaStreamAttrValue stream_attribute;

// Stream level attributes data structure

stream_attribute.accessPolicyWindow.base_ptr =

reinterpret_cast

<

void

*>(ptr);

// Global Memory data pointer

stream_attribute.accessPolicyWindow.num_bytes = num_bytes;

// Number of bytes for persistence access.

// (Must be less than cudaDeviceProp::accessPolicyMaxWindowSize)

stream_attribute.accessPolicyWindow.hitRatio =

0.6

;

// Hint for cache hit ratio

stream_attribute.accessPolicyWindow.hitProp = cudaAccessPropertyPersisting;

// Type of access property on cache hit

stream_attribute.accessPolicyWindow.missProp = cudaAccessPropertyStreaming;

// Type of access property on cache miss.

//Set the attributes to a CUDA stream of type cudaStream_t

cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &stream_attribute);

當內核隨后在CUDA stream 中執行時，全局內存范圍 [ptr..ptr+num_bytes] 內的內存訪問比訪問其他全局內存位置更有可能持久存在L2緩存中。

• CUDA GraphKernelNode Example

cudaKernelNodeAttrValue node_attribute;

// Kernel level attributes data structure

node_attribute.accessPolicyWindow.base_ptr =

reinterpret_cast

<

void

*>(ptr);

// Global Memory data pointer

node_attribute.accessPolicyWindow.num_bytes = num_bytes;

// Number of bytes for persistence access.

// (Must be less than cudaDeviceProp::accessPolicyMaxWindowSize)

node_attribute.accessPolicyWindow.hitRatio =

0.6

;

// Hint for cache hit ratio

node_attribute.accessPolicyWindow.hitProp = cudaAccessPropertyPersisting;

// Type of access property on cache hit

node_attribute.accessPolicyWindow.missProp = cudaAccessPropertyStreaming;

// Type of access property on cache miss.

//Set the attributes to a CUDA Graph Kernel node of type cudaGraphNode_t

cudaGraphKernelNodeSetAttribute(node, cudaKernelNodeAttributeAccessPolicyWindow, &node_attribute);

可以使用hitRatio參數指定接收hitProp屬性的訪問的比例。在上面的兩個示例中，全局內存區域中60%的內存訪問[ptr..ptr+num_bytes]具有持久化屬性，40%的內存訪問具有流屬性。哪些特定的內存訪問被分類為持久化(hitProp)是隨機的，概率近似于hitRatio;概率分布取決于硬件架構和內存大小。

例如，如果L2預留緩存大小為16KB，而accessPolicyWindow中的num_bytes為32KB:

• 當命中率為0.5時，硬件將隨機選擇32KB窗口中的16KB指定為持久化并緩存到預留的L2緩存區。
• 當hitRatio為1.0時，硬件將嘗試將整個32KB窗口緩存到預留的L2緩存區。由于預留區域比窗口小，緩存行將被刪除，以將最近使用的16KB數據保存在L2緩存的預留部分。

因此，可以使用hitRatio來避免緩存線的抖動，并從總體上減少移動到L2緩存和移出的數據量。

hitRatio值低于1.0可用于手動控制與并發CUDA流不同的accessPolicyWindows可以在L2中緩存的數據量。例如，設L2預留緩存大小為16KB;在兩個不同的CUDA流中的兩個并發內核，每個都具有16KB的accessPolicyWindow，并且都具有1.0的hitRatio值，在競爭共享的L2資源時，可能會驅逐彼此的緩存線。但是，如果兩個accessPolicyWindows的hitRatio值都是0.5，它們就不太可能驅逐自己的或彼此的持久化緩存行。

L2 Access Properties

為不同的全局內存數據訪問定義了三種類型的訪問屬性:

1. cudaAccessPropertyStreaming:帶有streaming屬性的內存訪問不太可能持久存在L2緩存中，因為這些訪問會優先被刪除。
2. cudaAccessPropertyPersisting:具有persisting屬性的內存訪問更有可能保存在L2緩存中，因為這些訪問優先保存在L2緩存的預留部分。
3. cudaAccessPropertyNormal: 這個訪問屬性強制重置之前應用的持久化訪問屬性到正常狀態。來自以前CUDA內核的具有持久化屬性的內存訪問可能會在預期使用之后很長時間內保留在L2緩存中。這種使用后持久化減少了不使用持久化屬性的后續內核可用的L2緩存量。使用cudaAccessPropertyNormal屬性重置訪問屬性窗口將刪除先前訪問的持久(優先保留)狀態，就像先前訪問沒有訪問屬性一樣。

L2 Persistence Example

下面的例子展示了如何為持久訪問預留L2緩存，通過CUDA流在CUDA內核中使用預留的L2緩存，然后重置L2緩存。

cudaStream_t stream;

cudaStreamCreate(&stream);

// Create CUDA stream

cudaDeviceProp prop;

// CUDA device properties variable

cudaGetDeviceProperties( &prop, device_id);

// Query GPU properties

size_t

size = min(

int

(prop.l2CacheSize *

0.75

) , prop.persistingL2CacheMaxSize );

cudaDeviceSetLimit( cudaLimitPersistingL2CacheSize, size);

// set-aside 3/4 of L2 cache for persisting accesses or the max allowed

size_t

window_size = min(prop.accessPolicyMaxWindowSize, num_bytes);

// Select minimum of user defined num_bytes and max window size.

cudaStreamAttrValue stream_attribute;

// Stream level attributes data structure

stream_attribute.accessPolicyWindow.base_ptr =

reinterpret_cast

<

void

*>(data1);

// Global Memory data pointer

stream_attribute.accessPolicyWindow.num_bytes = window_size;

// Number of bytes for persistence access

stream_attribute.accessPolicyWindow.hitRatio =

0.6

;

// Hint for cache hit ratio

stream_attribute.accessPolicyWindow.hitProp = cudaAccessPropertyPersisting;

// Persistence Property

stream_attribute.accessPolicyWindow.missProp = cudaAccessPropertyStreaming;

// Type of access property on cache miss

cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &stream_attribute);

// Set the attributes to a CUDA Stream

for

(

int

i =

0

; i <

10

; i++) {

cuda_kernelA<<

`

0

,stream>>>(data1);

// This data1 is used by a kernel multiple times

}

// [data1 + num_bytes) benefits from L2 persistence

cuda_kernelB<<,block_size,<>

0

,stream>>>(data1);

// A different kernel in the same stream can also benefit

// from the persistence of data1

stream_attribute.accessPolicyWindow.num_bytes =

0

;

// Setting the window size to 0 disable it

cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &stream_attribute);

// Overwrite the access policy attribute to a CUDA Stream

cudaCtxResetPersistingL2Cache();

// Remove any persistent lines in L2

cuda_kernelC<<,block_size,<>

0

,stream>>>(data2);

// data2 can now benefit from full L2 in normal mode

Reset L2 Access to Normal

來自上一個CUDA內核的持久化L2緩存線可能在它被使用后很長一段時間內持久化L2。因此，對于流或正常內存訪問來說，L2緩存的正常優先級重置為正常是很重要的。有三種方法可以將持久化訪問重置為正常狀態。

使用訪問屬性cudaAccessPropertyNormal設置先前的持久化內存區域。

通過調用cudaCtxResetPersistingL2Cache()將所有持久化L2緩存線重置為正常。

最終未碰觸的線路會自動重置為正常。由于自動復位發生所需的時間長度不確定，因此強烈不鼓勵依賴自動復位。

Manage Utilization of L2 set-aside cache

在不同的CUDA流中并發執行的多個CUDA內核可能會為它們的流分配不同的訪問策略窗口。然而， L2預留緩存部分在所有這些并發CUDA內核之間共享 。因此， 這個預留緩存部分的凈利用率是所有并發內核單獨使用的總和 。當持久化訪問的量超過預留的L2緩存容量時，將內存訪問指定為持久化訪問的好處就會減少。

為了管理預留的L2緩存部分的利用率，應用程序必須考慮以下因素:

L2預留緩存的大小。

可以并發執行的CUDA內核。

可并發執行的所有CUDA內核的訪問策略窗口。

需要在何時以及如何重置L2，以允許normal或streaming訪問以同等優先級利用之前設置的L2緩存。

Query L2 cache Properties

與L2緩存相關的屬性是cudaDeviceProp結構的一部分，可以使用CUDA運行時API cudaGetDeviceProperties查詢.

CUDA設備屬性包括:

l2CacheSize: GPU上可用的L2緩存量。

persistingL2CacheMaxSize:可為持久內存訪問預留的L2緩存的最大數量。

accessPolicyMaxWindowSize:訪問策略窗口的最大大小。

Control L2 Cache Set-Aside Size for Persisting Memory Access

使用CUDA運行時API cudaDeviceGetLimit查詢用于持久化內存訪問的L2預留緩存大小，并使用CUDA運行時API cudaDeviceSetLimit作為cudaLimit進行設置。該限制的最大值為cudaDeviceProp::persistingL2CacheMaxSize。

enum

cudaLimit {

/* other fields not shown */

cudaLimitPersistingL2CacheSize

};

`