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 = 1024*sizeof(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]內.(通常這個功能在計算圖片的顏色或灰度時十分受用)