CUDA Runtime

運(yùn)行時(shí)在cudart庫(kù)中實(shí)現(xiàn)，該庫(kù)通過(guò)cudart靜態(tài)地鏈接到應(yīng)用程序。

所有入口都有cuda的前綴。

正如在異構(gòu)編程中提到的，CUDA編程模型假設(shè)一個(gè)由主機(jī)和設(shè)備組成的系統(tǒng)，每個(gè)設(shè)備都有自己的獨(dú)立內(nèi)存。

Initialization

運(yùn)行時(shí)沒(méi)有顯式的初始化函數(shù)。 它在第一次調(diào)用運(yùn)行時(shí)函數(shù) (更確切地說(shuō)，是參考手冊(cè)中錯(cuò)誤處理和版本管理部分的函數(shù)以外的任何函數(shù)) 時(shí)初始化 。

運(yùn)行時(shí)為系統(tǒng)中的每個(gè)設(shè)備創(chuàng)建一個(gè)CUDA Context 。該上下文是該設(shè)備的primary context， 在該設(shè)備上需要活動(dòng)上下文的第一個(gè)運(yùn)行時(shí)函數(shù)時(shí)初始化 。 它在應(yīng)用程序的所有主機(jī)線程之間共享 。 作為創(chuàng)建上下文的一部分，如果需要的話，設(shè)備代碼將被實(shí)時(shí)編譯并加載到設(shè)備內(nèi)存中 。這一切都是透明的。如果需要，例如，為了驅(qū)動(dòng)API的互操作性，可以從驅(qū)動(dòng)API訪問(wèn)設(shè)備的主上下文。

當(dāng)主機(jī)線程調(diào)用cudaDeviceReset()時(shí)，這將銷毀主機(jī)線程當(dāng)前操作的設(shè)備的 primary context (即在device Selection中定義的當(dāng)前設(shè)備)。當(dāng)前擁有該設(shè)備的任何主機(jī)線程的下一個(gè)運(yùn)行時(shí)函數(shù)調(diào)用將為該設(shè)備創(chuàng)建一個(gè)新的 primary context。

注意:CUDA接口使用全局狀態(tài)，該狀態(tài)在主機(jī)程序啟動(dòng)時(shí)初始化，在主機(jī)程序終止時(shí)銷毀。CUDA運(yùn)行時(shí)和驅(qū)動(dòng)程序無(wú)法檢測(cè)此狀態(tài)是否無(wú)效，因此在程序啟動(dòng)或main后終止期間使用任何這些接口(隱式或顯式)將導(dǎo)致未定義的行為。

Device Memory

正如在異構(gòu)編程中提到的，CUDA編程模型假設(shè)一個(gè)由主機(jī)和設(shè)備組成的系統(tǒng)，每個(gè)設(shè)備都有自己的獨(dú)立內(nèi)存。內(nèi)核在設(shè)備內(nèi)存之外運(yùn)行，因此運(yùn)行時(shí)提供了分配、釋放和復(fù)制設(shè)備內(nèi)存的函數(shù)，以及在主機(jī)內(nèi)存和設(shè)備內(nèi)存之間傳輸數(shù)據(jù)。

設(shè)備內(nèi)存可以分配作為linear memory 或 CUDA arrays。

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

Graphics Interoperability 介紹了運(yùn)行時(shí)提供的與兩個(gè)主要圖形API，OpenGL和 Direct3D互操作的各種功能。

Texture and Surface Memory 提供了紋理和表面存儲(chǔ)器空間，提供了訪問(wèn)設(shè)備內(nèi)存的另一種方式;它們還公開(kāi)了GPU紋理硬件的一個(gè)子集。

Call Stack 提到了用于管理CUDA c++調(diào)用棧的運(yùn)行時(shí)函數(shù)。

Error Checking 描述如何正確檢查運(yùn)行時(shí)生成的錯(cuò)誤。

Multi-Device System 展示了編程模型如何擴(kuò)展到具有多個(gè)設(shè)備連接到同一主機(jī)的系統(tǒng)。

Asynchronous Concurrent Execution 描述了用于在系統(tǒng)的各個(gè)級(jí)別上支持異步并發(fā)執(zhí)行的概念和API。

Page-Locked Host Memory 引入了頁(yè)鎖定主機(jī)內(nèi)存，它需要在內(nèi)核執(zhí)行與主機(jī)和設(shè)備內(nèi)存之間的數(shù)據(jù)傳輸重疊。

Shared Memory演示了如何使用線程層次結(jié)構(gòu)中引入的共享內(nèi)存來(lái)最大化性能。

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

// 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 也可以通過(guò)cudaMallocPitch()和cudaMalloc3D()來(lái)分配。這些函數(shù)被推薦用于2D或3D數(shù)組的分配，因?yàn)樗_保分配被適當(dāng)填充，以滿足設(shè)備內(nèi)存訪問(wèn)中描述的對(duì)齊要求，因此在訪問(wèn)行地址或在2D數(shù)組和設(shè)備內(nèi)存的其他區(qū)域之間執(zhí)行復(fù)制時(shí)(使用cudaMemcpy2D()和cudaMemcpy3D()函數(shù))確保最佳性能。返回的pitch(或stride)必須用于訪問(wèn)數(shù)組元素。

• 下面的代碼示例分配了一個(gè) width x height 的二維浮點(diǎn)值數(shù)組，并展示了如何在設(shè)備代碼中循環(huán)遍歷數(shù)組元素:

// 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];

}

}

}

• 下面的代碼示例分配了一個(gè) width x height x depth 的浮點(diǎn)值3D數(shù)組，并展示了如何在設(shè)備代碼中循環(huán)遍歷數(shù)組元素:

// 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];

}

}

}

}

下面的代碼示例演示了通過(guò)運(yùn)行時(shí)API訪問(wèn)全局變量的各種方法:

__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()用于檢索指向分配給在全局內(nèi)存空間中聲明的變量的內(nèi)存的地址。所分配內(nèi)存的大小通過(guò)cudaGetSymbolSize()獲得。

Device Memory L2 Access Management

當(dāng)CUDA內(nèi)核重復(fù)訪問(wèn)全局內(nèi)存中的數(shù)據(jù)區(qū)域時(shí)，可以認(rèn)為這種數(shù)據(jù)訪問(wèn)是 persisting 。

另一方面，如果數(shù)據(jù)只被訪問(wèn)一次，則可以將這種數(shù)據(jù)訪問(wèn)視為 streaming 。

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

L2 cache Set-Aside for Persisting Accesses

L2緩存的一部分可以被預(yù)留出來(lái)，用于持久化對(duì)全局內(nèi)存的數(shù)據(jù)訪問(wèn) 。持久化訪問(wèn)優(yōu)先使用L2緩存的預(yù)留部分，而正常的或流的全局內(nèi)存訪問(wèn)只能在持久化訪問(wèn)未使用時(shí)使用L2的這部分。

用于持久化訪問(wèn)的L2緩存預(yù)留大小可以在限制范圍內(nèi)進(jìn)行調(diào)整:

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*/

當(dāng)GPU配置為MIG (Multi-Instance GPU)模式時(shí)，L2緩存預(yù)留功能不可用。

當(dāng)使用多進(jìn)程服務(wù)(MPS)時(shí)，L2緩存預(yù)留大小不能通過(guò)cudaDeviceSetLimit來(lái)改變。相反，只能在啟動(dòng)MPS服務(wù)器時(shí)通過(guò)環(huán)境變量CUDA_DEVICE_DEFAULT_PERSISTING_L2_CACHE_PERCENTAGE_LIMIT指定預(yù)留大小。

L2 Policy for Persisting Accesses

訪問(wèn)策略窗口指定全局內(nèi)存的連續(xù)區(qū)域和L2緩存中的持久性屬性，以便在該區(qū)域內(nèi)進(jìn)行訪問(wèn)。

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

• 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);

當(dāng)內(nèi)核隨后在CUDA stream 中執(zhí)行時(shí)，全局內(nèi)存范圍 [ptr..ptr+num_bytes] 內(nèi)的內(nèi)存訪問(wèn)比訪問(wèn)其他全局內(nèi)存位置更有可能持久存在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參數(shù)指定接收hitProp屬性的訪問(wèn)的比例。在上面的兩個(gè)示例中，全局內(nèi)存區(qū)域中60%的內(nèi)存訪問(wèn)[ptr..ptr+num_bytes]具有持久化屬性，40%的內(nèi)存訪問(wèn)具有流屬性。哪些特定的內(nèi)存訪問(wèn)被分類為持久化(hitProp)是隨機(jī)的，概率近似于hitRatio;概率分布取決于硬件架構(gòu)和內(nèi)存大小。

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

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

因此，可以使用hitRatio來(lái)避免緩存線的抖動(dòng)，并從總體上減少移動(dòng)到L2緩存和移出的數(shù)據(jù)量。

hitRatio值低于1.0可用于手動(dòng)控制與并發(fā)CUDA流不同的accessPolicyWindows可以在L2中緩存的數(shù)據(jù)量。例如，設(shè)L2預(yù)留緩存大小為16KB;在兩個(gè)不同的CUDA流中的兩個(gè)并發(fā)內(nèi)核，每個(gè)都具有16KB的accessPolicyWindow，并且都具有1.0的hitRatio值，在競(jìng)爭(zhēng)共享的L2資源時(shí)，可能會(huì)驅(qū)逐彼此的緩存線。但是，如果兩個(gè)accessPolicyWindows的hitRatio值都是0.5，它們就不太可能驅(qū)逐自己的或彼此的持久化緩存行。

L2 Access Properties

為不同的全局內(nèi)存數(shù)據(jù)訪問(wèn)定義了三種類型的訪問(wèn)屬性:

1. cudaAccessPropertyStreaming:帶有streaming屬性的內(nèi)存訪問(wèn)不太可能持久存在L2緩存中，因?yàn)檫@些訪問(wèn)會(huì)優(yōu)先被刪除。
2. cudaAccessPropertyPersisting:具有persisting屬性的內(nèi)存訪問(wèn)更有可能保存在L2緩存中，因?yàn)檫@些訪問(wèn)優(yōu)先保存在L2緩存的預(yù)留部分。
3. cudaAccessPropertyNormal: 這個(gè)訪問(wèn)屬性強(qiáng)制重置之前應(yīng)用的持久化訪問(wèn)屬性到正常狀態(tài)。來(lái)自以前CUDA內(nèi)核的具有持久化屬性的內(nèi)存訪問(wèn)可能會(huì)在預(yù)期使用之后很長(zhǎng)時(shí)間內(nèi)保留在L2緩存中。這種使用后持久化減少了不使用持久化屬性的后續(xù)內(nèi)核可用的L2緩存量。使用cudaAccessPropertyNormal屬性重置訪問(wèn)屬性窗口將刪除先前訪問(wèn)的持久(優(yōu)先保留)狀態(tài)，就像先前訪問(wèn)沒(méi)有訪問(wèn)屬性一樣。

L2 Persistence Example

下面的例子展示了如何為持久訪問(wèn)預(yù)留L2緩存，通過(guò)CUDA流在CUDA內(nèi)核中使用預(yù)留的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

來(lái)自上一個(gè)CUDA內(nèi)核的持久化L2緩存線可能在它被使用后很長(zhǎng)一段時(shí)間內(nèi)持久化L2。因此，對(duì)于流或正常內(nèi)存訪問(wèn)來(lái)說(shuō)，L2緩存的正常優(yōu)先級(jí)重置為正常是很重要的。有三種方法可以將持久化訪問(wèn)重置為正常狀態(tài)。

使用訪問(wèn)屬性cudaAccessPropertyNormal設(shè)置先前的持久化內(nèi)存區(qū)域。

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

最終未碰觸的線路會(huì)自動(dòng)重置為正常。由于自動(dòng)復(fù)位發(fā)生所需的時(shí)間長(zhǎng)度不確定，因此強(qiáng)烈不鼓勵(lì)依賴自動(dòng)復(fù)位。

Manage Utilization of L2 set-aside cache

在不同的CUDA流中并發(fā)執(zhí)行的多個(gè)CUDA內(nèi)核可能會(huì)為它們的流分配不同的訪問(wèn)策略窗口。然而， L2預(yù)留緩存部分在所有這些并發(fā)CUDA內(nèi)核之間共享 。因此， 這個(gè)預(yù)留緩存部分的凈利用率是所有并發(fā)內(nèi)核單獨(dú)使用的總和 。當(dāng)持久化訪問(wèn)的量超過(guò)預(yù)留的L2緩存容量時(shí)，將內(nèi)存訪問(wèn)指定為持久化訪問(wèn)的好處就會(huì)減少。

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

L2預(yù)留緩存的大小。

可以并發(fā)執(zhí)行的CUDA內(nèi)核。

可并發(fā)執(zhí)行的所有CUDA內(nèi)核的訪問(wèn)策略窗口。

需要在何時(shí)以及如何重置L2，以允許normal或streaming訪問(wèn)以同等優(yōu)先級(jí)利用之前設(shè)置的L2緩存。

Query L2 cache Properties

與L2緩存相關(guān)的屬性是cudaDeviceProp結(jié)構(gòu)的一部分，可以使用CUDA運(yùn)行時(shí)API cudaGetDeviceProperties查詢.

CUDA設(shè)備屬性包括:

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

persistingL2CacheMaxSize:可為持久內(nèi)存訪問(wèn)預(yù)留的L2緩存的最大數(shù)量。

accessPolicyMaxWindowSize:訪問(wèn)策略窗口的最大大小。

Control L2 Cache Set-Aside Size for Persisting Memory Access

使用CUDA運(yùn)行時(shí)API cudaDeviceGetLimit查詢用于持久化內(nèi)存訪問(wèn)的L2預(yù)留緩存大小，并使用CUDA運(yùn)行時(shí)API cudaDeviceSetLimit作為cudaLimit進(jìn)行設(shè)置。該限制的最大值為cudaDeviceProp::persistingL2CacheMaxSize。

enum

cudaLimit {

/* other fields not shown */

cudaLimitPersistingL2CacheSize

};

`