CUDA10.0官方文件的翻譯與學習之程式設計介面

coder_szc發表於2020-12-12
程式設計

目錄

背景

用nvcc編譯

編譯工作流

二進位制適配性

ptx適配性

應用適配性

C/C++適配性

64位適配性

cuda c執行時

初始化

裝置記憶體

共享記憶體

頁鎖主機記憶體

可移植記憶體

寫合併記憶體

對映記憶體

非同步併發執行

主機與裝置間的併發執行

併發核執行

資料遷移與核執行重疊

併發資料傳輸

流

多裝置系統

裝置列舉：

裝置選擇

流與事件行為

對等結點記憶體訪問

對等結點記憶體複製

統一虛地址空間

程式間通訊

錯誤檢查

呼叫棧

紋理記憶體和表面(surface)記憶體

紋理物件API

紋理引用API

兩位元組浮點紋理

分層紋理

立方圖紋理

紋理聚集

表面記憶體

表面引用API

立方圖表面

cuda陣列

讀寫一致性

圖形互動性

OpenGL互動性

Direct3D互動性

SLI互動

版本和適配性

計算模式

模式切換

針對Windows的Tesla計算叢集模式

結語

背景

在文章CUDA10.0官方文件的翻譯與學習之程式設計模型和CUDA10.0官方文件的翻譯與學習之介紹中我分別翻譯了CUDA10.0官方文件的前兩章，這篇文章我將翻譯此文件中最重要的部分——程式設計介面

cuda C為熟悉C程式語言的使用者提供了簡單的寫出可以被裝置執行的程式的方法，它由對C語言的簡單擴充套件和一個執行時的庫組成。核心的語言擴充套件已經在程式設計模型一章中介紹了，它允許程式設計師把核函式定義成C函式，然後使用新的語法來為每一次函式呼叫指定網格與執行緒塊的維度，對所有擴充套件的完整描述請參見官方手冊。任何包含這些擴充套件的原始檔都必須用nvcc編譯器編譯，詳情參見本章第一節——用nvcc編譯。

執行時庫在第一節第一部分編譯流程中介紹，它提供了在主機上執行的C函式，用來分配與回收裝置記憶體、在主機記憶體與裝置記憶體之間遷移資料、管理擁有多個裝置的系統等，關於執行時的完整描述請參見cuda參考手冊。

執行時庫是在底層C API、CUDA驅動API(也可以被應用獲取)上構建的，驅動API通過暴露底層概念（比如cuda上下文——類似於主機程式、cuda模組——類似於動態連結庫等）來提供另一個層面的控制。大多數應用不使用驅動API，因為他們不需要這種額外的控制，當使用執行時時，上下文、模組管理時隱式的，這樣程式碼會簡潔得多，驅動API的介紹和完整的描述請參見參考手冊

用nvcc編譯

核函式可以通過使用cuda指令集來編寫，cuda指令集又稱為ptx，它在ptx參考手冊中有所描述。當然，使用C這種高階程式語言效率更高，但不管怎樣，核函式都要通過nvcc編譯成二進位制程式碼以便在裝置上執行。nvcc是一個簡化編譯C或ptx程式碼的編譯驅動：它提供簡單且熟悉的命令列選項，然後通過呼叫實現不同編譯階段的工具集來執行它們。本節給出了nvcc工作流和命令選項的總覽，完成的介紹請參見nvcc使用者手冊。

編譯工作流

離線編譯：

    被nvcc編譯的原始檔由主機程式碼(在主機上執行的程式碼)和裝置程式碼(在裝置上執行的程式碼)組成，nvcc的基本工作流包括從主機程式碼中分離裝置程式碼，把裝置程式碼編譯成彙編形式(ptx程式碼)或二進位制形式(cubin物件)，然後通過把核函式呼叫時的<<<...>>>語法替換成有必要載入的cuda c執行時函式呼叫來修改主機程式碼，最後從ptx程式碼或者cubin物件中啟動每個編譯好的核函式。修改好的主機程式碼要麼以用來被別的工具編譯的C程式碼的形式輸出，要麼讓nvcc在最後的編譯階段呼叫主機編譯器輸出目的碼。而後應用就可以要麼連結編譯好的主機程式碼(一般如此)，要麼忽略所有修改過的主機程式碼，並使用cuda驅動API來載入ptx程式碼或者執行cubin檔案。

即時編譯：

    被應用載入的任何ptx程式碼在執行時被裝置驅動進一步編譯成二進位制程式碼，這就是即時編譯。即時編譯增大了應用載入時間，但允許應用從任何新裝置驅動對編譯器的改進中獲益，這也是應用執行在它編譯時還不存在的裝置上的唯一方法，詳情請參見應用適配性一節。當裝置驅動的即時編譯器為一些應用編譯ptx程式碼時，它會自動快取一份生成的二進位制程式碼的拷貝，以防止對此應用的在後續呼叫時重複編譯。這個快取（又稱之為計算快取）在裝置驅動更新時會自動失效，所以應用才能在安裝到裝置驅動中的新即時編譯器的改進中獲益。官方手冊中的cuda環境變數部分有關於可以用來控制即時編譯的環境變數的描述。

二進位制適配性

二進位制程式碼是架構相關的，cubin物件是通過使用指定目標架構的-code編譯選項生成的，例如用-code=sm_35編譯選項生成的二進位制程式碼是執行在計算能力為3.5的裝置上的。二進位制適應性在從小的副版本遷移到大的副版本的過程中被保證，但不會在從大的副版本遷移到小的副版本或者跨主版本的遷移過程中保證，換句話說，為計算能力X.y生成的cubin程式碼只會執行在計算能力X.z(z >= y)的裝置上。

ptx適配性

一些ptx指令只在計算能力更高的的裝置上被支援，例如偽執行緒混洗函式只在計算能力>=3.0的裝置上支援。編譯選項-arch指定了把C編譯成ptx程式碼時使用的計算能力。因此，包含偽執行緒混洗的程式碼必須用引數-arch=compute_30(或者更大)來進行編譯。被某個計算能力產生的ptx程式碼總是可以被編譯成相同或者更高計算能力的二進位制程式碼，注意從早期ptx版本編譯過來的二進位制程式碼可能不會使用某些硬體特徵，例如由計算能力6.0的ptx程式碼編譯成的計算能力7.0二進位制程式碼，就不會使用只能在7.0裝置上才能用的功能，比如張量核指令。所以，如果二進位制檔案是使用最新版本的ptx生成的話，它的效能可能會更好一些。

應用適配性

為了在指定計算能力的裝置上執行程式碼，應用必須載入和這種計算能力適配的二進位制或者ptx程式碼，參見上兩小節。特殊地，為了能夠在擁有更高計算能力的未來架構上執行程式碼(對於這種架構還不能生成二進位制程式碼)，應用就必須裝載可以被這些未來裝置即時編譯的ptx程式碼(請參見即時編譯部分)。哪種ptx或二進位制程式碼可以被嵌入到cuda c應用中可以由-arch和-code編譯選項分別指定，或者用-gencode包裝它們。但是，如果有多個arch引數，就只會嵌入值最大的ptx程式碼，舉例如下所示

nvcc x.cu -gencode arch=compute_35,code=sm_35 -gencode arch=compute_50,code=sm_50 -gencode arch=compute_60,code=\"compute_60,sm_60\"

這段命令將嵌入適配3.5、5.0和6.0的二進位制程式碼以及適配6.0的ptx程式碼，沒有顯式適配3.5和5.0，是因為6.0的ptx已經適配了這兩個版本了，另外最後面的code=\"compute_60,sm_60\"表示要生成的是6.0的胖二進位制(既可以即時編譯，也可以立馬執行)。主機程式碼的生成是為了在執行時自動選擇最合適的程式碼來載入和執行，例如在上面的例子中，要載入執行的程式碼就是：

    1、計算能力為3.5和3.7的裝置 ->3.5版本的二進位制程式碼；

    2、計算能力為5.0和5.2的裝置 ->5.0版本的二進位制程式碼；

    3、計算能力為6.0和6.1的裝置 ->6.0版本的二進位制程式碼；

    4、計算能力≥7.0的裝置 -> ptx程式碼，可以被進一步編譯成二進位制程式碼。

x.cu可以使用優化後的程式碼路徑，比如使用只在計算能力>=3.0的裝置上才能用的偽執行緒混洗操作，__CUDA_ARCH__巨集可以被用來區別基於不同計算能力的程式碼路徑，但它只在裝置程式碼中才能使用，比如當使用-arch=comput_35編譯時，__CUDA_ARCH__就等於350。

使用驅動API的應用必須編譯程式碼以分離檔案，然後在執行時明確載入執行最合適的檔案。Volta架構引入了能改變執行緒在GPU上排程的獨立執行緒排程機制，對於以之前架構上的smit排程執行的程式碼，獨立執行緒排程可能會改變參與的執行緒集合，從而導致錯誤的結果。為了使用包含獨立執行緒排程的路徑的同時保證正確遷移，Volta開發者可以用編譯選項-arch=compute_60 -code=sm_70來編譯pascal執行緒排程模組(引數意思是，生成適配6.0的ptx和直接能執行在7.0上的二進位制，如果要執行在6.0裝置上，6.0的ptx將會再編譯成6.0的二進位制，這時就會使用SIMT了)。

nvcc使用者手冊為-arch、-code和-gencode列出了各種簡寫，例如-arch=sm_35是arch=compute_35 -code=compute_35,sm_35的縮寫(也等於-gencode arch=comput_35, code=\"compute_35, sm_35\")，也就是要生成3.5的胖二進位制

C/C++適配性

編譯器前端來處理包括C++語法的cuda原始檔，在主機程式碼中支援所有的C++，但在裝置程式碼中只有部分C++的自己才被完全支援，請參見官方手冊中的C/C++語言支援部分

64位適配性

64位的nvcc編譯器以64位模式編譯裝置程式碼，比如指標都是64位的，被64位模式編譯的裝置程式碼只能支援同樣以64位模式編譯主機程式碼，對於32位的nvcc也是如此。但是，為了相容性，可以為32位編譯器加上-m64、為64位編譯器加上-m32來讓它們分別以64位和32位模式編譯裝置程式碼

cuda c執行時

執行時在cudart庫中被實現，這個庫通過靜態(cudart.lib或libcudart.a)或動態(cudart.dll或libcudart.so)的方式和應用連結，需用使用cudart.dll或者cudart.so做動態連結的應用經常把它們做為自己安裝包的一部分，只有在連結到相同的cuda執行時例項的元件之間進行的cuda執行時符號地址傳遞才是安全的，另外所有cuda執行時函式的字首都是cuda。

初始化

執行時庫裡沒有明確的初始化函式，當一個執行時函式(準確地說是除了裝置管理或版本管理之外的任何執行時函式)被呼叫時就會初始化，當我們需要為執行時函式計時或者把第一次呼叫生成的錯誤碼解釋到執行時時要記住這一點。初始化時，執行時會為系統中的每個裝置建立cuda上下文，這就是裝置的主上下文，它被應用的所有主機執行緒共享。作為主執行緒建立過程的一部分，裝置程式碼會被載入到裝置記憶體中，必要時會在載入前作即時編譯。所有這一切都是隱式進行的，執行時不會把主上下文暴露給應用。

當一個主機執行緒呼叫了cudaDeviceReset()函式，當前主機執行緒所在的裝置主上下文就會隨之銷燬，此裝置上任何主機執行緒進行下一次執行時函式呼叫會為這個裝置建立一個新主上下文。

裝置記憶體

如前文異構程式設計中所言，cuda程式設計模型假設系統由主機和裝置組成，兩者擁有自己獨立的記憶體。核函式並不直接操作記憶體，所以執行時提供了分配、回收、複製裝置記憶體的函式，以及在主機記憶體與裝置記憶體之間遷移資料的函式。

裝置記憶體能夠以線性記憶體或者cuda陣列的形式被分配。cuda陣列是用來做紋理讀取的不透明的記憶體佈局，我們會在紋理和表面記憶體一節中講述；線性記憶體存在於裝置中的一個40位的地址空間中，所以獨立分配的實體可以通過例如二叉樹這樣的指標相互引用。線性記憶體通常使用cudaMalloc()函式分配，用cudaFree()函式回收，主機與裝置之間的資料遷移一般用cudaMemcpy()函式進行，在向量加法樣例中剩餘程式碼裡，向量需要從主機記憶體複製到裝置記憶體中

// vecAdd.cu


#include <stdlib.h>
#include <stdio.h>
#include <cuda_runtime.h>

__global__ void VecAdd(float* A, float* B, float* C, int N) {
    int x = blockDim.x * blockIdx.x + threadIdx.x;
    if (x < N) {
        C[x] = A[x] + B[x];
    }
}

int main() {
    int N = 10;
    size_t size = N * sizeof(float);

    float* h_A = (float*)malloc(size);
    float* h_B = (float*)malloc(size);
    float* h_C = (float*)malloc(size);

    ... // 初始化h_A和h_B，賦值

    float* d_A;
    float* d_B;
    float* d_C;

    cudaMalloc(&d_A, size);
    cudaMalloc(&d_B, size);
    cudaMalloc(&d_C, size);

    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;

    VecAdd<<<blocksPerGrid, trheadsPerBlock>>>(d_A, d_B. d_C, N);

    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    ... // 釋放主機記憶體
}

線性記憶體也可以通過cudaMallocPitch()和cudaMalloc3D()函式分配，當要分配2維或3維陣列時我們推薦使用這兩個函式，因為這種分配的填充方式可以近似滿足裝置記憶體訪問一節描述的對齊要求，從而保證了訪問低地址或執行二維陣列在裝置記憶體其他區域進行復制(使用cudaMemcpy2D()或cudaMemcpy3D()函式)時的最好表現，返回的行寬或者步長必須被用來分配陣列，下面的例子展示了為二維浮點陣列分配空間以及如何迴圈遍歷：

int width = 64, height = 64;
float* devPtr;
size_t pitch; // 得到的行寬

cudaMallocPitch(&devPtr, &pitch, width * sizeof(float), height);

MyKernel<<<100, 512>>>(devPtr, pitch, width, height);

__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];
        }
    }
}

下面是分配並迴圈三維陣列的例子：

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

__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];
            }
        }
    }
}

參考手冊列舉了各種用來在cudaMalloc()、cudaMallocPith()和cudaMalloc3D()分配的線性空間、cuda陣列和為宣告在全域性或常量記憶體中的記憶體之間進行記憶體複製的函式。下面的例子描述了通過執行時API訪問全域性記憶體變數的各種方法：

__constant__ float constData[256]; // 裝置中的常量記憶體變數，必須寫到所有函式外面
float data[256];

cudaMemcpyToSymbol(constData, data, sizeof(data)); // symbol是指全域性或裝置記憶體中的變數，因此這個函式是寫入記憶體函式
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()得到

共享記憶體

我們可以通過使用__shared__記憶體識別符號來為一個變數分配共享記憶體空間，如執行緒層次一節中所述，共享記憶體要比全域性記憶體快得多，只要有機會就應該用共享記憶體訪問代替全域性記憶體訪問，我們用下面的矩陣相乘例子來講述。下面的程式碼是不使用共享記憶體來實現的矩陣乘法樣例：

typedef struct {
    int width;
    int height;
    float* data;
} Matrix;

#define BLOCK_SIZE 2 // 要確保能被A的高與B的寬整除

__global__ void MatMulKernel(const Matrix, const Matrix, Matrix);

void MatMaul(const Matrix A, const Matrix B, Matrix C) {
    Matrix dA;
    dA.height = A.height;
    dA.width = A.width;
    size_t size = dA.width * dA.height * sizeof(float);

    cudaMalloc(&dA.data, size);
    cudaMemcpy(dA.data, A.data, size, cudaMemcpyHostToDevice);

    Matrix dB;

    dB.height = B.height;
    dB.width = B.width;

    size = dB.width * dB.height * sizeof(float);

    cudaMalloc(&dB.data, size);
    cudaMemcpy(dB.data, B.data, size, cudaMemcpyHostToDevice);

    Matrix dC;

    dC.height = C.height;
    dC.width = C.width;

    size = dC.width * dC.height * sizeof(float);

    cudaMalloc(&dC.data, size);

    dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
    dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);

    MatMulKernel<<<dimGrid, dimBlock>>>(dA, dB, dC);

    cudaMemcpy(C.data, dC.data, size, cudaMemcpyDeviceToHost);

    cudaFree(dA.data);
    cudaFree(dB.data);
    cudaFree(dC.data);

}

__global__ void MatMulKernel(const Matrix A, const Matrix B, Matrix C) {
    float value = 0;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;

    for (int i = 0; i < A.width; i++) {
        value += A.data[row * A.width + i] * B.data[i * B.width + col];
    }

    C.data[row * C.width + col] = value;

}

這裡的矩陣乘法圖示如下

然後我們再寫一個使用共享記憶體的實現，在這個實現中，每一個執行緒塊要計算C中的一個子方陣Csub，如下圖所示

Csub等於A中子矩陣(A.width, block_size)和B中子矩陣(block_size, B.height)的乘積，為了把A和B載入到裝置資源中，這兩個矩陣被分解成了儘可能多的block_size方陣，因此Csub就是這些方陣的乘積和。每次計算Csub都是先把兩個相關的方陣從全域性記憶體中載入到共享記憶體中，此時每個執行緒載入每個矩陣的一個元素並計算Csub中的一個元素，然後把結果累加到一個暫存器中，所有的結果計算完成後就把結果寫回全域性記憶體裡。

typedef struct {
    int width;
    int height;
    int stride; // 遍歷步長，其實這裡就是行寬
    float* data;
} Matrix;

#define BLOCK_SIZE 2 // 子方陣邊長，也是執行緒塊邊長

__global__ void MatMulKernel(const Matrix, const Matrix, Matrix);

__device__ float GetElement(const Matrix A, int row, int col) {
    return A.data[row * A.stride + col]; // 讀元素，注意記憶體的連續分配
}

__device__ void SetElement(Matrix A, int row, int col, float value) {
    A.data[row * A.stride + col] = value; // 寫元素
}

__device__ Matrix GetSubMatrix(Matrix A, int row, int col) { // 獲取第r行第c列開始的子矩陣，注意這裡的r和c都是被子方陣縮放過的
    Matrix sub;

    sub.width = BLOCK_SIZE;
    sub.height = BLOCK_SIZE;
    sub.stride = A.stride;
    sub.data = &A.data[A.stride * BLOCK_SIZE * row + BLOCK_SIZE * col];

    return sub;
}

void MatMaul(const Matrix A, const Matrix B, Matrix C) {
    Matrix dA;

    dA.height = A.height;
    dA.width = A.width;
    dA.stride = A.width;

    size_t size = dA.width * dA.height * sizeof(float);

    cudaMalloc(&dA.data, size);
    cudaMemcpy(dA.data, A.data, size, cudaMemcpyHostToDevice);

    Matrix dB;

    dB.height = B.height;
    dB.width = B.width;
    dB.stride = B.width;

    size = dB.width * dB.height * sizeof(float);

    cudaMalloc(&dB.data, size);
    cudaMemcpy(dB.data, B.data, size, cudaMemcpyHostToDevice);

    Matrix dC;

    dC.height = C.height;
    dC.width = C.width;
    dC.stride = C.width;

    size = dC.width * dC.height * sizeof(float);

    cudaMalloc(&dC.data, size);

    dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
    dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);

    MatMulKernel<<<dimGrid, dimBlock>>>(dA, dB, dC);

    cudaMemcpy(C.data, dC.data, size, cudaMemcpyDeviceToHost);

    cudaFree(dA.data);
    cudaFree(dB.data);
    cudaFree(dC.data);
}

__global__ void MatMulKernel(const Matrix A, const Matrix B, Matrix C) {
    int blockRow = blockIdx.y;
    int blockCol = blockIdx.x;

    Matrix Csub = GetSubMatrix(C, blockRow, blockCol);

    float value = 0;
    int row = threadIdx.y; // 每個執行緒處理一個元素，所以執行緒的塊內橫縱id就是元素在子矩陣中的位置
    int col = threadIdx.x;

    for (int i = 0; i < A.width / BLOCK_SIZE; i++) { //遍歷A的子矩陣
        Matrix Asub = GetSubMatrix(A, blockRow, i);
        Matrix Bsub = GetSubMatrix(B, i, blockCol);

        __shared__ float As[BLOCK_SIZE][BLOCK_SIZE + 1]; // 子方陣就是一個共享記憶體
        __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE + 1];

        As[row][col] = GetElement(Asub, row, col);
        Bs[row][col] = GetElement(Bsub, row, col);

        __syncthreads(); // 對多塊共享記憶體IO時要注意同步

        for (int j = 0; j < BLOCK_SIZE; j++) {
            value += As[row][j] * Bs[j][col];
        }

        __syncthreads(); // 對多塊共享記憶體IO時要注意同步

    }

    SetElement(Csub, row, col, value);
}

通過這種塊式計算，A只用讀取B.width / block_size次全域性記憶體，B只用讀取A.height / block_size次，所以大大節省了全域性記憶體的開銷

頁鎖主機記憶體

執行時提供了使用了頁鎖主機記憶體(也稱為釘主機)主機記憶體的函式，這種記憶體和使用malloc()分配的可分頁主存是對立的：

    1、cudaHostAlloc()和cudaFreeHost()用來對頁鎖主機記憶體進行分配和回收；

    2、cudaHostRegister()將指定範圍內被malloc()函式分配的記憶體的分頁鎖住。

使用頁鎖記憶體有以下幾個好處：

    1、對於一些裝置，頁鎖主機記憶體和裝置記憶體之間的複製是可以和核函式併發執行的；

    2、在一些裝置上，頁鎖主機記憶體可以被對映到裝置上的地址空間中，直接消除了從資料在主機與裝置記憶體之間複製的必要；

    3、在有用前端匯流排的裝置上，通過把主機記憶體分配成頁鎖記憶體可以提高主機和裝置之間的頻寬，如果頁鎖記憶體再被分配成寫合併，那麼頻寬會更高。

但是，頁鎖記憶體是一種稀缺資源，所以它的分配會比可分頁記憶體的分配失敗的更早。另外，減少能夠讓作業系統分頁的記憶體數量、使用太多的頁鎖記憶體會降低系統的整體效能。

可移植記憶體

一塊頁鎖記憶體可以和系統中任何裝置共存，但是在預設情況下，頁鎖記憶體的好處只有和它被分配時所在的裝置(以及共享同一塊虛擬地址空間的裝置，如果有的話)共事時才能體現出上面的優勢。為了讓所有裝置都享有那些好處，需要在分配記憶體時給cudaHostAlloc()函式傳入cudaHostAllocPortable標誌，或者呼叫cudaHostRegister()函式時傳入cudaHostRegisterPortable標誌

寫合併記憶體

預設情況下被分配的頁鎖主機記憶體是可快取的，它可以通過給cudaHostAlloc()函式傳遞cudaHostAllocWriteCombined引數被改成寫合併記憶體。寫合併記憶體不會使用主機的L1和L2快取，從而讓應用別的部分有更多的快取使用。另外，當寫合併記憶體通過PCI傳輸匯流排傳遞資料時，他不會被監視，從而提高了大約40%的傳輸效能。但是從主機中的寫合併記憶體中讀取資料不可避免地緩慢，所以寫合併用到只被主機寫資料的記憶體上。

對映記憶體

一塊頁鎖記憶體也可以通過給cudaHostAlloc()函式傳遞標誌cudaHostAllocMapped或者給cudaHostRegister()函式傳遞標誌cudaHostRegisterMapped標誌被對映到裝置記憶體的記憶體空間中去，因此這種記憶體塊就有了兩個地址：通過cudaHostAlloc()或者malloc()函式返回的主機地址和可以通過cudaHostGetDevicePointer()函式返回的裝置地址，故而此記憶體可以在核函式中被直接訪問。只是當主機和裝置使用同一塊虛擬地址空間時，我們直接使用cudaHostAlloc()返回的地址即可，不必呼叫cudaHostGetDevicePointer()函式了。

從核函式內直接訪問主機記憶體有以下幾個好處：

    1、不需要在裝置中分配記憶體然後紮在主機和裝置的記憶體之間來回複製了，此時如果核函式需要資料，會發生隱式地資料傳輸；

    2、將資料傳輸和核函式執行重疊時不需要使用流了，面向核函式的資料傳輸會自動和核函式執行重疊。

但是，由於對映頁鎖記憶體在主機和裝置間是共享的，應用必須使用流或事件同步記憶體訪問，以便避免任何的寫後讀、讀後寫、寫後寫等風險。為了讀取任何對映頁鎖記憶體的裝置地址，頁鎖記憶體對映必須在其他任何cuda呼叫執行前通過呼叫cudaSetDeviceFlags()並傳遞cudaDeviceMapHost引數來使能，否則，cudaHostGetDevicePointer()函式會返回一個錯誤。當然，如果裝置壓根兒不支援對映頁鎖記憶體，cudaHostGetDevicePointer()函式也會返回錯誤。所以我們要通過查詢canMapHostMemory裝置屬性來檢查這項功能是否開啟（1表示支援）。

注意，在對映頁鎖記憶體上進行的原子函式從主機或裝置的角度來說，將不再是原子的。而且，cuda執行時要求對從裝置端初始化的記憶體進行的1位元組、2位元組、4位元組和8位元組的自然對齊讀寫要從主機或裝置的角度保留成單向訪問，在一些平臺上，記憶體的原子性可能被硬體分解成讀操作和寫操作，這些操作對自然對齊訪問有著同樣的保留要求(單向)。例如，cuda執行時就不會支援存在把8位元組的自然對齊寫分割成主機與裝置間的兩個4位元組寫(只能是一個讀一個寫)的傳遞橋的PCI傳遞匯流排。

非同步併發執行

cuda將以下操作暴露為可以和其他任務併發執行的獨立任務：

    1、主機上的計算；

    2、裝置上的計算；

    3、主機向裝置的記憶體遷移；

    4、裝置向主機的記憶體遷移；

    5、裝置內部的與裝置之間的記憶體遷移。

這些操作能達到併發度將取決於裝置的特徵集與計算能力

主機與裝置間的併發執行

主機併發執行是通過在裝置完成指定任務前就把控制權返回給主機執行緒的非同步庫函式實現的，使用非同步呼叫時，當合適的裝置資源可用是很多的裝置操作將被cuda驅動進行排隊以待執行。這樣就減輕了主機執行緒很多管理裝置的責任，讓它能夠執行別的任務。對於主機而言，下面的裝置操作是非同步的：

    核心啟動、從一個裝置記憶體的資料複製、從主機向裝置≤64KB的記憶體的複製、使用字首Async的函式進行的複製、記憶體設定函式的呼叫

程式設計師可以通過設定CUDA_LAUNCH_BLOCKING環境變數為1來關閉所有執行在系統上的cuda應用的核函式的非同步啟動，但這一特徵只對除錯行為開放，不應該用來讓產品軟體進行可信賴執行。

除非併發核分析開啟，當通過分析器來收集硬體計數器時，核函式將會同步啟動。如果記憶體複製涉及非頁鎖主機記憶體的話，也將是非同步的

併發核執行

一些計算能力>=2.X的裝置可以併發執行多個核函式，應用可以通過檢查concurrentKernels屬性來看看是否開啟了這項功能（1表示開啟）。一臺裝置可以併發執行的最大核啟動數量取決於它的計算能力，如下表所示

從cuda上下文啟動的核不能與另外一個從cuda上下文啟動的核併發執行，使用大量紋理或區域性記憶體的核函式也不太可能與其他核函式併發執行。

資料遷移與核執行重疊

一些裝置可以將GPU與主機之間的記憶體複製與核執行併發執行，應用可以通過asyncEngineCount裝置屬性(>0為支援)來檢查此功能是否開啟，如果使用這種複製方式涉及的主機記憶體必須是頁鎖的。當裝置使能concurrentKernels裝置屬性時可以將裝置內複製與核執行並行，如果使能asyncEngineCount時可以將裝置內複製與裝置外(向裝置複製與從裝置複製)複製並行，裝置內部複製可以通過將標準記憶體複製函式的源地址與目的地址都設成同一塊裝置來初始化

併發資料傳輸

一些計算能力≥2.X的裝置可以將裝置外複製並行執行，應用可以通過asyncEngineCount裝置屬性(2為支援)來檢查此功能是否開啟，如果使用這種複製方式涉及的主機記憶體必須是頁鎖的

流

應用通過流來管理上述的併發操作，流是一組有序執行的命令集合（可能由不同的主機執行緒指定），另一方面，不同的流可能不按彼此間的順序來執行或者併發執行，這種操作無從保證，因此不能指望它們的正確性（比如，核函式之間的通訊是未定義的）

1、建立與銷燬：

流通過建立一個流物件並把它指定為核啟動與主機-裝置間記憶體複製序列的引數來被定義，以下的程式碼建立了兩個流，並且在頁鎖記憶體中分配了一個float陣列hostPtr：

cudaStream_t stream[2];

for (int i = 0; i < 2; i++) {
    cudaStreamCreate(&stream[i]);
}

float* hostPtr;
int size = 16;

cudaMallocHost(&hostPtr, 2 * size * sizeof(float));

下面程式碼會把每個流建立成從主機向裝置的記憶體複製、核函式啟動與從裝置向主機的記憶體複製的工作流：

__global__ void MyKernel(float* dev0Ptr, float* dev1Ptr, int size) {

}

int main() {

    .....

    float* dev0Ptr;
    cudaMalloc(&dev0Ptr, 2 * size * sizeof(float));

    float* dev1Ptr;
    cudaMalloc(&dev1Ptr, 2 * size * sizeof(float));

    for (int i = 0; i < 2; i++) {
        cudaMemcpyAsync(&dev0Ptr[i], &hostPtr[i], size * sizeof(float), cudaMemcpyHostToDevice, stream[i]);
        MyKernel<<<100, 512, 0, stream[i]>>>(&dev0Ptr[i], &dev1Ptr[i], size); // <<<>>>內第三個引數為共享記憶體數
        cudaMemcpyAsync(&hostPtr[i], &dev1Ptr[i], size * sizeof(float), cudaMemcpyDeviceToHost, stream[i]);
    }

    return 0;
}

每個流從輸入陣列hostPtr中將自己的部分粗知道裝置記憶體中的dev0陣列中，通過呼叫MyKernel在裝置上處理dev0陣列，然後把結果dev1陣列複製到hostPtr的對應部分上。重疊行為描述了這個例子中流是怎樣根據裝置的計算能力進行重疊的，注意這裡使用的hostPtr必須指向一塊頁鎖主機記憶體，以便進行重疊執行。

可以通過呼叫cudaStreamDestroy()函式銷燬流：

cudaStreamDestroy(stream[i]);

考慮到當呼叫cudaStreamDestroy()時裝置可能還在做工作，這個函式會立刻返回，當流上的任務全部執行完後，流會自動釋放它的資源。

2、預設流：

不指定任何流引數或者設定流引數為0的核啟動和主機裝置間的記憶體複製會被分配到預設流上執行，從而會被序列有序執行。對於使用--default-stream per-thread編譯選項編譯的或者在包含cuda標頭檔案前定義了CUDA_API_PER_THREAD_DEFAULT_STREAM巨集的程式碼，預設流是一個常規流，每個主機執行緒都會自己的預設流。不指定--default-stream編譯引數值時，--default-stream預設為legacy。

  • 顯式同步：

    有幾個顯式地讓流進行彼此間同步的函式：

        1）、cudaDeviceSynchronize()：等待所有主機執行緒中所有流的所有前驅命令執行完畢；

        2）、cudaStreamSynchronize()：需要一個流作為引數，等待其中所有前驅命令完成。此函式可以被用來將主機和指定流進行同步，而允許別的流繼續在裝置上執行；

        3）、3cudaStreamWaitEvent()：需要一個流和一個時間作為引數，此函式被呼叫之後，此後所有被新增到這個流上的命令(呼叫前新增的不算)會延遲執行，直到指定的事件完成。這個流可以是0號預設流，這時呼叫cudaStreamWaitEvent()之後新增到所有流上的所有命令都要等待指定的事件完成；

        4）、cudaStreamQuery()：給應用提供了了解是否指定流上所有的前驅命令都已經完成的方式。

  為了避免不必要的降速，所有這些同步函式最好是為了計時或隔離失敗的啟動或者失敗的記憶體複製來被使用。

  • 隱式同步：

    當主機執行緒在不同的流之間執行下列操作中的任何一個，那麼來自這些不同的流的兩個命令將不能併發執行(也就是如果主機執行緒在流A上執行命令0，在流B上執行命令1，那麼命令0和命令1必須序列執行)：

        頁鎖主機記憶體分配；裝置記憶體分配與設定；同一塊記憶體裝置中兩個不同地址的記憶體複製；對NULL流的任何cuda命令；L1和共享記憶體配置間的切換

    對於支援併發核執行並且計算能力<=3.0的裝置來說，任何需要進行依賴檢查以檢視某個使用流的核啟動是否完成的操作，只有當cuda上下文中所有前置核啟動的所有執行緒塊開始執行之後才能執行，並且阻塞來自任何cuda上下文流的後置核啟動，直到被檢查的核啟動完成。

    需要做依賴性檢查的操作包含要檢查的核啟動流上的任何其他命令，以及在那個流上的任何cudaStreamQuery()呼叫。因此應用需要遵守以下指南來提供他們在併發核執行方面的潛力：

        1）、在依賴性操作前，應該執行所有的獨立操作；

        2）、任何型別的同步應該被儘可能延遲。

3、重疊行為：

兩個流之間重疊執行的數量取決於每個流執行的命令的順序，以及裝置是否支援資料遷移與核執行的併發、核併發執行以及併發資料傳輸。例如，在不支援併發資料傳輸的裝置上，建立與銷燬中的樣例程式碼根本不會重疊執行，因為流1發起的從主機向裝置的記憶體複製是在流0發起的從裝置向主機的記憶體複製之後進行的，所以只有流0的從裝置向主機的記憶體複製完成之後，流1的從主機向裝置的記憶體複製才能開始。如果以下面的方式重寫程式碼，並且假設裝置支援資料遷移與核執行的重疊的話，那麼流1的從主機向裝置的記憶體複製將和流0的核啟動重疊執行：

for (int i = 0; i < 2; i++) {
    cudaMemcpyAsync(&dev0Ptr[i], &hostPtr[i], size * sizeof(float), cudaMemcpyHostToDevice, stream[i]);

}

for (int i = 0; i < 2; i++) {
    MyKernel<<<100, 512, 0, stream[i]>>>(&dev0Ptr[i], &dev1Ptr[i], size);
}

for (int i = 0; i < 2; i++) {
    cudaMemcpyAsync(&hostPtr[i], &dev1Ptr[i], size * sizeof(float), cudaMemcpyDeviceToHost, stream[i]);
}

在支援併發資料傳輸的裝置上，建立與銷燬中的樣例程式碼會重疊執行：流1的從主機向裝置的記憶體複製與流0的從裝置向主機的記憶體複製甚至是流0的核啟動(假設裝置支援資料遷移與核執行的並行)將會重疊執行。但是，對於計算能力<=3.0的裝置，核的執行可能並不會重疊，因為流1的核啟動是在流0的裝置向主機複製記憶體之後進行的(更是在流0的核啟動之後)，並且兩個核函式都要訪問相同的裝置記憶體，所以它會向隱式同步中所描述的一樣，等待流0的核啟動完成之後再停止自己的阻塞。對於以上的重寫程式碼，假設裝置又支援併發的核啟動，那麼核的執行會重疊，因為流1的核啟動是在流0的從裝置向主機的記憶體複製完成之前進行的，但是在這種情況下，流0的裝置向主機的記憶體複製只會和流1核啟動的最後一個執行緒塊(佔核執行總時間的一小部分)重疊，如隱式同步中所述，這與重寫程式碼前流0和流1核啟動不重疊的情況一樣，因為流0的裝置向主機記憶體的複製與流1的核啟動都訪問了裝置記憶體1.

4、回撥：

執行時提供了通過cudaStreamAddCallback()函式在任何時間點向一個流中來插入回撥的方法，回撥是在插入回撥前所有被分配到流上的任務執行完之後，才在主機上執行的函式。0號流上的回撥會在插入回撥前的被分配到所有流上的所有前置任務執行完之後再執行。在下面的例子中，一個回撥函式MyCallback被分別插入到了兩個流中主機向裝置的記憶體複製、核函式啟動和裝置向主機記憶體的複製之後，因此當每個流的裝置向主機記憶體的複製完成之後，這個回撥函式就會各執行一次：

void CUDART_CB MyCallback(cudaStream_t stream, cudaError_t status, void *data){
    printf("Inside callback %d with status %d\n", (size_t)data, status);
}

....

int main() {

    .....

    int i;

    for (i = 0; i < 2; i++) {
        cudaMemcpyAsync(&dev0Ptr[i], &hostPtr[i], size * sizeof(float), cudaMemcpyHostToDevice, stream[i]);
        MyKernel<<<100, 512, 0, stream[i]>>>(&dev0Ptr[i], &dev1Ptr[i], size);
        cudaMemcpyAsync(&hostPtr[i], &dev1Ptr[i], size * sizeof(float), cudaMemcpyDeviceToHost, stream[i]);
        cudaStreamAddCallback(stream[i], MyCallback, (void*) i, 0);
        cudaStreamDestroy(stream[i]);

    }

   cudaDeviceSynchronize(); // 註冊回撥後，必須呼叫另一個cuda函式，否則會阻塞到回撥函式裡
   .....
   return 0;
}

被分配到一個流上的命令，或者回撥被分配到流0上時所有流上的所有命令都不會在回撥執行完之前開始執行，cudaStreamAddCallback()函式中的最後一個引數是以備後用的。回撥不能直接或間接地呼叫任何cuda api，因為如果它呼叫會導致死鎖的api時會自行中斷阻塞，也就是如果一個函式在等待回撥的執行，而這個回撥呼叫了會導致死鎖的api，那麼等它的這個函式將不會等回撥執行完就立刻執行。

5、流優先順序：

流的相對優先順序可以在建立時通過cudaStreamCreateWithPriority()函式指定，可用的優先順序範圍可以通過cudaDeviceGetStreamPriorityRange()函式以[最低, 最高]的形式獲取，在執行時，高優先順序的執行緒塊停止後，第優先順序流中的執行緒才能開始排程。下面程式碼就是獲取優先順序返回、指定流優先順序的示例

int priority_high, priority_low;
cudaDeviceGetStreamPriorityRange(&priority_low, &priority_high); // low為0，high為-2


cudaStream_t st_high, st_low;
cudaStreamCreateWithPriority(&st_high, cudaStreamNonBlocking, priority_high);
cudaStreamCreateWithPriority(&st_low, cudaStreamNonBlocking, priority_low);

6、圖：

圖表達了一種新的cuda任務提交模型，圖是由一系列類似核啟動這樣的操作用依賴相連(依賴的定義要獨立於執行)，而且允許圖的一次定義、重複執行。把圖的定義從執行中分離開來可以支援大量的優化：首先，CPU啟動圖的代價比流要小，因為大部分設定工作已經完成了；第二，向cuda提交整個工作流可以支援流的分段提交工作機制不支援的優化。為了看到這些只可能存在於圖中的優化，請考慮流中的這樣的場景：當我們把核函式放入流中，主機平臺會執行一系列操作以準備在GPU上執行核函式，這些對於設定與啟動核函式很有必要的操作就成了提交每個核時必須要付出的代價。對於執行時間較短的GPU核，這種代價可能是整個端到端耗時的主要部分。

使用圖的共提交可以被劃分為三個主要階段：定義、例項化、執行。

在定義階段，程式會在圖中建立對操作以及操作間依賴的描述；例項化會取圖模板的一個快照，驗證之，並且執行啟動所需最少的設定與初始化的大部分工作，例項化的結果又稱為可執行圖；一個可執行圖可以被啟動到一個流中，就像別的cuda任務一樣，可以一次例項化、多次啟動。

  • 圖的結構：

    操作對應一個結點，操作間的依賴對應邊，這些依賴限制了操作的執行序列。一旦依賴完成，結點對應的操作可以被隨時排程，排程就交給cuda系統。

    結點型別：進核函式、CPU函式呼叫、記憶體複製、memset、空、子圖（執行一個獨立的巢狀圖，如下圖中的Y所示）

  • 使用圖API建立圖：

    圖的建立可以通過兩種機制實現：顯式API和流捕獲，我們以建立並執行下面的圖為例：

    用圖API建立圖的程式碼如下：

cudaGraphCreate(&graph, 0); // 建立空圖

cudaGraphAddKernelNode(&a, graph, NULL, 0, &nodeParams); // 建立圖結點a、b、c、d
cudaGraphAddKernelNode(&b, graph, NULL, 0, &nodeParams);
cudaGraphAddKernelNode(&c, graph, NULL, 0, &nodeParams);
cudaGraphAddKernelNode(&d, graph, NULL, 0, &nodeParams); 

cudaGraphAddDependencies(graph, &a, &b, 1); // A->B，邊的建立也可以在建立結點時進行
cudaGraphAddDependencies(graph, &a, &c, 1); // A->C
cudaGraphAddDependencies(graph, &b, &d, 1); // B->D
cudaGraphAddDependencies(graph, &c, &d, 1); // C->D
  • 用流捕獲建立圖：

    流捕獲提供了一種從現有的基於流的api中建立圖的機制，可用cudaStreamBeginCapture()和cudaStreamEndCapture()包裹住將工作放入流中等現有程式碼，如下所示：

cudaGraph_t graph;

cudaStreamBeginCapture(stream);

kernel_A<<< ..., stream >>>(...);
kernel_B<<< ..., stream >>>(...);
libraryCall(stream);
kernel_C<<< ..., stream >>>(...);

cudaStreamEndCapture(stream, &graph);

    cudaStreamBeginCapture()的呼叫將流放入捕獲模式下，當流被捕獲時，寫入流中的工作不會去排隊等著執行，而是被追加到一個正在被建立的中間圖中，這個圖會通過cudaStreamEndCapture()的呼叫被返回(&graph)，同時也會結束流的捕獲模式。通過流捕獲正在被積極構建的圖稱之為捕獲圖。流捕獲可以在任何cuda流(cudaStreamLegacy除外，又稱為null流)上使用，也可以在cudaStreamPerThread上使用。如果程式正在使用legacy流，它可以把0號流重新定義成per-thread流，同時不改變自己的功能性，請參見預設流部分

   1）、跨流的依賴和事件：

    流捕獲可以處理用cudaEventRecord()和cudaStreamWaitEvent()傳遞的跨流依賴，只要被等待的事件被記錄到了同一張捕獲圖中。當一個事件被在捕獲模式下的流記錄的時候，它就成了捕獲事件，表示捕獲圖中的一系列結點。

    當一個捕獲事件被流所等待，它就會把不處於捕獲模式下的等它的流放入捕獲模式下，而且流中下一項將會對捕獲事件中的結點產生額外的依賴，而後這兩條流就都會捕獲到一張捕獲圖中。

    當跨流依賴在捕獲流中存在時，cudaStreamEndCapture()還是必須被呼叫cudaStreamStartCature()的流呼叫，這叫原始流。由於事件依賴的關係，任何其他被捕獲到同一張捕獲圖的流必須加入到原始流的後面。而且，同一張捕獲圖裡的所有捕獲流當呼叫cudaStreamEndCapture()時會退出捕獲模式，如果加入原始流失敗，整個捕獲操作就會隨之失敗

cudaStreamBeginCapture(stream1); // 流1是原始流

kernel_A<<< ..., stream1 >>>(...);

cudaEventRecord(event1, stream1); // 事件1是捕獲事件
cudaStreamWaitEvent(stream2, event1); // 流2加入原始流

kernel_B<<< ..., stream1 >>>(...); // 流1執行B
kernel_C<<< ..., stream2 >>>(...); // 流2執行C

cudaEventRecord(event2, stream2); // 流2記錄事件2
cudaStreamWaitEvent(stream1, event2); // 流1等待事件2，事件發生後流2併入流1中

kernel_D<<< ..., stream1 >>>(...); // 流1執行D

cudaStreamEndCapture(stream1, &graph); // 流1和流2退出捕獲模式

    得到的圖如下所示：

    2）、相關的無效操作：

    查詢正在被捕獲的流或事件的執行狀態或者對其同步是無效的，因為他們沒有展現可以被排程執行的項。對包含任何處於捕獲模式下的流的更寬控制程式碼(比如裝置、上下文等)進行的同步和狀態查詢也是無效的。當上下文中存在任何沒有使用cudaStreamNoBlocking建立的流正在被捕獲時，嘗試使用null流也是無效的，這是因為null流總是會包含其他流的引用，加入null流會建立對正在被捕獲的流的依賴，對null流的查詢或同步也會作用到被捕獲的流上。所以在這種情況下使用同步API也是無效的，類似cudaMemcpy()這樣的同步API會把任務加入到null流中，並且同步等待api的返回。

    如果一個捕獲事件來自於一條捕獲流，並且這條流與另外一張捕獲圖相關聯，也就是這個捕獲事件處一條位於兩張捕獲圖交集的捕獲流上，那麼通過這個事件來進行兩張捕獲圖的合併，那是無效的。同樣無效的還有在捕獲流上等待非捕獲事件。

    把非同步操作入隊到流中的少量API(例如cudaStreamAttachAsync())目前不支援圖，如果被捕獲流呼叫將會返回錯誤。

    3）、失效：

    當以上一個失效操作在流捕獲期間被嘗試執行，相關的任何捕獲圖就會隨之失效。當一個捕獲圖無效後，對於相關捕獲流和捕獲事件的後續使用將會無效並且返回一個錯誤，直到流捕獲被cudaStreamEndCapture()中斷，這個函式會把相關的流帶出捕獲模式，但也會返回一個錯誤和一個NULL圖指標。

  • 圖API的使用：

    cudaGraph_t物件不是執行緒安全的，使用者應該保證多個執行緒不會併發訪問一個graph_t物件；cudaGraphExec_t不能和自己併發執行，一個cudaGraphExec_t的啟動將在啟動同一個可執行圖之後進行；和其他非同步任務排序時，圖的執行將在流中進行，但是這個流只是用來排序的，它不會限制圖的內部並行度，也不會影響圖結點在哪條流上執行

7、事件：

執行時也提供了近距離監視裝置程式和執行準確計時的方法，那就是讓應用在程式的任何部分非同步記錄時間，然後當事件完成後進行查詢。當流上的所有前置任務或者命令完成後，事件也就完成了。在0號流上的事件則在所有流上的所有前置任務完成後才會完成

  • 建立與銷燬：

    事件的建立與銷燬的程式碼如下所示：

cudaEvent_t start, stop;

cudaEventCreate(&start);
cudaEventCreate(&stop);

cudaEventDestroy(start);
cudaEventDestroy(stop);
  • 執行時間：

    上面例子中建立的事件可以用下面的方式統計執行時間

cudaEventRecord(start, 0);

for (int i = 0; i < 2; ++i) {
    cudaMemcpyAsync(&dev0Ptr[i * size], &hostPtr[i * size],
                    size * sizeof(float ), cudaMemcpyHostToDevice, stream[i]);
    MyKernel<<<1, 1, 0, stream[i]>>>
            (dev0Ptr, dev1Ptr, i * size, size);
    cudaMemcpyAsync(&hostPtr[i * size], &dev1Ptr[i * size],
                    size * sizeof(float ), cudaMemcpyDeviceToHost, stream[i]);
}

cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);

float elapsedTime;
cudaEventElapsedTime(&elapsedTime, start, stop);

std::cout << "time: " << elapsedTime << std::endl; // time: 1.1407 單位為毫秒

8、同步呼叫：

當一個同步函式呼叫時，裝置完成指定任務之前是不會把控制權返回給主機執行緒的。在主機執行緒執行任何cuda呼叫前，可以通過呼叫cudaSetDeviceFlags()並傳入具體的引數來開啟主機執行緒在這種情況下是讓出控制權(讓別的主機執行緒搶佔CPU)、阻塞系統還是自己繼續執行。

多裝置系統

裝置列舉：

 一個主機系統可以有多個裝置，下面的樣例程式碼展示瞭如何列舉這些裝置、查詢他們的屬性，並確定支援cuda的裝置：

int deviceCount;

cudaGetDeviceCount(&deviceCount);

int device;

for (device = 0; device < deviceCount; ++device) {
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, device);
    printf("Device %d has compute capability %d.%d.\n",device, deviceProp.major, deviceProp.minor);
    // Device 0 has compute capability 6.1.
}

裝置選擇

主機執行緒隨時可以通過呼叫cudaSetDevice()來選擇它要操作的裝置，裝置記憶體分配和核心啟動、流和事件的建立將在當前選擇的裝置上進行。如果沒有呼叫cudaSetDevice()，那麼當前裝置就是裝置0。

下面的樣例程式碼展示了設定當前裝置會怎樣影響記憶體分配與核執行：

size_t size = 1024 * sizeof(float);

cudaSetDevice(0); // 設定裝置0為當前裝置
float* p0;
cudaMalloc(&p0, size); // 在裝置0上分配記憶體
MyKernel<<<1000, 128>>>(p0); // 在裝置0上執行核函式

cudaSetDevice(1); // 設定裝置1位當前裝置
float* p1;
cudaMalloc(&p1, size); // 在裝置1上分配記憶體
MyKernel<<<1000, 128>>>(p1); // 在裝置1上執行核函式

流與事件行為

如果核函式被分配到了不屬於當前裝置上的流時，它就會啟動失敗，如下所示

cudaSetDevice(0); 
cudaStream_t s0;
cudaStreamCreate(&s0); // s0在裝置0上
MyKernel<<<100, 64, 0, s0>>>(); 

cudaSetDevice(1);
cudaStream_t s1;
cudaStreamCreate(&s1); // s1在裝置1上
MyKernel<<<100, 64, 0, s1>>>();

MyKernel<<<100, 64, 0, s0>>>();  // 當前裝置為裝置1，在s0上啟動核函式會失敗

但是，即便把記憶體分配操作分配到不在當前裝置上的流時，它依舊會成功。

如果輸入的流和輸入事件不在同一臺裝置上，cudaEventRecord()會失敗；如果兩個輸入的事件不再同一臺裝置上，cudaEventElapsedTime()會失敗；如果輸入的事件不屬於當前裝置，cudaEventSynchronize()和cudaEventQuery()函式依舊成功；如果輸入流和輸入事件不屬於同一臺裝置，cudaStreamWaitEvent()還是會成功，因此此函式可以被用來多臺裝置間的同步。

每臺裝置都有自己的預設流，因此被分配到不同裝置上預設流的命令之間會併發執行。

對等結點記憶體訪問

當應用執行在Tesla系列、計算能力>=2.0且為64位的裝置上時，它可以對不同裝置的記憶體空間進行取址（例如，在某臺裝置上執行的核函式可以析構另一臺裝置上的記憶體指標）。如果兩臺裝置的cudaDeviceCanAccessPeer()都返回true的話，這種對等結點記憶體訪問特徵將在這兩臺裝置上受到支援。

必須通過呼叫cudaDeviceEnablePeerAccess()函式來在開啟與目標裝置的對等結點記憶體訪問，如下面程式碼所示：

cudaSetDevice(0); // 設定裝置0為當前裝置
float* p0;
size_t size = 1024 * sizeof(float);

cudaMalloc(&p0, size); // 在裝置0上分配記憶體
MyKernel<<<1000, 128>>>(p0); // 啟動核函式

cudaSetDevice(1); // 設定裝置1為當前裝置
cudaDeviceEnablePeerAccess(0, 0); // 開啟與裝置1的對等結點記憶體訪問

MyKernel<<<1000, 128>>>(p0); // 此時可以在裝置1上位於裝置0記憶體中的p0了

在沒有支援NVSwitch的系統上，每個裝置可以最大支援8個系統級別的對等連線。如果兩臺裝置上都使用了統一虛地址空間，那麼相同的指標就可以被用來指向兩臺裝置上的記憶體地址。

對等結點記憶體複製

兩臺裝置間也可以進行記憶體複製。當兩臺裝置都使用統一地址空間時，對等記憶體複製可以通過常規記憶體複製來完成；否則的話，我們就需要使用cudaMemcpyPeer()、cudaMemcpyPeerAsync()、cudaMemcpy3DPeer()或cudaMemcpy3DPeerAsync()來進行了，如下程式碼所示：

cudaSetDevice(0); // 設定裝置0為當前裝置
float* p0;
size_t size = 1024 * sizeof(float);
cudaMalloc(&p0, size); // 在裝置0上分配記憶體

cudaSetDevice(1); // 設定裝置1為當前裝置
float* p1;
cudaMalloc(&p1, size); // 在裝置1上分配記憶體

cudaSetDevice(0); // 設定裝置0為當前裝置
MyKernel<<<1000, 128>>>(p0); // 在裝置0上啟動核函式

cudaSetDevice(1); // 設定裝置1為當前裝置

cudaMemcpyPeer(p1, 1, p0, 0, size); // 從p0記憶體複製到p1

MyKernel<<<1000, 128>>>(p1); // 在裝置1上啟動核函式

預設在null流上進行的不同裝置間的記憶體複製要等到兩臺裝置上所有的前置命令完成之前才會開始，而且兩臺裝置上任何在其之後的命令都要等待記憶體複製完成之後才能開始。

和流的普通行為一致，兩臺裝置間的記憶體非同步複製可以和其他流上的核函式或複製重疊執行。注意如果兩臺裝置間的對等記憶體訪問通過cudaDeviceEnablePeerAccess()開啟了，這兩臺裝置間的記憶體訪問就不需要通過主機進行了，從而更快。

統一虛地址空間

當應用以一個64位程式執行時，主機和所有計算能力>=2.0裝置將使用一塊地址空間。通過cuda api進行的所有主機記憶體訪問和裝置記憶體訪問都會在這個虛地址空間中進行，因此：

    1）、通過cuda分配的主機記憶體位置或使用統一記憶體空間的裝置記憶體位置都可以通過cudaPointerGetAttributes()函式由指標的值得到；

    2）、當和任何使用統一記憶體空間的裝置進行記憶體複製時，cudaMemcpyXXX()函式的cudaMemcpyKind引數可以設定成cudaMemcpyDefault來通過指標確定位置，即便這個指標時沒有使用cuda分配的也沒關係，只要要互動的裝置使用了統一地址即可；

    3）、通過cudaHostAlloc()函式進行的分配是自動可以在所有使用統一記憶體地址裝置間進行移植的，而且此函式返回的指標可以直接在這些裝置上的核函式裡使用(不需要向在對映記憶體中記載的那樣，通過cudaHostGetDevicePoiner()來獲得裝置指標了)

應用可以通過檢查unifiedAddressing引數(1為使用)來檢視某個裝置是否使用了統一記憶體空間，檢查方法參見多裝置系統中的裝置列舉一節

程式間通訊

被主機執行緒建立的任何裝置記憶體指標或事件控制程式碼可以被同一程式中的其他任何執行緒引用，但是出了這個程式就不行了，也就是不能被別的程式中的執行緒引用。

為了在程式間共享裝置記憶體指標和事件，應用就必須使用程式間通訊的API了，這在參考手冊中有詳細的介紹。IPC API只支援執行計算能力>=2.0的裝置上的64位Linux程式，但是不支援通過cudaMallocManaged()分配的記憶體。

使用這些IPC API，應用可以使用cudaIpcGetMemHandle()來獲取一個給定裝置記憶體指標的IPC控制程式碼，使用標準的IPC機制(共享記憶體或檔案等)來把它傳給另一個程式，再使用cudaIpcOpenMemHandle()來從IPC控制程式碼中獲取其他程式中有效的指標。事件共享的方法類似。

使用IPC API的一個例子就是當一個主程式生成一批輸入資料，可以通過IPC在不進行任何重新生成或者複製的情況下讓這些資料對別的子程式可用。

錯誤檢查

所有的執行時函式都會返回一個錯誤碼，但是對於非同步函式，這個錯誤碼可能不會報告發生在裝置上的錯誤，因為函式在裝置完成任務之前就返回了；錯誤碼只會在任務執行前給主機報告錯誤，主要是和引數驗證相關的錯誤；如果非同步錯誤發生，它將會通過一些下游不相關的執行時函式呼叫來進行報告。

在一些非同步函式呼叫之後檢查非同步錯誤的唯一方法因此就是同步，通過呼叫cudaDeviceSynchronize()來進行，或者使用非同步併發執行一節中提及的其他同步機制，然後檢查cudaDeviceSynchronize()等函式返回的錯誤碼。

執行時為每一個主機執行緒都分配了一個初始值為cudaSuccess的錯誤變數，它的值可以在每個錯誤發生時被修改(要麼是非同步錯誤，要麼是一個引數驗證錯誤)。cudaPeekAtLastError()函式會返回這個錯誤，cudaGetLastError()也會返回這個錯誤，但同時會把錯誤變數重置為cudaSuccess。

核函式啟動不會返回任何錯誤碼，因此cudaPeekAtLastError()或cudaGetLastError()必須在核函式啟動後立刻呼叫來獲取任何的預啟動錯誤。為了保證由這兩個函式返回的錯誤不是來源於核函式啟動之前的某個操作，我們要確保執行時錯誤變數在核函式啟動前被設定成了cudaSuccess，比如在核函式之前呼叫cudaGetLastError()。核函式啟動是非同步的，所以必須在核函式啟動和其之後的cudaPeekAtLastError()或cudaGetLastError()呼叫之間進行同步，以檢查這種非同步錯誤。

注意，cudaStreamQuery()和cudaEventQuery()返回的cudaErrorNotReady錯誤不會被當成一個錯誤，因此不會被cudaPeekAtLastError()或cudaGetLastError()返回。

呼叫棧

在計算能力大於等於2.X的裝置上，呼叫棧的大小可以通過cudaDeviceGetLimit()來得到，也可以通過cudaDeviceSetLimit()來設定。當呼叫棧溢位時，核函式呼叫就會失敗，並且產生一個棧溢位錯誤(如果使用cuda-gdb、Nsight等cuda除錯工具執行應用的話)或者未指定的的啟動錯誤。

紋理記憶體和表面(surface)記憶體

cuda支援一種GPU可以用來為影像訪問紋理和表面記憶體的紋理硬體子集，從紋理或表面記憶體而不是全域性記憶體中讀取資料可以有幾個效能上的好處，這在裝置記憶體訪問中會提到。

兩種API可以用來訪問紋理和表面記憶體：所有裝置都支援的紋理引用API和只在計算能力為3.x裝置上支援的紋理物件API，前者有更多的限制

紋理物件或紋理引用會指定以下屬性

屬性

描述

紋理

指定哪塊紋理要被獲取。紋理物件在執行時建立，而紋理在紋理物件建立時被建立；紋理引用在編譯期被建立，但紋理是在紋理引用執行時通過執行時函式被繫結到紋理時被建立的。一些特殊的紋理引用可能被繫結到同一塊紋理或者在記憶體中重疊的紋理上。紋理可以是線性記憶體中的任何部分或者是一個cuda陣列。

維度

維度指明紋理是用一個（兩個、三個）紋理座標表示的一維（二維、三維）陣列，陣列中的元素稱為texels(紋素)，是紋理元素texture elements的簡稱。紋理寬度、高度和深度表示陣列在每個維度上的尺寸。

型別

紋理元素的型別受限於基本整型、單精度浮點型別和在char、short、int、long、longlong、float、double中定義的一維、二維和四維向量型別，這些向量定義也來源於整型和單精度浮點型別

讀取模式

可取cudaReadModeNormalizedFloat或cudaReadModeElementType。如果是前者，並且元素型別是2位元組或單位元組整型，被紋理獲取返回的值將是浮點數型別，或者全範圍整型，但是值域會被對映到[0.0, 1.0]和[-1.0, 1.0]之間，前者針對無符號，後者針對有符號，例如，一個值為0xff的無符號單位元組紋理元素會被讀成1；如果是後者，就不會有任何轉換

座標是否被正規化

預設情況下，紋理是被定義域為[0, N - 1]的浮點座標引用的，其中N為與座標相關的紋理維度最大值。例如，尺寸為64 * 32的紋理的座標範圍就是([0, 63], [0, 31])。正規化的座標會導致座標被對映到[0.0, 1.0 - 1 / N]而不是[0, N - 1]，所以同一個64 * 32的紋理會被範圍為([0.0, 1.0 - 1 / 64], [0.0, 1.0 - 1 / 31])的正規化座標引用。正規化座標可以天然適應某些應用的要求，特別是要求紋理座標獨立於紋理大小的情況下

取址模式

當座標超出定義域時，呼叫紋理函式依舊是有效的，其結果取決於取址模式。預設的取址模式是把座標固定到定義域中：非正規化為[0, N)，正規化為[0.0, 1.0)。如果邊界模式被指定，讀取座標超出定義域的紋理會返回0。對於正規化的座標，我們還可以使用包裹模式和映象模式。當使用包裹模式時，每個座標x會被轉換成x * floor(x)，floor(x)表示不超過x的最大整數；當使用映象模式時，座標x會根據floor(x)的奇偶性進行轉換：若為偶數，則為x * floor(x)，否則就是 1 - x * floor(x)。取址模式通過一個三維陣列表示，每個元素表示每個紋理座標維度的取址模式，可用的有cudaAddressModeBorder、cudaAddressModeClamp、cudaAddressModeWrap和cudaAddressModeMirror，後兩個只支援正規化的紋理座標

過濾模式

過濾模式指定讀取紋理時返回值是怎麼基於輸入的紋理座標計算的。線性紋理過濾只適用於配置為返回浮點型別資料的紋理，它在相鄰紋理元素間執行低精度的插值。當使用線性紋理過濾時，首先會讀取紋理獲取座標附近的紋理元素，然後會紋理獲取的返回值會根據落在這些紋理元素之間的座標進行插值運算得到。

一維紋理會執行線性插值；二維紋理會執行雙線性插值；三維紋理會進行三線性插值。紋理獲取一節會更進行詳細的介紹。

過濾模式可以取cudaFilterModePoint或cudaFilterModeLinear。如果是前者，返回值將是座標最接近輸入紋理座標的紋理元素；如果是後者，返回值將是一維(二維、三維)紋理的2（4、8）個座標最接近輸入紋理座標的線性插值結果，但是這種模式只支援浮點型別的返回值。

下表列舉了不同計算能力的裝置支援的最大紋理寬度、高度和深度：

關於char、short、int、long、longlong、float、double中定義的一維、二維和四維向量的的定義如下：來源於整型和浮點型別的向量型別是結構體，其第1、2、3、4個元素可以通過欄位x、y、z和w得到。這些向量都通過形如make_<type_name>的建構函式得到，比如int2 make_int2(int x, int y);會建立int2型別的向量，帶有欄位(x, y)。

紋理物件API

紋理物件通過cudaResourceDesc型別結構體的資源描述符的cudaCreateTextureObject()方法建立，這種結構體指定了紋理的屬性，其內容如下：

struct cudaTextureDesc
{
    enum cudaTextureAddressMode addressMode[3];   // 取址模式
    enum cudaTextureFilterMode  filterMode;       // 過濾模式
    enum cudaTextureReadMode    readMode;         // 讀取模式
    int                         sRGB;
    int                         normalizedCoords; // 是否正規化座標
    unsigned int                maxAnisotropy;
    enum cudaTextureFilterMode  mipmapFilterMode;
    float                       mipmapLevelBias;
    float                       minMipmapLevelClamp;
    float                       maxMipmapLevelClamp;
}; // sRGB、maxAnisotropy等欄位請參考手冊

下面的程式碼把簡單的轉換核函式應用到了紋理上：

#include <stdio.h>
#include <cuda_runtime.h>

__global__ void transformKernel(float *output, cudaTextureObject_t texObj, int width, int height, float theta) {
    unsigned int x = blockIdx.x * blockIdx.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockIdx.y + threadIdx.y;

    float u = x / (float) width - 0.5f;
    float v = y / (float) height - 0.5f;

    float tu = u * cosf(theta) - v * sinf(theta) + 0.5f; // cuda取樣時會偏移5畫素，因此要偏移回去
    float tv = v * cosf(theta) + u * sinf(theta) + 0.5f;

    output[y * width + x] = tex2D<float>(texObj, tu, tv);
}

int main() {
    cudaChannelFormatDesc desc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat); // 建立cuda陣列需要用到管道，這裡採用float型別的，向量為(32, 0, 0, 0)，也就是隻有x方向上有4個位元組
    cudaArray *cudaArray;
    int width = 5, height = 2;
    int size = width * height;
    float *h_data = (float *) malloc(size * sizeof(float));

    for (int i = 0; i < size; i++) { // 初始化資料
        h_data[i] = i;
    }

    printf("data initialized.\n");

    cudaMallocArray(&cudaArray, &desc, width, height); // 分配cuda陣列
    cudaMemcpyToArray(cudaArray, 0, 0, h_data, size, cudaMemcpyHostToDevice); // 從主機複製陣列給cuda陣列

    struct cudaResourceDesc resDesc; // 資源描述符定義

    memset(&resDesc, 0, sizeof(resDesc));
    resDesc.resType = cudaResourceTypeArray;
    resDesc.res.array.array = cudaArray;

    struct cudaTextureDesc texDesc; // 紋理描述符定義

    memset(&texDesc, 0, sizeof(texDesc));
    texDesc.addressMode[0] = cudaAddressModeWrap;
    texDesc.addressMode[1] = cudaAddressModeWrap;
    texDesc.filterMode = cudaFilterModeLinear;
    texDesc.readMode = cudaReadModeElementType;
    texDesc.normalizedCoords = 1;

    cudaTextureObject_t texObj = 0; // 紋理物件建立
    cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL);

    printf("Texture object created.\n");

    float* output;

    cudaMalloc(&output, size * sizeof(float ));

    dim3 dimBlock(16, 16);
    dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y - 1) / dimBlock.y);

    printf("Ready to call kernel.\n");

    transformKernel<<<dimGrid, dimBlock>>>(output, texObj, width, height, 0.5f); // 呼叫核函式

    float *h_result = (float *) malloc(size * sizeof(float));

    cudaMemcpy(h_result, output, size * sizeof(float), cudaMemcpyDeviceToHost); // 複製結果到主機

    printf("Kernel is finished and here`s the result:\n");

    for (int i = 0; i < size; i++) {
        printf("%f\n", h_result[i]);
    }

    cudaDestroyTextureObject(texObj); // 釋放資源
    cudaFreeArray(cudaArray);
    cudaFree(output);

    return 0;
}

輸出如下：

紋理引用API

紋理引用的一些屬性是不可改變的，並且必須在編譯期確定，因此必須在宣告紋理引用時指定。一個此檔案可用的紋理引用可以被定義成紋理型別的變數：

texture<DataType, Type, ReadMode> texRef;

其中，DataType指定紋理元素的型別；Type指定紋理引用的型別(cudaTextureType1D、cudaTextureType2D、cudaTextureType3D、cudaTextureType1DLayered和cudaTextureType2DLayered，分別對應一維、二維、三維、分層一維和分層二維的紋理)，預設值為cudaTextureType1D；ReadMode指定了讀取模式，預設值為cudaReadModeElementType。紋理引用只能被宣告為靜態全域性變數，不能作為引數傳給函式。

紋理引用的其他屬性是可改變的，其值可以在主機執行時被修改。根據參考手冊中的記載，執行時API有一個低階的C風格介面和一個高階的C++風格介面。而紋理型別是作為一種公有結構體在高階API中被定義，而這種公有結構體來源於低階API中的textureReference型別，其定義如下：

struct __device_builtin__ textureReference
{
    int                          normalized;        // 紋理座標是否正規化，可在主機程式碼中直接修改
    enum cudaTextureFilterMode   filterMode;        // 過濾模式，可在主機程式碼中直接修改
    enum cudaTextureAddressMode  addressMode[3];    // 取址模式，可在主機程式碼中直接修改
    struct cudaChannelFormatDesc channelDesc;       // 紋理元素格式，必須和紋理引用宣告時的DataType引數匹配。這個欄位的型別如下：
/*

struct __device_builtin__ cudaChannelFormatDesc
{
    int                        x; // x、y、z、w對應四維向量型別的每一維的位元組數量，請參見上文紋理記憶體的屬性表部分
    int                        y;
    int                        z; 
    int                        w;
    enum cudaChannelFormatKind f; // cudaChannelFormatKindSigned(紋理元素為有符號整型)、cudaChannelFormatKindUnsigned(紋理元素為無符號整型)、cudaChannelForamtKindFloat(紋理元素為浮點型)
};
*/
    int                          sRGB;
    unsigned int                 maxAnisotropy;
    enum cudaTextureFilterMode   mipmapFilterMode;
    float                        mipmapLevelBias;
    float                        minMipmapLevelClamp; 
    float                        maxMipmapLevelClamp;
    int                          __cudaReserved[15];
};

在核函式可以使用紋理引用來從紋理記憶體中讀取資料之前，紋理引用必須通過cudaBindTexture()或cudaBindTexture2D()來繫結一塊線性記憶體，或者使用cudaBindTextureToArray()來繫結一個cuda陣列。cudaUnbindTexture()用來為紋理引用解綁，紋理引用解綁後可以安全地繫結到新的陣列上，即便使用老的繫結紋理的核函式還沒有完成。我們建議使用cudaMallocPitch()線上性空間中分配二維紋理，然後把此函式返回的物件作為引數傳遞給cudaBindTexture2D()。下面的程式碼簡單地把一個二維紋理引用繫結到了devPtr指向的線性記憶體中：

  • 低階API：
texture<float, cudaTextureType2D, cudaReadModeElementType> texRef;
const textureReference *texRefPtr;

cudaGetTextureReference(&texRefPtr, &texRef);
cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();

float *devPtr;
size_t pitch;
cudaMallocPitch((void **) &devPtr, &pitch, width * sizeof(float), height);

size_t offset;
cudaBindTexture2D(&offset, texRefPtr, devPtr, &desc, width, height, pitch);
  • 高階API：
texture<float, cudaTextureType2D, cudaReadModeElementType> texRef;
cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();

float *devPtr;
size_t pitch;
cudaMallocPitch((void **) &devPtr, &pitch, width * sizeof(float), height);

size_t offset;
cudaBindTexture2D(&offset, texRef, devPtr, desc, width, height, pitch);

下面的例子把一個二維紋理引用繫結到了cuda陣列cuArray上：

  • 低階API：
texture<float, cudaTextureType2D, cudaReadModeElementType> texRef;
const textureReference* texRefPtr;

cudaGetTextureReference(&texRefPtr, &texRef);

cudaChannelFormatDesc desc;

cudaArray *cudaArray;
cudaMallocArray(&cudaArray, &desc, width, height);

cudaGetChannelDesc(&desc, cudaArray);

cudaBindTextureToArray(texRef, cudaArray);
  • 高階API：
cudaChannelFormatDesc desc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);

cudaArray *cudaArray;

cudaMallocArray(&cudaArray, &desc, width, height);

texture<float, cudaTextureType2D, cudaReadModeElementType> texRef;

cudaBindTextureToArray(texRef, cudaArray);

把紋理繫結到紋理引用時指定的格式必須和紋理引用宣告時的DataType引數匹配，否則紋理獲取的結果就不得而知。每個核函式可以繫結的紋理數量如下表所示

下面的程式碼同樣是使用核函式對紋理進行簡單的轉換：

#include "cuda_runtime.h"
#include "texture_fetch_functions.h"
#include <stdio.h>

void textureObjTest();
void bindTextureRef();

int width = 5, height = 2;
int size = width * height;
float *h_data = (float *) malloc(size * sizeof(float));
texture<float, cudaTextureType2D, cudaReadModeElementType> texRef;

__global__ void transformKernelRef(float *output, int width, int height, float theta) {
    unsigned int x = blockIdx.x * blockIdx.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockIdx.y + threadIdx.y;

    float u = x / (float) width - 0.5f;
    float v = y / (float) height - 0.5f;

    float tu = u * cosf(theta) - v * sinf(theta) + 0.5f;
    float tv = v * cosf(theta) + u * sinf(theta) + 0.5f;

    output[y * width + x] = tex2D(texRef, tu, tv);
}

int main() {
    bindTextureRef();
    return 0;
}

void bindTextureRef() {
    cudaChannelFormatDesc desc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);

    for (int i = 0; i < size; i++) {
        h_data[i] = i;
    }

    cudaArray *cudaArray;
    cudaMallocArray(&cudaArray, &desc, width, height);
    cudaMemcpyToArray(cudaArray, 0, 0, h_data, size, cudaMemcpyHostToDevice);

    texRef.addressMode[0] = cudaAddressModeWrap;
    texRef.addressMode[1] = cudaAddressModeWrap;
    texRef.filterMode = cudaFilterModeLinear;
    texRef.normalized = 1;

    cudaBindTextureToArray(texRef, cudaArray);

    float *output;
    cudaMalloc(&output, size * sizeof(float));

    dim3 dimBlock(16, 16);
    dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y - 1) / dimBlock.y);

    printf("Ready to call kernel.\n");
    transformKernelRef<<<dimGrid, dimBlock>>>(output, width, height, 0.5f);

    float *h_result = (float *) malloc(size * sizeof(float));
    cudaMemcpy(h_result, output, size * sizeof(float), cudaMemcpyDeviceToHost);

    printf("Kernel is finished and here`s the result:\n");
    for (int i = 0; i < size; i++) {
        printf("%f\n", h_result[i]);
    }

    cudaFreeArray(cudaArray);
    cudaFree(output);
}

輸出結果和使用紋理物件的一樣

兩位元組浮點紋理

cuda陣列支援的兩位元組浮點或半(half)格式和IEEE 754-2008的binary2format格式一樣。cuda c不支援匹配資料型別，但是提供了指令級函式來通過無符號短整型進行四位元組浮點型別的紋理和兩位元組紋理之間的轉換，這些函式有__float2half_rn(float)和__half2float(unsigned short)。這些函式只能在裝置程式碼中使用，主機程式碼中的等價函式可以在OpenEXR庫中找到。

在執行紋理獲取的過濾前，兩位元組浮點元素應該被轉換成四位元組浮點數。另外，可以通過cudaCreateChannelDescHalfXX()系列函式來建立兩位元組浮點格式的通道描述符。

分層紋理

一維二維分層紋理（在Direct3D中也稱為紋理陣列，OpenGL中則是陣列紋理）是由一系列層組成的紋理，所有層的紋理都是維度、大小和資料型別一致的常規紋理，只在計算能力>=2.0的裝置上支援。一維的分層紋理通過一個整數索引和一維浮點數紋理座標來取址，前者表示層數，後者表示層內座標；二維分層紋理通過一個整數索引和二維浮點數座標取址，前者表示層數，後者表示層內座標。

分層紋理只能是一個cuda陣列，通過給cudaMalloc3DArray()函式傳遞cudaArrayLayered標誌建立，如果是一維分層紋理的話，高度引數為0。在裝置函式中獲取分層紋理可以參見官方手冊中tex1DLayered()和tex2DLayered()函式，紋理過濾(參見官方手冊紋理獲取)只能在層內進行，而不能跨層進行。

立方圖紋理

立方圖紋理是一種特殊的二維分層紋理，這種紋理有六層，分別表示立方圖的每個面，而層的寬度等於其高度；

立方圖使用三個紋理座標x、y、z來取址，這三個座標可以被解釋為以立方體中心為原點，指向立方體某一面和和那個面對應的層上的某一紋理元素的方向向量。更具體地，面的選擇方法為：選擇座標中的最大值m，然後使用座標(s / m + 1) / 2和(t / m + 1) / 2來對對應的層進行取址，m、s和t的取值以及面的選擇如下表所示：

  • 立方圖分層紋理：

    立方圖分層紋理(只支援計算能力>=2.0的裝置)是一種泛化的立方圖紋理，唯一的區別就是立方圖分層紋理的層數不固定，但每層還是一個立方體。它可以使用一個整型索引和三個浮點紋理座標來取址，前者表示層數，後者表示層內的座標。

    這種立方圖分層紋理只能是一個cuda陣列，通過給cudaMalloc3DArray()函式傳遞cudaArrayCubemap標誌和cudaArrayLayered標誌建立，其在裝置函式中可通過texCubemapLayered()函式進行獲取，紋理過濾(參見官方手冊紋理獲取部分)只能在層內進行，而不能跨層進行。

紋理聚集

紋理聚集是一種特殊的紋理獲取，只能用在二維紋理上，可通過tex2Dgather()函式進行，其引數只比tex2D()函式多了一個com引數(可取0、1、2、3)。這個函式返回四個四位元組的數，每一個數對應四個已經在常規紋理獲取時用來做雙線性過濾的紋理元素之一的由comp引數指定的向量元素。比如，如果這些文理元素的值為(253, 20, 31, 255)、(250, 25, 29, 254)、(249, 16, 37, 253)和(251, 22, 30, 250)，當comp = 2時，tex2Dgather()函式會返回(31,29, 37, 30)。

注意紋理座標只能和精確到小數點後8位的引數正常工作，因此如果tex2D()函式使用1.0作為其引數之一的話(α或β)，tex2Dgather()函式可能返回異常值。例如，當紋理座標x為2.49805，xB = x - 0.5 = 1.99805，但是xB的小數部分是以固定八位的格式儲存的，因此0.99805更接近256.f / 256.f，而不是255.f / 256.f，因此xB的值就是2。tex2Dgather在這種情況下會返回x座標2和3，而非1和2.

紋理聚集只支援用cudaArrayTextureGather標誌建立的cuda陣列(以及計算能力>=2.0的裝置)，並且其寬高最大值要小於下表所示的值，這些值要比常規的紋理獲取要小。

表面記憶體

對計算能力≥2.0的裝置來說，用cudaArraySurfaceLoadStore標誌建立的cuad陣列可以通過表面物件或者表面引用相關的函式進行讀寫，下表列舉了不同計算能力的裝置的最大寬高和深度：

  • 表面物件API：

    使用型別為struct cudaResourceDesc的資源描述符的cudaCreateSurfaceObject()函式建立的是表面物件，下面的程式碼把簡單的轉換應用到了紋理中：

#include "texture_fetch_functions.h"
#include "surface_indirect_functions.h"
#include <stdio.h>

int width = 5, height = 2;
int size = width * height;
float *h_data = (float *) malloc(size * sizeof(float));
texture<float, cudaTextureType2D, cudaReadModeElementType> texRef;

__global__ void transformKernelSurfaceObj(cudaSurfaceObject_t inputSurObj, cudaSurfaceObject_t outputSurObj, int width, int height) {
    unsigned int x = blockIdx.x * blockIdx.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockIdx.y + threadIdx.y;

    if (x < width && y < height) {
        uchar4 data;
        surf2Dread(&data, inputSurObj, 4 * x, y); // 把資料從inputSurObj讀到data中
                                                  // 4 * x是因為每個執行緒讀取的資料data，要按4位元組排列到表面記憶體中(uchar4大小就是四位元組)
        surf2Dwrite(data, outputSurObj, 4 * x, y); // 把資料從data寫到outputObj中
    } 
}

int main() {
    for (int i = 0; i < size; i++) {
        h_data[i] = i;
    }

    // 分配cuda陣列
    cudaChannelFormatDesc des = cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsigned);
    cudaArray* cuInputArray;
    cudaMallocArray(&cuInputArray, &des, width, height, cudaArraySurfaceLoadStore);
    cudaArray* cuOutputArray;
    cudaMallocArray(&cuOutputArray, &des, width, height, cudaArraySurfaceLoadStore);

    cudaMemcpyToArray(cuInputArray, 0, 0, h_data, size * sizeof(float), cudaMemcpyHostToDevice);

    // 配置表面記憶體
    struct cudaResourceDesc resourceDesc;
    memset(&resourceDesc, 0, sizeof(resourceDesc));
    resourceDesc.resType = cudaResourceTypeArray;
    resourceDesc.res.array.array = cuInputArray;

    cudaSurfaceObject_t cuInputObj = 0, cuOutputObj = 0;
    cudaCreateSurfaceObject(&cuInputObj, &resourceDesc);
    resourceDesc.res.array.array = cuOutputArray;
    cudaCreateSurfaceObject(&cuOutputObj, &resourceDesc);

    // 呼叫核函式
    dim3 blockDim(16, 16);
    dim3 gridDim((width + blockDim.x - 1) / blockDim.x, (height + blockDim.y - 1) / blockDim.y);

    transformKernelSurfaceObj<<<gridDim, blockDim>>>(cuInputObj, cuOutputObj, width, height);

    // 把結果複製到主機中
    float *h_result = (float *) malloc(size * sizeof(float));
    cudaMemcpyFromArray(h_result, cuOutputArray, 0, 0, size * sizeof(float), cudaMemcpyDeviceToHost);

    for (int i = 0; i < size; i++) {
        printf("%f\n", h_result[i]);
    }

    // 釋放資源
    cudaDestroySurfaceObject(cuInputObj);
    cudaDestroySurfaceObject(cuOutputObj);
    cudaFreeArray(cuInputArray);
    cudaFreeArray(cuOutputArray);

    return 0;
}

    執行結果如下

表面引用API

當前檔案可用的表面引用可以通過surface型別的變數宣告：surface<void, Type> surfRef;，其中Type指定了表面引用的型別，可以是cudaSurfaceType1D、cudaSurfaceType2D、cudaSurfaceType3D、cudaSurfaceTypeCubemap、cudaSurfaceType1DLayered、cudaSurfaceType2DLayered和cudaSurfaceType3DLayered，但這是一個可選的引數，預設值為cudaSurfaceType1D。表面引用只能宣告為靜態全域性變數，而且不能作為引數傳給函式。在一個核函式可以使用表面引用來訪問cuda陣列前，此表面引用必須通過cudaBindSurfaceToArray()函式來繫結給這個cuda陣列。下面的程式碼將一個表面引用和一個cuda陣列cuArray繫結了起來：

  • 低階API：
const surface<void, cudaSurfaceType2D> surfRef;
const surfaceReference* surRefPtr;
cudaGetSurfaceReference(&surRefPtr, "surRef");
cudaChannelFormatDesc desc;
cudaArray* cuArray;
cudaGetChannelDesc(&desc, cuArray);
cudaBindSurfaceToArray(surfRef, cuArray);
  • 高階API：
cudaArray* cuArray;
cudaChannelFormatDesc des = cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsigned);
cudaMallocArray(&cuArray, &des, width, height, cudaArraySurfaceLoadStore);
surface<void, cudaSurfaceType2D> surfRef;
cudaBindSurfaceToArray(surfRef, cuArray);

cuda陣列必須使用維度和型別匹配的表面函式和維度匹配的表面引用來進行讀寫，否則讀寫的結果就會有異常。

不像紋理記憶體，表面記憶體使用位元組取址，這意味著通過紋理函式訪問紋理元素時使用的x座標需要被元素的位元組大小整除（這也是為何我在transformKernelSurfaceObj()函式中使用4 * x來定位資料的原因），以便通過表面函式來訪問這個紋理元素。比如，一個一維浮點cuda陣列的某個元素的紋理座標為x，而且此元素和紋理引用texRef、表面引用surfRef進行了繫結，所以應該通過tex1d(texRef, x)和surf1Dread(surfRef, x * 4)分別從texRef和surfRef中讀取此資料。類似地，某二維浮點cuda陣列中的某個元素紋理座標為(x, y)，而且此元素和紋理引用texRef、表面引用surfRef進行了繫結，所以應該通過tex2d(texRef, x， y)和surf2Dread(surfRef, x * 4, y)分別從texRef和surfRef中讀取此資料（y座標的位元組偏移量會通過cuda陣列的對應行自動計算）。下面程式碼也是通過核函式進行簡單的複製，但使用的是表面繫結機制：

#include "cuda_runtime.h"
#include "texture_fetch_functions.h"
#include "surface_indirect_functions.h"
#include "surface_functions.h"
#include <stdio.h>

int width = 5, height = 2;
int size = width * height;
float *h_data = (float *) malloc(size * sizeof(float));
const surface<void, cudaSurfaceType2D> inputSurf, outputSurf;

__global__ void transformKernelSurfaceRef(int width, int height) {
    unsigned int x = blockIdx.x * blockIdx.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockIdx.y + threadIdx.y;

    if (x < width && y < height) {
        uchar4 data;
        surf2Dread(&data, inputSurf, 4 * x, y);
        surf2Dwrite(data, outputSurf, 4 * x, y);
    }
}

int main() {
    for (int i = 0; i < size; i++) {
        h_data[i] = i;
    }

    // 分配cuda陣列
    cudaChannelFormatDesc des = cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsigned);
    cudaArray* cuInputArray;
    cudaMallocArray(&cuInputArray, &des, width, height, cudaArraySurfaceLoadStore);
    cudaArray* cuOutputArray;
    cudaMallocArray(&cuOutputArray, &des, width, height, cudaArraySurfaceLoadStore);

    cudaMemcpyToArray(cuInputArray, 0, 0, h_data, size * sizeof(float), cudaMemcpyHostToDevice);

    // 繫結表面引用
    cudaBindSurfaceToArray(inputSurf, cuInputArray);
    cudaBindSurfaceToArray(outputSurf, cuOutputArray);

    // 執行核函式
    dim3 blockDim(16, 16);
    dim3 gridDim((width + blockDim.x - 1) / blockDim.x, (height + blockDim.y - 1) / blockDim.y);

    transformKernelSurfaceRef<<<gridDim, blockDim>>>(width, height);

    // 複製結果到主機
    float *h_result = (float *) malloc(size * sizeof(float));
    cudaMemcpyFromArray(h_result, cuOutputArray, 0, 0, size * sizeof(float), cudaMemcpyDeviceToHost);

    for (int i = 0; i < size; i++) {
        printf("%f\n", h_result[i]);
    }

    // 釋放資源
    cudaFreeArray(cuInputArray);
    cudaFreeArray(cuOutputArray);

    return 0;
}

執行結果如下：

立方圖表面

使用surfCubemapred()和surfCubemapwrite()函式可以把立方圖表面記憶體作為二維分層表面來訪問，也就是說使用整數索引表示面、二維浮點紋理座標表示和這個面對應的層上紋理元素的座標，面的排序方法如下表所示(和立方圖紋理中的方法一樣)

分層立方圖表面：

        通過surfCubemapLayeredRead()和surfCubemapLayeredWrite()函式可以把分層立方圖表面當成二維分層表面來訪問，也就是說使用一個整數表示某一立方圖的某個面，二維浮點紋理座標來定位和這個面對應的層上的紋理元素，面的排序方法和立方圖表面一樣。比如索引(2 * 6) + 3就表示第3個立方圖(cubemap 2)的第4張面(face 4)

cuda陣列

cuda陣列是為紋理獲取優化的不透明記憶體佈局，他們可以是一維、二維或三維的，由擁有單位元組、雙位元組或四位元組無符號整數(或者是雙位元組、四位元組浮點數)組成的1、2、4維向量的元素組成，一維向量元素對應單位元組整數和雙位元組浮點數，二維向量對應雙位元組整數和浮點數，三維向量對應四位元組整數和浮點數。cuda陣列只能通過核函式的紋理獲取或表面記憶體的讀寫來訪問，具體請參見上兩小節。

讀寫一致性

紋理和表面記憶體是快取的，在同一個核函式呼叫中，涉及全域性記憶體寫和表面記憶體寫的快取不會保持一致性，所以在同一核函式呼叫中對已經通過全域性寫或表面寫的地址進行紋理讀和表面讀的話，會返回未知值。換言之，只有某記憶體位置被之前的核函式呼叫或記憶體複製更新(而非被同一核函式的任何呼叫更新)之後，一個執行緒才可以安全地進行紋理或表面讀取。

所以在上面寫過的幾個紋理或表面記憶體核函式中，我要麼是隻對紋理記憶體或表面記憶體讀，要麼是讀寫不同的記憶體地址，但是記憶體地址的紋理座標都是一致的，故而保證了讀寫一致性。

圖形互動性

一些來自OpenGL、Direct3D這種源的資源可以被對映到cuda的地址空間中，從而支援cuda讀取被OpenGL或Direct3D寫入的資料或者寫入可以被OpenGL或Direct3D消費的資料。

在可以使用和OpenGL或Direct3D互動的函式之前，這種資源必須現在cuda中註冊。那些互動函式會返回一個指向cuda影像的struct cudaGraphicsResource型別的指標，而註冊資源是代價昂貴的，因此每個資源只能註冊一次，另外可以通過cudaGraphicsUnregisterResource()函式來為資源登出。每個要使用資源的cuda上下文都需要對資源單獨註冊。

一旦一個資源被註冊到了cuda中，它就可以多次對映和去對映，通過使用cudaGraphicsMapResources()和cudaGraphicsUnmapResources()函式，cudaGraphicsResourcesSetMapFlags()函式可以指明cuda驅動用來優化資源管理的資源訪問許可權(只讀、只寫等)。

對映好的資源可以通過使用裝置記憶體地址在核函式中進行讀寫，這種裝置記憶體地址可以通過cudaGraphicsResourcesGetMappedPointer()或者cudaGraphicsSubResourcesGetMappedArray()函式返回得到，前者對應快取的地址，後者對應cuda陣列的地址。

通過未註冊的cuda上下文或者OpenGL、Direct3D直接訪問對映好的資源會產生未知結果。

OpenGL互動性

可以被對映到cuda地址空間的OpenGL資源有OpenGL快取、紋理和渲染快取(renderbuffer)物件。快取物件可以通過cudaGraphicsGLRegisterBuffer()函式註冊，在cuda中它表現為裝置指標，因此可以通過核函式或者cudaMemcpy()函式進行讀寫；紋理或者渲染快取物件通過cudaGraphicsGLRegisterImage()函式註冊，在cuda中它們表現為繫結到紋理或表面引用的cuda陣列，如果在註冊時使用了cudaGraphicsRegisterFlagsSurfaceLoadStore標誌的話，它們就可以通過表面寫函式來寫了。這些陣列也可以通過cudaMemcpy2D()函式進行讀寫。cudaGraphicsGLRegisterImage()函式支援一維、二維或四維向量的紋理和OpenGL的浮點型別(GL_RGBA_FLOAT32等)、正規化整數(GL_RGDA8、GL_INTENSITY16等)和非正規化整數(GL_RGBA8UI)，注意因為非正規化整數格式要求OpenGL版本為3.0，那麼它們只能通過著色器(shaders)而不是固定函式流程(fixed function pipeline)來進行寫。

資源正在被著色的OpenGL上下文必須是當前使用OpenGL互動API的主機執行緒。注意，當一個OpenGL紋理是無繫結(bindless)建立時(比如通過glGetTextureHandleX/glGetImageHandleX系列API請求紋理或影像控制程式碼時)的，那它不能在cuda中註冊，我們需要在請求紋理或影像控制程式碼前為紋理進行互動註冊。下面的程式碼使用核函式動態修改一個儲存在結點快取物件中的尺寸為width * height的二維網格：

GLuint positionsVBO;
struct cudaGraphicsResource* positionsVBO_CUDA;

int main() {
    .....
    // 為裝置0初始化OpenGL和GLUT，並設定OpenGL的上下文為當前上下文
    glutDisplayFunc(display);

    // 明確使用裝置0
    cudaSetDevice(0);

    // 建立buffer物件，並在cuda中註冊
    glGenBuffers(1, &positionsVBO);
    glBindBuffer(GL_ARRAY_BUFFER, positionsVBO);

    unsigned int size = width * height * 4 * sizeof(float);
    glBufferData(GL_ARRAY_BUFFER, size, 0, GL_DYNAMIC_DRAW);
    glBindBuffer(GL_ARRAY_BUFFER, 0);

    cudaGraphicsGLRegisterBuffer(&positionsVBO_CUDA, positionsVBO, cudaGraphicsMapFlagsWriteDiscard);

    // 啟動渲染迴圈
    glutMainLoop();
    ...
}

void display() {
    // 對映cuda要寫入的快取物件
    float4* positions;
    cudaGraphicsMapResources(1, &positionsVBO_CUDA, 0);

    size_t num_bytes;
    cudaGraphicsResourceGetMappedPointer((void**)&positions, &num_bytes, positionsVBO_CUDA));

    // 執行核函式
    dim3 dimBlock(16, 16, 1);
    dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1);

    createVertices<<<dimGrid, dimBlock>>>(positions, time, width, height);

    // 快取物件去對映
    cudaGraphicsUnmapResources(1, &positionsVBO_CUDA, 0);

    // 從快取物件中渲染
    glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);
    glBindBuffer(GL_ARRAY_BUFFER, positionsVBO);
    glVertexPointer(4, GL_FLOAT, 0, 0);
    glEnableClientState(GL_VERTEX_ARRAY);
    glDrawArrays(GL_POINTS, 0, width * height);
    glDisableClientState(GL_VERTEX_ARRAY);

    // 交換快取
    glutSwapBuffers();
    glutPostRedisplay();
}

void deleteVBO() {
    cudaGraphicsUnregisterResource(positionsVBO_CUDA);
    glDeleteBuffers(1, &positionsVBO);
}

__global__ void createVertices(float4* positions, float time, unsigned int width, unsigned int height) {
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;

    // 計算uv座標
    float u = x / (float)width;
    float v = y / (float)height;
    u = u * 2.0f - 1.0f;
    v = v * 2.0f - 1.0f;

    // 計算簡單正弦波模板
    float freq = 4.0f;
    float w = sinf(u * freq + time) * cosf(v * freq + time) * 0.5f;

    // 寫入資料
    positions[y * width + x] = make_float4(u, w, v, 1.0f);
}

在windows和Quadro GPU上，cudaWGLGetDevice()可以被用來讀取和wglEnumGpusNV()函式返回的控制程式碼相關的cuda裝置。在OpenGL渲染在Quadro GPU上執行、cuda計算在系統中其他GPU上執行的多GPU配置下，Quadro GPU關於提供比GeForce和Tesla GPU更好的OpenGL互動性表現。

Direct3D互動性

Direct3D互動性支援Direct3D 9Ex、Direct3D 10和Direct3D 11。cuda上下文只能和完全滿足這些條件的Direct3D裝置互動：Direct3D 9Ex裝置必須通過把DeviceType設定成D3DDEVTYPE_HAL、BehaviorType設定成D3DCREATE_HARDWARE_VERTEXPROCESSING來建立；Direct3D 10和Direct3D 1裝置則需要通過把DriverType設定成D3D_DRIVER_TYPE_HARDWARE來建立。

可以對映到cuda地址空間的Direct3D資源有Direct3D快取、紋理和表面，這些資源通過cudaGraphicsD3D9RegisterResource()、cudaGraphicsD3D10RegisterResource()和cudaGraphicsD3D11RegisterResource()註冊，下面的程式碼使用核函式動態修改一個儲存在結點快取物件的二維width * height網格：

  • Direct3D 9版本：
IDirect3D9 *D3D;
IDirect3DDevice9 *device;

struct CUSTOMVERTEX {
    FLOAT x, y, z;
    DWORD color;
};

IDirect3DVertexBuffer9 *positionsVB;
struct cudaGraphicsResource *positionsVB_CUDA;

int main() {
    int dev;

    // 初始化Direct3D
    D3D = Direct3DCreate9Ex(D3D_SDK_VERSION);

    // 獲取支援cuda的介面卡
    unsigned int adapter = 0;
    for (; adapter < g_pD3D->GetAdapterCount(); adapter++) {
        D3DADAPTER_IDENTIFIER9 adapterId;
        g_pD3D->GetAdapterIdentifier(adapter, 0, &adapterId);
        if (cudaD3D9GetDevice(&dev, adapterId.DeviceName) == cudaSuccess)
            break;
    }

    // 建立裝置
    ...
    D3D->CreateDeviceEx(adapter, D3DDEVTYPE_HAL, hWnd, D3DCREATE_HARDWARE_VERTEXPROCESSING, &params, NULL, &device);
    // 使用裝置
    cudaSetDevice(dev);

    // 建立並註冊結點快取
    unsigned int size = width * height * sizeof(CUSTOMVERTEX);
    device->CreateVertexBuffer(size, 0, D3DFVF_CUSTOMVERTEX, D3DPOOL_DEFAULT, &positionsVB, 0);
    cudaGraphicsD3D9RegisterResource(&positionsVB_CUDA, positionsVB, cudaGraphicsRegisterFlagsNone);
    cudaGraphicsResourceSetMapFlags(positionsVB_CUDA, cudaGraphicsMapFlagsWriteDiscard);

    // 啟動渲染迴圈
    while (...) {
        ...
        Render();
        ...
    }
    ...
}

void Render() {
    // 對映結點快取，以便cuda寫
    float4 *positions;
    cudaGraphicsMapResources(1, &positionsVB_CUDA, 0);
    size_t num_bytes;
    cudaGraphicsResourceGetMappedPointer((void **) &positions, &num_bytes, positionsVB_CUDA);

    // 執行核函式
    dim3 dimBlock(16, 16, 1);
    dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1);
    createVertices<<<dimGrid, dimBlock>>>(positions, time, width, height);

    // 去對映節點快取
    cudaGraphicsUnmapResources(1, &positionsVB_CUDA, 0);
    
    // 畫圖並展示
    ...
}

void releaseVB() {
    cudaGraphicsUnregisterResource(positionsVB_CUDA);
    positionsVB->Release();
}

__global__ void createVertices(float4 *positions, float time, unsigned int width, unsigned int height) {
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;

    // 計算uv座標
    float u = x / (float) width;
    float v = y / (float) height;
    u = u * 2.0f - 1.0f;
    v = v * 2.0f - 1.0f;

    // 計算簡單正弦波模板
    float freq = 4.0f;
    float w = sinf(u * freq + time) * cosf(v * freq + time) * 0.5f;

    // 寫入資料
    positions[y * width + x] = make_float4(u, w, v, __int_as_float(0xff00ff00));
}
  • Direct3D 10版本：
ID3D10Device *device;

struct CUSTOMVERTEX {
    FLOAT x, y, z;
    DWORD color;
};

ID3D10Buffer *positionsVB;
struct cudaGraphicsResource *positionsVB_CUDA;

int main() {
    int dev;

    // 獲取支援cuda的介面卡
    IDXGIFactory *factory;
    CreateDXGIFactory(__uuidof(IDXGIFactory), (void **) &factory);
    IDXGIAdapter *adapter = 0;
    for (unsigned int i = 0; !adapter; ++i) {
        if (FAILED(factory->EnumAdapters(i, &adapter))
            break;
        if (cudaD3D10GetDevice(&dev, adapter) == cudaSuccess)
            break;
        adapter->Release();
    }
    factory->Release();

    // 建立交換鏈和裝置
    ...
    D3D10CreateDeviceAndSwapChain(adapter, D3D10_DRIVER_TYPE_HARDWARE, 0, D3D10_CREATE_DEVICE_DEBUG, D3D10_SDK_VERSION, &swapChainDesc, &swapChain, &device);
    adapter->Release();

    // 使用裝置
    cudaSetDevice(dev);

    // 建立並註冊結點快取
    unsigned int size = width * height * sizeof(CUSTOMVERTEX);
    D3D10_BUFFER_DESC bufferDesc;
    bufferDesc.Usage = D3D10_USAGE_DEFAULT;
    bufferDesc.ByteWidth = size;
    bufferDesc.BindFlags = D3D10_BIND_VERTEX_BUFFER;
    bufferDesc.CPUAccessFlags = 0;
    bufferDesc.MiscFlags = 0;

    device->CreateBuffer(&bufferDesc, 0, &positionsVB);
    cudaGraphicsD3D10RegisterResource(&positionsVB_CUDA, positionsVB, cudaGraphicsRegisterFlagsNone);
    cudaGraphicsResourceSetMapFlags(positionsVB_CUDA, cudaGraphicsMapFlagsWriteDiscard);

    // 啟動渲染迴圈
    while (...) {
        ...
        Render();
        ...
    }
    ...
}

void Render() {
    // 對映結點快取，以便cuda寫資料
    float4 *positions;
    cudaGraphicsMapResources(1, &positionsVB_CUDA, 0);
    size_t num_bytes;
    cudaGraphicsResourceGetMappedPointer((void **) &positions, &num_bytes, positionsVB_CUDA);

    // 執行核函式
    dim3 dimBlock(16, 16, 1);
    dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1);

    createVertices<<<dimGrid, dimBlock>>>(positions, time, width, height);

    // 結點快取去對映
    cudaGraphicsUnmapResources(1, &positionsVB_CUDA, 0);

    // 繪畫並展示
    ...
}

void releaseVB() {
    cudaGraphicsUnregisterResource(positionsVB_CUDA);
    positionsVB->Release();
}

__global__ void createVertices(float4 *positions, float time, unsigned int width, unsigned int height) {
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;

    // 計算uv座標
    float u = x / (float) width;
    float v = y / (float) height;

    u = u * 2.0f - 1.0f;
    v = v * 2.0f - 1.0f;

    // 計算簡單的正弦波模板
    float freq = 4.0f;
    float w = sinf(u * freq + time) * cosf(v * freq + time) * 0.5f;

    // 寫入資料
    positions[y * width + x] = make_float4(u, w, v, __int_as_float(0xff00ff00));
}
  • Direct3D 11版本：
ID3D11Device *device;

struct CUSTOMVERTEX {
    FLOAT x, y, z;
    DWORD color;
};

ID3D11Buffer *positionsVB;
struct cudaGraphicsResource *positionsVB_CUDA;

int main() {
    int dev;

    // 得到支援cuda的介面卡
    IDXGIFactory *factory;
    CreateDXGIFactory(__uuidof(IDXGIFactory), (void **) &factory);
    IDXGIAdapter *adapter = 0;
    for (unsigned int i = 0; !adapter; ++i) {
        if (FAILED(factory->EnumAdapters(i, &adapter))
            break;
        if (cudaD3D11GetDevice(&dev, adapter) == cudaSuccess)
            break;
        adapter->Release();
    }
    factory->Release();

    // 建立交換鏈和裝置
    ...
    sFnPtr_D3D11CreateDeviceAndSwapChain(adapter, D3D11_DRIVER_TYPE_HARDWARE, 0, D3D11_CREATE_DEVICE_DEBUG, featureLevels, 3, D3D11_SDK_VERSION, &swapChainDesc, &swapChain, &device, &featureLevel, &deviceContext);
    adapter->Release();

    // 使用裝置
    cudaSetDevice(dev);

    // 建立並註冊結點快取
    unsigned int size = width * height * sizeof(CUSTOMVERTEX);
    D3D11_BUFFER_DESC bufferDesc;
    bufferDesc.Usage = D3D11_USAGE_DEFAULT;
    bufferDesc.ByteWidth = size;
    bufferDesc.BindFlags = D3D11_BIND_VERTEX_BUFFER;
    bufferDesc.CPUAccessFlags = 0;
    bufferDesc.MiscFlags = 0;
    device->CreateBuffer(&bufferDesc, 0, &positionsVB);
    cudaGraphicsD3D11RegisterResource(&positionsVB_CUDA, positionsVB, cudaGraphicsRegisterFlagsNone);
    cudaGraphicsResourceSetMapFlags(positionsVB_CUDA, cudaGraphicsMapFlagsWriteDiscard);

    // 啟動渲染迴圈
    while (...) {
        ...
        Render();
        ...
    }
    ...
}

void Render() {
    // 對映結點快取，以便cuda寫入
    float4 *positions;
    cudaGraphicsMapResources(1, &positionsVB_CUDA, 0);
    size_t num_bytes;
    cudaGraphicsResourceGetMappedPointer((void **) &positions, &num_bytes, positionsVB_CUDA));

    // 執行核函式
    dim3 dimBlock(16, 16, 1);
    dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1);

    createVertices<<<dimGrid, dimBlock>>>(positions, time, width, height);

    // 去對映結點快取
    cudaGraphicsUnmapResources(1, &positionsVB_CUDA, 0);

    // 繪圖並展示
    ...
}

void releaseVB() {
    cudaGraphicsUnregisterResource(positionsVB_CUDA);
    positionsVB->Release();
}

__global__ void createVertices(float4 *positions, float time, unsigned int width, unsigned int height) {
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;

    // 計算uv座標
    float u = x / (float) width;
    float v = y / (float) height;

    u = u * 2.0f - 1.0f;
    v = v * 2.0f - 1.0f;

    // 計算簡單正弦波模板
    float freq = 4.0f;
    float w = sinf(u * freq + time) * cosf(v * freq + time) * 0.5f;

    // 寫入資料
    positions[y * width + x] = make_float4(u, w, v, __int_as_float(0xff00ff00));
}

SLI互動

在多核GPU系統中，所有的支援cuda的GPU都可以通過cuda驅動和執行時作為獨立裝置訪問，但是當系統處在SLI模式下時有這麼幾點特殊的考慮：

    首先，一個GPU上的一個cuda裝置的記憶體分配將會消耗Direct3D或OpenGL裝置SLI配置中的GPU記憶體，所以記憶體分配可能比預期要更早失敗。

    第二，應用應該為SLI配置中的每個GPU都建立一個cuda上下文。儘管這不是硬性要求，但它避免了裝置間不必要的資料遷移。我們可以為Direct3D使用cudaD3d[9|10|11]()、為OpenGL使用cudaGLGetDevices()系列方法來為裝置指定渲染當前和下一幀的cuda裝置控制程式碼。有了由這些函式返回的裝置資訊，當deviceList引數被設定成cudaD3D[9|10|11]DeviceListCurrentFrame或cudaGLDeviceListCurrentFrame時，應用可以選擇合適的裝置並把Diretct3D或OpenGL資源對映到cuda裝置上。

注意，cudaGraphicsD3D[9|10|11]RegisterResource()和cudaGraphicsGLRegister[Buffer|Image]()只能用在發生註冊的裝置上，因此當SLI配置中，不同幀的資料在不同cuda裝置上計算時，為每臺裝置單獨註冊資源就是有必要的。

關於cuda執行時如何與Direct3D和OpenGL互動的細節，請參見上兩小節。

版本和適配性

當開發一個cuda應用時，開發者需要注意兩個版本號：描述一般標準和特徵的計算能力和描述驅動、執行時API支援的特徵的cuda驅動API版本號。驅動API的版本號在驅動標頭檔案中被定義為CUDA_VERSION，允許開發者檢查他們的應用是否要求一個比已安裝的裝置驅動更新的裝置驅動，這一點很重要，因為驅動API是向後相容的，意味著用特定版本的驅動API編譯的外掛、庫(包括C執行時)、應用可以繼續在隨後發行的裝置驅動上工作，如下圖所示

但是驅動API不是向前相容的，意味著用特定版本的驅動API編譯的外掛、庫(包括C執行時)、應用不會在之前版本的裝置驅動上工作。再就是需要注意，支援多版本的混合與匹配有一些限制：

    1、因為系統中一次只能安裝一個版本的cuda驅動，那麼被安裝的驅動的版本必須>=任何應用、外掛或庫編譯時用的驅動API最大版本；

    2、被一個應用使用的所有外掛和庫必須使用相同版本的cuda執行時，除非它們想要和cuda執行時進行動態連結，在這種情況下，統一程式空間可以存在多個版本的執行時。注意如果nvcc被用來連結應用，cuda執行時庫的靜態版本將會被預設使用，所有的cuda工具包庫將會和cuda執行時靜態連結；

    3、被一個應用使用的所有外掛和庫如果在執行時要使用一些庫(比如cuFFT、cuBLAS等)，那麼這些庫必須使用相同版本

計算模式

在執行於Windows Server 2008及以後的Windows或者Linux上的Tesla裝置上，我們可以使用nvidia-smi給系統中任何裝置設定三種模式之一，見下表

計算模式

描述

預設

多個主機執行緒可以通過使用執行時API時在裝置上呼叫cudaSetDevice()或者使用驅動API時讓和此裝置相關的上下文為當前上下文的方法同時共享裝置

程式獨佔

系統中的所有程式只能在裝置上建立一個cuda上下文，這個上下文可以是建立這個上下文的程式裡多個執行緒共有的當前上下文。換句話說，一次只能有一個程式佔用裝置。

禁止

此裝置上不允許 cuda上下文

也就是說，如果一個使用執行時API的主機執行緒沒有明確呼叫cudaSetDevice()的話，它可能和一個不是device0的裝置關聯，如果device0處在禁止模式下或者處於程式獨佔模式下且正在被別的程式使用時。cudaSetValidDevices()可以被用來從一個裝置優先順序表中對裝置進行設定。

還要注意的是，對於使用先進的Pascal架構的裝置(計算能力主版本號>=6)來說，搶佔式計算是被支援且預設開啟的。這允許任務在指令級別被搶佔，而不是在之前的Maxwell和Kepler GPU架構了使用的執行緒塊級別，這給程式帶來的好處是：避免了長時間執行的核函式獨佔系統或執行超市。然而，使用搶佔式計算時有著一些切換上下文的效能損耗。正在使用的裝置是否支援搶佔式計算可以通過cudaDeviceGetAttribute()函式查到的結構體裡cudaDevAttrComputePreemptionSupported欄位進行判斷，希望避免使用不同程式帶來的上下文切換效能損耗的使用者可以通過選擇程式獨佔模式來確保GPU上一次只有一個活動程式。應用可以通過檢查computeMode裝置屬性來查詢裝置的計算模式。

模式切換

擁有螢幕輸出的GPU會把一些記憶體專門用在所謂的主表面(primary surface)上，這個主表面用來重新整理輸出被使用者觀看的顯示器裝置。當使用者通過改變分別率或者顯示位深度(使用英偉達控制皮膚或者windows上的顯示控制皮膚)來初始化顯示器的模式切換(mode switch)時，主表面需要的記憶體數量也會隨之變化。例如如果使用者把解析度從1280 * 1024 * 32位切換到1600 * 1200 * 32位時，主表面需要的記憶體會從5.24MB增加到7.68MB（使用反鋸齒的全屏影像應用的主表面可能需要更多）。在windows上，可能導致顯示模式切換的事件還包括啟動全屏DirectX應用、使用Alt + Tab從全屏DirectX應用中移除任務或者使用Ctrl + Alt + Del進行鎖屏

如果模式切換增加了主表面需要的記憶體數量，系統可能會調撥一些分配給cuda應用的記憶體，所以模式切換會導致所有cuda執行時API呼叫失敗並且返回一個無效上下文錯誤。

針對Windows的Tesla計算叢集模式

使用nvidia-smi，windows裝置驅動可以為計算能力>=2.0的Tesla和Quadro系列裝置進入Tesla計算叢集(TCC)模式，這個模式的主要好處如下：

    1、它讓叢集中沒有整合英偉達的影像裝置使用這些GPU；

    2、它讓GPU通過遠端桌面直接可用，或者可以通過依賴於遠端桌面的叢集管理系統訪問；

    3、它讓作為windows服務的應用(比如會話0中的應用)可以使用GPU；

但是，TCC模式移除了對任何影像功能的支援。

結語

以上就是第三章程式設計介面的翻譯，內容龐雜，下一章將翻譯硬體實現部分。

相關文章