CUDA 7.0 速查手冊

Create by Jane/Santaizi 03:57:00 3/14/2016
All right reserved.

速查手冊基於 CUDA 7.0 toolkit documentation 並對原文進行了精簡.

手冊專一於CUDA的GPU計算方面,不涉及圖形顯示.如需完整檔請查原文http://docs.nvidia.com/cuda/index.html#axzz42oaojUNj

3.2.4 Page-Locked Host Memory

在Host CPU程序中劃出的內存區域供多GPU設備共享使用

使用方法：

1. cudaHostAlloc() , cudaFreeHost() 分配,釋放 page-locked host 內存
2. cudaHostRegister() page-locks 一個由 malloc 獲得的內存塊

優勢:

1. 在 page-locked 的內存和 GPU 內存之間能夠在 kernel 執行時異步拷貝
2. 一些GPU設備能夠直接映射 page-locked 的CPU內存，跳過拷貝步驟
3. 在一些有 front-side bus(前端總線)的設備上, host 內存和 GPU內存能夠以更高速度拷貝,用 write-combining 特性的話,速度將更快.

缺點:

1. Page-locked host 內存是稀缺資源,因此在分配時容易失敗.
2. 分配大量page-locked 內存將致使pageable 內存減小,影響整體性能.

3.2.4.1 Portable memory

在多GPU設備之間充當共享內存角色.是一個 Unified Virtual Address Space.

使用方法:

1. cudaHostAlloc(), 傳入 flag cudaHostAllocPortable
2. cudaHostRegister(), 傳入 flag cudaHostRegisterPortable

3.2.4.2 Write-Combining Memory

默認 page-locked host 內存是以 cacheable 方式分配的.你能夠用 Write-Combining 方式分配. Write-Combining 內存釋放 host 的L1,L2緩存資源, 在通過PCI總線時提升最多40%的速度.

使用方法：

1. cudaHostAlloc(), 傳入 flag cudaHostAllocWriteCombined

優勢:

1. 增長高速緩存的容量,使得CPU到GPU內存之間的內存拷貝加速

缺點:

1. 從 Host 環境中讀取 write-combining 內存很是慢,因此只適合 Host 往裏寫數據(而不讀取)的狀況.

3.2.4.3 Mapped Memory

host CPU內存和GPU內存之間的內存地址映射.
host 和GPU有對應的內存指針. 函數返回的是 host 指針, GPU內存指針需用 cudaHostGetDevicePointer() 獲取,獲取的GPU內存指針能夠在 kernel中去使用.

使用方法:

1. cudaHostAlloc(), 傳入 flag cudaHostAllocMapped
2. cudaHostRegister(), 傳入 flag cudaHostRegisterMapped

優勢:

1. 不用在CPU-GPU之間拷貝內存數據
2. There is no need to use streams (see Concurrent Data Transfers) to overlap data transfers with kernel execution; the kernel-originated data transfers automatically overlap with kernel execution.

缺點:

1. 內存映射破壞了數據的原子性, 應用程序必須使用 stream 或 events 來避免數據讀寫順序控制和數據同步問題.

注意:在獲取GPU內存指針以前必須使用 cudaSetDeviceFlags(), 傳入 flag cudaDeviceMapHost.不然 cudaHostGetDevicePointer() 會致使錯誤. cudaHostGetDevicePointer() 錯誤也會在 設備GPU不支持內存映射時產生.

屬性查看:

1. 使用設備屬性 canMapHostMemory = 1(支持)查詢設備支持狀況.

一樣注意: Atomic Functions 對映射內存的原子操做對host 和 GPU設備來講也是非原子的.

3.2.5 Asynchronous Concurrent Execution

CUDA中如下操做是互相獨立且併發的：

1. Host(CPU) 中的計算
2. Device(GPU) 中的計算
3. 從 Host 到 Device 的數據傳遞
4. 從 Device 到 Host 的數據傳遞
5. 在單個 Device 內存中的數據傳遞
6. 在多個 Device 內存之間的數據傳遞

除了Host(CPU)環境內部的數據傳遞是順序同步的以外,一切和GPU有關的數據傳遞都是獨立併發的(異步).

3.2.5.1 Concurrent Execution between Host and Device

Host 中的併發操做是經過異步函數庫方法實現的,並在啓動後直接移交控制權回 Host 主線程,且並不保證GPU設備已經計算完相應任務.這個模式相似於 event loop,任務在異步啓動後排隊等待被處理,而不阻塞主線程.下面幾種操做對 host 來講是異步調用的：

1. Kernel launch (kernel 函數的啓動)
2. 在單個GPU設備中的內存傳遞
3. Host 內存拷貝至 Device 內存 (64KB甚至更少的數據塊傳遞也是異步的)
4. 任何以 Async 爲後綴的內存拷貝函數
5. Memory set function calls

能夠設置環境變量 CUDA_LAUNCH_BLOCKING = 1 來禁止 kernel 函數的異步啓動. 這個特性只能用來 debug (Notice: Debug Only!).
另外在使用 Visual Profiler Nsight 採集硬件計數器的時候 kernel 的啓動也是同步的, 除非 concurrent kernel profiling 選項被開啓. 以 Async 後綴的內存拷貝一樣在 not page-locked 的 host 內存中是同步的.

3.2.5.2 Concurrent Kernel Execution

設備計算能力超過2.X均可以併發執行 kernel 函數. 在附錄表13中可查. 不一樣CUDA context中的kernel 不能併發. 使用大量 texture 和 內存的 kernel 也不太可能與其餘併發.

屬性查看:

1. 設備屬性 concurrentKernels=1 查詢設備支持狀況(see Device Enumeration).

3.2.5.3 Overlap of Data Transfer and Kernel Execution

一些設備可併發執行 kernel函數和異步GPU內存拷貝操做. Host 內存塊必須是 page-locked的. Device內存內部的多個內存拷貝(intra-device)和 kernal 函數甚至能夠同時執行.

屬性查看:

1. 設備屬性 asyncEngineCount > 0 查詢設備支持狀況(see Device Enumeration).
2. concurrentKernels = 1, 而且 asyncEngineCount > 0 查詢多個Device內部內存拷貝和 kernal 的併發操做支持.

3.2.5.4 Concurrent Data Transfers

設備計算能力超過2.X 能夠執行併發內存拷貝.Host 內存必須爲 page-locked.

屬性查看:

1. 設備屬性 asyncEngineCount = 2 查詢設備支持狀況(see Device Enumeration).

3.2.5.5 Streams

應用程序使用 streams 來管理上述全部併發操做.一個 stream 就是一串順序命令. 不一樣 streams 之間是亂序或同步執行的.

3.2.5.5.1 Creation and Destruction

使用方法：

下例中建立了2個 stream 並分配了一個 float array 的 page-locked 內存塊給 hostPtr

cudaStream_t stream[2];
for (int i = 0; i < 2; ++i)
    cudaStreamCreate(&stream[i]);
float * hostPtr;
cudaMallocHost(&hostPtr, 2*size);

每一個 stream 都被指定順序執行下述操做:

1. Host -> Device 的內存拷貝
2. kernel 啓動

3. Device -> Host 的內存拷貝

   for (int i = 0; i < 2; ++i)
   {
   cudaMemcpyAsync(inputDevPtr + isize, hostPtr + isize, size, cudaMemcpyHostToDevice, stream[i]);
   MyKernel <<<100, 512, 0, stream[i]>>>(outputDevPtr + isize, inputDevPtr + isize, size);
   cudaMemcpyAsync(hostPtr + isize, outputDevPtr + isize, size, cudaMemcpyDeviceToHost, stream[i]);
   }

釋放 streams 使用 cudaStreamDestroy().

for (int i = 0; i < 2; ++i)
    cudaStreamDestory(stream[i]);

cudaStreamDestory() 等待全部 stream 中的命令執行完畢後再銷燬 stream 並返回控制權給 host 主線程,也就是說它是一個阻塞的強制同步函數.

3.2.5.5.2 Default Stream

kernel 啓動和 host-device 之間的內存拷貝不須要設置特殊 stream 參數(默認設置爲 0 ), 他們在stream中順序執行.

使用方法:

1. 使用 flag --default-stream per-thread 編譯或者在 include cuda.h和cuda_runtime.h頭以前定義宏 CUDA_API_PER_THREAD_DEFAULT_STREAM 那麼一般 stream 將都是默認的 stream, 且每一個host 線程都有本身的 stream.
2. 使用 flag --default-stream legacy 編譯, 那麼默認 stream 將會是特殊的,名叫 NULL stream ,且每一個 device 對每一個 host 線程來講都有一個單獨的 stream. NULL stream 由於它隱含的同步特性而比較特別.詳細描述在 Implicit Synchronization之中
3. 對那些沒有設置 flag --default-stream 的編譯來講 --default-stream legacy 爲默認的設置.

3.2.5.5.3 Explicit Synchroonization

下面列舉了幾種顯式同步各個 streams 的方法. 爲了不運算性能下降, 全部同步函數都應在須要時間控制和分離啓動與內存拷貝(順序控制)時使用.

使用方法:

1. cudaDeviceSynchronize() 暫停主線程並等待全部 host 線程中的 streams 中的全部命令都執行完畢,再把控制權還給主線程.
2. cudaStreamSynchronize() 接受一個 stream 爲參數,等待該 stream 中全部命令執行完畢. 它被用來同步 host 中的某一個 stream,並容許其餘 stream 異步處理.
3. cudaStreamWaitEvent() 接受一個 stream 和一個 event 爲參數, 使得全部以後加入該 stream 的事件都等待相關 event 結束以後再開始執行. stream 參數能夠爲 0,代表任何命令在cudaStreamWaitEvent()執行以後,不管被加入哪一個 stream 之中都必須等待 event 結束才能開始執行.
4. cudaStreamQuery() 能夠用來查詢在某個 stream 中全部命令是否已經所有執行完畢.

3.2.5.5.4 Implicit Synchronization

若是碰到如下狀況, 兩個 stream 中的命令是不能併發執行的：

1. page-locked 的 Host 內存分配
2. device(GPU) 內存分配
3. device(GPU) 內存設置(賦值)
4. 在同一個 Device 內存中不一樣地址之間的內存拷貝
5. 任何在 NULL stream 上的 CUDA命令
6. L1/shared 內存的設置切換

對於那些支持併發 kernel 執行的設備來講, 任何操做都須要附加一個檢查來查看 streamed kernel launch是否已經完成：

1. 只有在CUDA context中全部stream 中全部 thread blocks 的kenel 啓動以後才能執行.
2. 只有在CUDA context中全部kernel 啓動被確認完成以後才能執行

由於操做須要作一個 cudaStreamQuery()檢查,因此爲了提升性能應遵循下面兩個習慣:

1. 全部互相獨立的操做應該放在非獨立操做以前完成
2. 任何形式的同步都應放到最後.

3.2.5.5.5 Overlapping Behavior

兩個 stream 上的命令能夠根據設備的支持狀況進行重疊(併發)執行. 對於3.2.5.5.1 Creation and Destruction 例子

for (int i = 0; i < 2; ++i)
{
    cudaMemcpyAsync(inputDevPtr + i*size, hostPtr + i*size, size, cudaMemcpyHostToDevice, stream[i]);
    MyKernel <<<100, 512, 0, stream[i]>>>(outputDevPtr + i*size, inputDevPtr + i*size, size);
    cudaMemcpyAsync(hostPtr + i*size, outputDevPtr + i*size, size, cudaMemcpyDeviceToHost, stream[i]);
}

對於 stream[0]、 stream[1] 來講,2次循環前一次中 stream[0]裏的 cudaMemcpyAsync DeviceToHost 和後一次循環中 stream[1]裏的 cudaMemcpyAsync HostToDevice 操做能夠重疊(併發), 固然這要求設備支持併發數據傳輸(Concurrent Data Transfer). 可是就上述代碼而言，即便設備支持併發Kernel執行(Concurrent Kernel and Kernel Execution),它也不太可能跳過兩次內存拷貝過程使 stream[0]和stream[1]的 kernel執行併發,因此是隱式同步(Implicit Synchronization).爲了充分利用 併發數據傳輸(Concurrent Data Transfer)和併發Kernel執行(Concurrent Kernel and Kernel Execution)這兩個特性,重寫代碼以下

for (int i = 0; i < 2; ++i)
    cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size,
    size, cudaMemcpyHostToDevice, stream[i]);
for (int i = 0; i < 2; ++i)
    MyKernel<<<100, 512, 0, stream[i]>>>
    (outputDevPtr + i * size, inputDevPtr + i * size, size);
for (int i = 0; i < 2; ++i)
    cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size,
    size, cudaMemcpyDeviceToHost, stream[i]);

上述代碼即便在設備不支持 kernel併發執行的狀況下：
stream[0]的 kernel執行和 stream[1]的 cudaMemcpyAsync HostToDevice內存拷貝能夠重疊, stream[0]的 cudaMemcpyAsync DeviceToHost內存操做和 stream[1]的kernel執行也能夠重疊.

上述代碼在設備支持 kernel併發及 data transfer併發的狀況下:
stream[0] 和 stream[1]中 cudaMemcpyAsync HostToDevice/DeviceToHost 併發 ,kernel 執行併發.

兩種方法比較之下後一種充分利用了設備的任務重疊併發特性(從一次增長到三次).即便設備不支持,也增長了一次重疊併發(從一次併發增長到兩次).

3.2.5.5.6 Callbacks

CUDA-runtime 提供了在stream中的函數回調.

使用方法:

1. cudaStreamAddCallback() 若是參數傳入 stream = 0 則表明等待全部在callback以前的 streams中指令完結以後函數回調.

下例添加 MyCallback函數回調至每一個 stream DeviceToHost內存拷貝操做以後:

void CUDART_CB MyCallback(cudaStream_t stream, cudaError_t status, void *data){
    printf("Inside callback %d\n", (size_t)data);
}
...
for (size_t i = 0; i < 2; ++i) {
    cudaMemcpyAsync(devPtrIn[i], hostPtr[i], size, cudaMemcpyHostToDevice,
    stream[i]);
    MyKernel<<<100, 512, 0, stream[i]>>>(devPtrOut[i], devPtrIn[i], size);
    cudaMemcpyAsync(hostPtr[i], devPtrOut[i], size, cudaMemcpyDeviceToHost,
    stream[i]);
    cudaStreamAddCallback(stream[i], MyCallback, (void*)i, 0);
}

cudaStreamAddCallback 函數最後一個參數爲 0 ,是CUDA保留爲了未來新功能的加入.

注意: 回調中絕對不能調用CUDA API(直接或間接), 這會致使自我調用的死循環.

3.2.5.5.7 Stream Priorities

設置 stream的優先級.

使用方法:

1. 在建立 stream時使用 cudaStreamCreateWithPriority() 函數
2. 使用 cudaDeviceGetStreamPriorityRange() 獲取可取優先級範圍 [ highest priority, lowest priority ]

例子:

// get the range of stream priorities for this device
int priority_high, priority_low;
cudaDeviceGetStreamPriorityRange(&priority_low, &priority_high);
// create streams with highest and lowest available priorities
cudaStream_t st_high, st_low;
cudaStreamCreateWithPriority(&st_high, cudaStreamNonBlocking, priority_high);
cudaStreamCreateWithPriority(&st_low, cudaStreamNonBlocking, priority_low);

3.2.5.6 Events

events 提供了能夠監控設備進程的方法.和回調同樣,它在特定的 stream中被觸發.
傳入參數 stream = 0 表示等待全部 stream 中的命令完成後觸發該事件.

3.2.5.6.1 Creation and Destruction

例子:

建立:
cudaEvent_t start, stop;
cudaEventCreat(&start);
cudaEventCreat(&stop);

銷燬:
cudaEventDestroy(start);
cudaEventDestroy(stop);

3.2.5.6.2 Elapsed Time

下例使用 event 記錄時間:

// 添加 start event 至全部 streams中
cudaEventRecord(start, 0);
for (int i = 0; i < 2; ++i) {
    cudaMemcpyAsync(inputDev + i * size, inputHost + i * size,
        size, cudaMemcpyHostToDevice, stream[i]);
    MyKernel<<<100, 512, 0, stream[i]>>>
        (outputDev + i * size, inputDev + i * size, size);
    cudaMemcpyAsync(outputHost + i * size, outputDev + i * size,
        size, cudaMemcpyDeviceToHost, stream[i]);
}
// 在全部命令添加完以後往 streams中添加 end event
cudaEventRecord(stop, 0);
// 同步等待全部 streams中命令完成後到達 stop event
cudaEventSynchronize(stop);
float elapsedTime;
// 記錄 start event 至 stop event的時間消耗
cudaEventElapsedTime(&elapsedTime, start, stop);

3.2.5.7 Synchronous Calls

當同步函數被調用以後, 直達全部相關命令執行結束後才返回控制權.使用 cudaSetDeviceFlags() 決定在同步結束後 host 線程行爲是 yield,block仍是spin.

3.2.6 Multi-Device System

3.2.6.1 Device Enumeration

一個 host 系統能夠擁有多個設備Device. 例子中遍歷設備並獲取他們的屬性.

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

3.2.6.2 Device Selection

一個 Host線程能夠在任什麼時候候使用 cudaSetDevice() 來指配設備進行運算.並切換全部執行環境.分配內存,kernel launch,streams,events等,都在最近指定的設備GPU上運行. 若是沒有指定則當前選擇設備號 = 0.

例子:
size_t size = 1024sizeof(float);
cudaSetDevice(0); //切換到設備0
float p0;
cudaMalloc(&p0, size); //在設備0 上分配global內存
MyKernel<<<1000, 128>>>(p0); //在設備0 上執行kernel函數
cudaSetDevice(1); //切換到設備 1
float * p1;
cudaMalloc(&p1, size); //在設備1 上分配global內存
MyKernel<<<1000, 128>>>(p1); //在設備1 上執行kernel函數

在多GPU設備的條件下,耗時的任務能夠指派給多個GPU進行運算.這是很好的.(SLI技術是多GPU完成單個任務,與這個不一樣)

3.2.6.3 Stream and Event Behavior

注意: kernel launch在 stream與當前 device沒有關聯的狀況下會失敗.
失敗例子:
cudaSetDevice(0); //切換到設備0
cudaStream_t s0;
cudaSreamCreate(&s0); //在當前設備0 中建立 stream s0
MyKernel<<<100,64,0,s0>>>(); //在當前設備0 中的 stream s0 中加入(異步) kernel launch指令

cudaSetDevice(1);                    //切換到設備1
cudaStream_t s1;
cudaSreamCreate(&s1);                //在當前設備1 中建立 stream s1
MyKernel<<<100,64,0,s1>>>();         //在當前設備1 中的 stream s1 中加入(異步) kernel launch指令

// 上述代碼是正確的
// 下面這行代碼會失敗
MyKernal<<<100,64,0,s0>>>();  #Error //在當前設備1 中試圖往設備0 中的 stream s0加入kernel launch指令

而內存拷貝指令卻與當前設備選擇無關:
// 下述代碼是正確的
cudaSetDevice(0); //切換到設備0
cudaStream_t s0;
cudaSreamCreate(&s0); //在當前設備0 中建立 stream s0

cudaSetDevice(1);                    //切換到設備1
cudaMemcpyAsync(devMemPtr, hostMemPtr, size, cudaMemcpyHostToDevice, s0); //This is OK

cudaEventRecord() 在 stream與當前 device沒有關聯的狀況下會失敗.
cudaEventElapsedTime() 在 stream與當前 device沒有關聯的狀況下會失敗.

cudaEventSynchronize() , cudaEventQuery() ,cudaStreamWaitEvent() 與當前設備選擇無關
所以 cudaStreamWaitEvent() 能夠在多個GPU設備之間作同步.

每一個設備擁有本身的默認 stream (see Default Stream).因此不一樣 GPU設備之間的任務執行是獨立無序的,你須要本身控制設備間的同步問題.

3.2.6.4 Peer-to-Peer Memory Access

應用程序若是在 64位處理器上執行的話,計算能力超過2.0的 Tesla系列顯卡能夠互相引用他們的內存地址(i.e. 一個kernel可使用另外一個設備內存地址中的數據來執行運算) 這個點對點的內存獲取特性可使用 cudaDeviceCanAccessPeer() = true檢查支持狀況.

點對點的內存獲取功能必須使用函數 cudaDeviceEnablePeerAccess() 開啓.每一個設備能夠支持全局最多 8個點的內存連接.

下例爲兩個設備之間的數據傳遞:
cudaSetDevice(0);
float p0;
size_t size = 1024sizeof(float);
cudaMalloc(&p0,size);
MyKernel<<<1000,128>>>(p0);
cudaSetDevice(1);
cudaDeviceEnablePeerAccess(0,0); //開啓對設備0 的點對點通道

// 在設備0 上launch kernel ,且該kernel使用設備0 中的地址 p0
MyKernel<<<1000,128>>>(p0);

3.2.6.5 Peer-to-Peer Memory Copy

兩個設備之間的點對點內存拷貝.
例子:
cudaSetDevice(0); // Set device 0 as current
float* p0;
size_t size = 1024 * sizeof(float);
cudaMalloc(&p0, size); // Allocate memory on device 0
cudaSetDevice(1); // Set device 1 as current
float* p1;
cudaMalloc(&p1, size); // Allocate memory on device 1
cudaSetDevice(0); // Set device 0 as current
MyKernel<<<1000, 128>>>(p0); // Launch kernel on device 0
cudaSetDevice(1); // Set device 1 as current
cudaMemcpyPeer(p1, 1, p0, 0, size); // Copy p0 to p1
MyKernel<<<1000, 128>>>(p1); // Launch kernel on device 1
A copy (in the implicit NULL stream) between the memories of two different devices

部分掠過詳細請查閱CUDA7.5 toolkit Documentation

3.2.7 Unified Virtual Address Space

當程序運行在 64位處理器上時, 一個64位的內存地址能夠供全部2.0以上設備和host所使用. 全部使用 CUDA API分配的 host 內存和全部 device 內存都在這個虛擬地址範圍內.(換句話說64位處理器提供的指針地址範圍夠大了),咱們稱爲這個虛擬地址爲通用的(unified).
咱們稱它爲通用虛擬地址是由於它並不表明真實的內存地址,而是一個虛擬地址到真實地址的內存地址映射(真實的內存地址是malloc出來的內存地址),爲了編程方便咱們須要多個設備和host統一使用同一個內存地址規範,而通用虛擬地址解決了這個問題.

1. 使用 cudaPointerGetAttributes() 來判斷是否內存地址是否使用了通用虛擬地址技術.
2. 當從通用地址中讀寫值的時候 cudaMemcpy() 函數的參數cudaMemcpyKind 應設爲flag cudaMemcpyDefault. 而且只要當前設備使用了通用地址,那麼即便 host 的內存不是從CUDA API中分配的,一樣也可使用(malloc/new).
3. 經過 cudaHostAlloc() 函數分配的 host 內存直接就是使用通用地址的 page-locked 內存塊(可供GPU直接讀取Host內存),因此也無需使用cudaHostGetDevicePointer()來獲取設備內存指針了.

優勢:

1. 使用cudaHostAlloc 分配的 page-locked 內存塊將自動提高 cudaMemcpy 等拷貝函數的帶寬和速度,別忘了以 cudaFreeHost 釋放.
2. 由於是 page-locked 因此GPU設備可直接讀取內容.

缺點:

1. 過多分配將下降應用程序可以使用內存,因此大多用來進行CPU和GPU之間的內存傳遞.

可使用設備屬性 unifiedAddressing = 1查看設備是否使用了通用內存地址.

3.2.8 Interprocess Communication

全部由 host線程建立分配的 Device內存指針或者 event handle 均可以在程序進程中全部的線程使用,但不能跨進程.
若是想要跨進程使用指針和事件,必須使用 InterProcess Communication API.詳細可查閱 reference manual. 而且該功能只在64位 Linux系統上受到支持.(部份內容略)

3.2.9 Error Checking

全部 run-time 函數均返回 error code.但對於異步併發(Async)函數來講,返回錯誤是不可能的(基於一些緣由).因此必須使用一些 host run-time 函數來獲得相關錯誤.

檢查異步錯誤的惟一方法是使用對應同步函數. 使用 cudaDeviceSynchronize() 函數來同步設備已得到在設備上發生的異步錯誤.
你也可使用不一樣級別的同步函數,好比cudaStreamSynchronize(), cudaStreamWaitEvent(), __syncthreads()等.
通常 run-time函數返回 cudaSuccess做爲異常指示標誌.

1. cudaPeekAtLastError() 用來獲取錯誤
2. cudaGetLastError() 獲取到錯誤後重置 last error = cudaSuccess.

kernel launch並不像其餘 run-time函數那樣返回錯誤標識,因此必須使用上述兩種方法獲取錯誤. 而且這兩個函數必須緊跟 kernel launch函數,來得到 pre-launch errors. 由於全局只有一個Error，而咱們不但願當中有任何函數引發的 Error 覆蓋了它.爲了保險起見，在 kernel launch以前也使用 cudaGetLastError()來獲取以前的異常並重置爲 cudaSuccess.
注意： cudaStreamQuery() 和 cudaEventQuery() 可能返回 cudaErrorNotReady ,它並不被認爲是一種異常錯誤,因此不會被上述方法所捕捉到.

3.2.10 Call Stack

在計算能力超過2.0的設備上可使用 cudaDeviceGetLimit(), cudaDeviceSetLimit() 查詢和設置調用棧的大小.
當棧溢出的時候, kernel call會失敗並返回一個棧溢出錯誤.
數據採集自GeForce-GTX760:
cudaLimitStackSize: 1024 bytes cudaLimitPrintfFifoSize: 1048576 bytes cudaLimitMallocHeapSize: 8388608 bytes cudaLimitDevRuntimeSyncDepth: 8388608 cudaLimitDevRuntimePendingLaunchCount: 8388608

3.2.11 Texture and Surface Memory

CUDA支持一些具備 texturing功能(Tesla系列就沒有)的GPU設備使用 texture 和 surface內存. 從texture 或者 surface內存中讀取數據比從 global內存中讀取有的優點在於如下幾點：

1. texture 和 surface內存爲讀取二維數據所優化,因此在讀取二維數據上能提供更高的帶寬速度
2. 地址計算由專門的計算單元進行,而無須放在 kernel中去處理.
3. 打包的數據能夠用一條指令操做來賦值給多個變量.相似於SIMD
4. 8-bit 和 16-bit 的 integer input data 能夠選擇性的轉換成 32-bit 的 floating-point value 於範圍[0.0, 1.0] or [-1.0, 1.0]內.(一般這個功能在計算圖片的顏色或灰度時十分受用)