CUDA知識點
- CUDA裝置屬性
- 記憶體讀寫
- 修飾符
-
- \_\_global\_\_
- \_\_device\_\_
-
- 修飾函數
- 修飾變量
- \_\_constant\_\_
- \_\_shared\_\_
- 并行程式設計
-
- 樣例代碼
- 内置變量
- 用事件測量性能
- 紋理記憶體
-
- 簡介
- 使用
-
- 一維紋理記憶體
- 二維紋理記憶體
- 原子性
-
- 計算功能集
- 原子操作
- 頁鎖定主機記憶體
- CUDA流
- 零拷貝主機記憶體
-
- 簡介
- 使用
- 零拷貝記憶體的性能
- 使用多個GPU
-
- 樣例代碼
- 可移動的頁鎖定記憶體
CUDA裝置屬性
struct cudaDeviceProp
{
char name[256]; /**< 裝置的ASCII辨別 */
size_t totalGlobalMem; /**< 可用的全局記憶體量，機關位元組 */
size_t sharedMemPerBlock; /**< 每個block可用的共享記憶體量，機關位元組 */
int regsPerBlock; /**< 每個block裡可用32位寄存器數量 */
int warpSize; /**< 線程束大小*/
size_t memPitch; /**< 允許的記憶體複制最大修正，機關位元組*/
int maxThreadsPerBlock; /**< 每個block最大線程數量 */
int maxThreadsDim[3]; /**< 每個block裡每個次元最大線程量 */
int maxGridSize[3]; /**< 一格裡每個次元最大數量 */
int clockRate; /**< 時鐘頻率，機關千赫khz */
size_t totalConstMem; /**< 裝置上可用的常量記憶體，機關位元組 */
int major; /**< 計算功能主版本号*/
int minor; /**< 計算功能次版本号*/
size_t textureAlignment; /**< 對齊要求的紋理 */
int deviceOverlap; /**< 判斷裝置是否可以同時拷貝記憶體和執行核心。已過時。改用asyncEngineCount */
int multiProcessorCount; /**< 裝置上的處理器數量 */
int kernelExecTimeoutEnabled; /**< 核心函數是否運作受時間限制*/
int integrated; /**< 裝置是不是獨立的 */
int canMapHostMemory; /**< 裝置能否映射主機cudaHostAlloc/cudaHostGetDevicePointer */
int computeMode; /**< 計算模式，有預設，獨占，禁止，獨占程序(See ::cudaComputeMode) */
int maxTexture1D; /**< 1D紋理最大值 */
int maxTexture2D[2]; /**< 2D紋理最大維數*/
int maxTexture3D[3]; /**< 3D紋理最大維數 */
int maxTexture1DLayered[2]; /**< 最大的1D分層紋理尺寸 */
int maxTexture2DLayered[3]; /**< 最大的2D分層紋理尺寸 */
size_t surfaceAlignment; /**< 表面的對齊要求*/
int concurrentKernels; /**< 裝置是否能同時執行多個核心*/
int ECCEnabled; /**< 裝置是否支援ECC */
int pciBusID; /**< 裝置的PCI總線ID */
int pciDeviceID; /**< PCI裝置的裝置ID*/
int pciDomainID; /**<PCI裝置的域ID*/
int tccDriver; /**< 如果裝置是使用了TCC驅動的Tesla裝置則為1，否則就是0 */
int asyncEngineCount; /**< 異步Engine數量 */
int unifiedAddressing; /**< 裝置是否與主機共享統一的位址空間*/
int memoryClockRate; /**<峰值記憶體時鐘頻率，機關khz*/
int memoryBusWidth; /**< 全局記憶體總線寬度，機關bit*/
int l2CacheSize; /**< L2 cache大小，機關位元組 */
int maxThreadsPerMultiProcessor;/**< 每個多處理器的最大的常駐線程 */
};
#include "stdio.h"
#include <cuda_runtime.h>
int main(){
cudaDeviceProp prop;
int count;
cudaGetDeviceCount(&count);
for(int i=0;i<count;i++){
cudaGetDeviceProperties(&prop,i);
printf("Name: %s\n",prop.name);
}
return 0;
}
記憶體讀寫
- 在主機代碼中調用cudaMalloc配置設定裝置記憶體，并可以将指向裝置記憶體的指針傳遞給裝置函數或者主機函數，但不可在主機代碼中通路裝置記憶體
- 主機代碼中隻能通路主機記憶體（堆，棧）
- 要通路已經配置設定的裝置記憶體，隻有将裝置記憶體的指針傳遞給裝置函數，在裝置函數上進行通路
- 裝置函數中無法通路主機記憶體，是以不能給裝置函數傳遞指向主機記憶體的指針，但是可以直接值傳遞數值參數
- 裝置函數中無法調用主機函數，但是計算功能集大于等于2.0支援在裝置函數中調用printf函數
- 要使用裝置函數處理主機記憶體中的資料，可使用cudaMemcpy函數将主機記憶體中的資料拷貝到裝置記憶體進行處理，處理好之後再拷貝回主機記憶體
- 使用cudaMalloc和malloc配置設定的記憶體一定要用cudaFree和free進行釋放
修飾符
__global__
- 函數僅可在裝置上執行，僅可被主機函數調用
- 函數不支援遞歸
- 函數體内無法聲明靜态變量
- 函數參數數量不可變
- 函數傳回值必須為void
- 函數的調用是異步的，也就是說它會在裝置執行完成之前傳回
- 函數執行後需要調用函數cudaDeviceSynchronize進行同步，函數cudaMemcpy會隐式自動同步
- 函數參數将同時通過共享存儲器傳遞給裝置，且限制為 256 位元組（不了解）
__device__
修飾函數
- 函數僅可在裝置上執行，僅可被裝置函數調用
- 函數不支援遞歸
- 函數體内無法聲明靜态變量
- 函數參數數量不可變
- 函數的位址無法擷取
修飾變量
- 變量位于全局存儲器空間中，與應用程式具有相同的生命周期
- 變量可被所有線程通路
- 變量值可通過函數cudaMemcpyToSymbol和cudaMemcpyFromSymbol進行拷貝傳遞
__constant__
- 變量位于固定存儲器空間中，與應用程式具有相同的生命周期
- 變量可被所有線程通路，通路權限為隻讀，是以不能在裝置函數中進行初始化，隻能在主機函數中進行初始化
- 變量值可通過函數cudaMemcpyToSymbol進行初始化，初始化後不可改變，用函數cudaMemcpyFromSymbol進行拷貝
- 性能提升原因：
- 線程束指一個包含32個線程的集合，每個線程集合步調一緻地執行，線程束中的每個線程都将在不同的資料上執行相同的指令
- 硬體能夠将單次讀取常量記憶體的值廣播到半個線程束，即16個線程。如果半個線程束都讀取相同位址的資料，則隻要一次讀操作，再将資料廣播到其他”鄰近“線程，進而節約15次讀操作
- 常量記憶體的資料将緩存起來，對相同位址的連續讀操作不會産生額外的記憶體通信量
__shared__
- 變量位于線程塊的共享存儲器空間中
- 變量與塊具有相同的生命周期
- 變量可被塊内的所有線程通路
- 通路共享記憶體的延遲要遠低于通路普通緩沖區的延遲
- 使用函數__syncthreads()保證對一個線程塊中所有線程對共享數組的寫入操作在讀取之前完成
并行程式設計
樣例代碼
#define DIM 128
__global__ kernel(const float *a, float *b){
int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
int offset = x + y * blockDim.x * gridDim.x;
//......
}
int main(){
dim3 blocks(DIM/16, DIM/16);
dim3 threads(16, 16);
floag *a, *b;
CHECK(cudaMalloc((void **)&a, DIM * sizeof(float)));
CHECK(cudaMalloc((void **)&b, DIM * sizeof(float)));
kernel<<<blocks, threads>>>(a, b);
//......
CHECK(cudaFree(a));
CHECK(cudaFree(b));
}
内置變量
- threadIdx.x: 線程塊中的x方向的線程号
- blockIdx.x: 線程格中的x方向的線程塊号
- blockDim.x: 線程塊中的x方向的線程的數量
- gridDim.x: 線程格中的x方向的線程塊的數量
- y同理
- 不同線程根據各自不同的内置變量值獲得各自需要處理的資料位置，進而實作對一個數組所有元素的并行處理
用事件測量性能
事件的本質是一個GPU時間戳
cudaEvent_t start, stop;
CHECK(cudaEventCreate(&start));
CHECK(cudaEventCreate(&stop));
CHECK(cudaEventRecord(start, 0));
//......
CHECK(cudaEventRecord(stop, 0));
CHECK(cudaEventSynchronize(stop));
float time_diff;
CHECK(cudaEventElapsedTime(&time_diff, start, stop));
CHECK(cudaEventDestroy(start));
CHECK(cudaEventDestroy(stop));
cudaEventSynchronize(stop)的作用：確定stop事件之前的所有GPU工作已經完成
紋理記憶體
簡介
紋理記憶體是一種隻讀記憶體，能夠為記憶體通路存在大量空間局部性的程式減少對記憶體的請求并提供更高效的記憶體帶寬。
空間局部性：同一個線程或鄰近多個線程讀取資料的位址相近。
使用
一維紋理記憶體
//主機函數中初始化紋理記憶體
//......
texture<float> tex;
CHECK(cudaMalloc((void **)&a, sizeof(float) * N));
CHECK(cudaBindTexture(NULL, tex, a, N));
//......
//裝置函數中對紋理記憶體進行讀取
//......
float t = tex1Dfetch(tex, index);
//......
//在主機函數中釋放紋理記憶體
//......
cudaUnbindTexture(tex);
//......
二維紋理記憶體
//主機函數中初始化紋理記憶體
//......
texture<float> tex;
CHECK(cudaMalloc((void **)&a, sizeof(float) * N));
cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>(); //通道格式描述符
CHECK(cudaBindTexture(NULL, tex, a, desc, DIM, DIM, sizeof(float) * DIM));
//......
//裝置函數中對紋理記憶體進行讀取
//......
float t = tex2Dfetch(tex, x, y);
//......
//在主機函數中釋放紋理記憶體
//......
cudaUnbindTexture(tex);
//......
原子性
計算功能集
NVIDIA将GPU支援的各種功能統稱為計算功能集
編譯指定計算功能集不能低于某一個版本：
nvcc -arch=sm_12
原子操作
原子性：一次性對某個記憶體空間進行讀寫操作，在執行過程中不會被其他線程中斷
atomicAdd(&a, 1); //+1
注意：當數千個線程嘗試通路少量記憶體時，将發生大量競争，為保持原子性需要付出大量開銷，是以會降低性能
解決措施：線上程塊内設定一個共享記憶體，隻讓塊内的多個線程競争，将結果暫時存在共享記憶體中，最後将所有線程塊中共享記憶體中的結果綜合到全局記憶體中
頁鎖定主機記憶體
malloc()：配置設定可分頁的主機記憶體
cudaHostAlloc()：配置設定不可分頁的主機記憶體
不可分頁主機記憶體：作業系統不會對這塊記憶體分頁并交換到磁盤上，確定該記憶體始終駐留在實體記憶體中
用cudaHostAlloc()提升性能的原因：GPU知道記憶體的實體位址，可以通過直接記憶體通路(DMA)技術在GPU和主機之間複制資料，無需CPU介入。當記憶體為分頁記憶體時，CPU可能會在DMA執行過程中将目标記憶體交換到磁盤上，或通過更新作業系統的可分頁表來 重新定位目标記憶體的實體位址，進而對DMA操作造成延時。而不可分頁記憶體的使用不會造成延時，提高了性能。另一方面，在可分頁記憶體資料拷貝到GPU的過程中，複制操作會執行兩次，第一次是将可分頁記憶體複制到一塊臨時的頁鎖定記憶體，再從這個頁鎖定記憶體複制到GPU上，兩次複制增大了開銷
注意：過多不可分頁記憶體的使用會使主機記憶體耗盡，影響其他程式運作
float *a;
CHECK(cudaHostAlloc((void **)&a, sizeof(float) * N));
//......
CHECK(cudaFreeHost(a));
CUDA流
cudaMemcpy()：同步複制，函數傳回時複制操作已經完成
cudaMemcpyAsync()：異步複制，函數傳回時複制操作不一定完成，結合stream使用
stream的作用：在主機函數有序調用多個異步裝置函數時，保證這些異步裝置函數能夠有序執行
cudaStreamSynchronize(stream)：保證異步方式在stream流中執行的裝置函數全部完成
stream0中複制C的操作要等待核函數執行完成，是以阻塞了stream1中複制A的操作
改進：使用寬度優先方式将操作放入各個流的隊列，加速運作速度
零拷貝主機記憶體
簡介
在cudaHostAlloc()函數中傳入cudaHostAllocMapped參數，能夠配置設定得到在裝置函數中直接通路的頁鎖定主機記憶體，進而可以免去主機和裝置之間的資料拷貝
使用
CHECK(cudaSetDeviceFlags(cudaDeviceMapHost)); //設定為裝置映射主機記憶體
float *a, *dev_a;
CHECK(cudaHostAlloc((void**)&a, size * sizeof(float), cudaHostALlocWriteCombined | cudaHostAllocMapped)); //配置設定可在裝置函數通路的頁鎖定主機記憶體
CHECK(cudaHostGetDevicePointer(&dev_a, a, 0)); //獲得在GPU上的有效指針
//......
CHECK(cudaFreeHost(a));
标志cudaHostAllocWriteCombined：将記憶體配置設定為“合并式寫入”記憶體，主要用于主機到裝置的傳輸或者通過映射頁鎖定空間CPU寫而裝置讀的情況，可以顯著提升GPU讀取記憶體的性能，但CPU讀取該記憶體時會比較低效
标志cudaHostAllocMapped：将主機記憶體配置設定為GPU可通路的記憶體
零拷貝記憶體的性能
內建GPU：裝置記憶體和主機記憶體在實體上共享，是以使用零拷貝記憶體可避免不必要的資料拷貝，提升性能
獨立GPU：當程式滿足“僅讀取/寫入一次”這個限制條件時，在獨立GPU上使用零拷貝記憶體可以獲得性能提升。但是由于GPU不會緩存零拷貝記憶體，當裝置函數需要多次讀寫記憶體時，會降低性能，還不如把資料複制到GPU上
注意：由于零拷貝記憶體時頁鎖定記憶體，是以申請過多零拷貝記憶體會減少可使用的實體記憶體，影響其他程式的運作
使用多個GPU
樣例代碼
/*
* Copyright 1993-2010 NVIDIA Corporation. All rights reserved.
*
* NVIDIA Corporation and its licensors retain all intellectual property and
* proprietary rights in and to this software and related documentation.
* Any use, reproduction, disclosure, or distribution of this software
* and related documentation without an express license agreement from
* NVIDIA Corporation is strictly prohibited.
*
* Please refer to the applicable NVIDIA end user license agreement (EULA)
* associated with this source code for terms and conditions that govern
* your use of this NVIDIA software.
*
*/
#include <book.h>
#define imin(a,b) (a<b?a:b)
#define N (33*1024*1024)
const int threadsPerBlock = 256;
const int blocksPerGrid =
imin( 32, (N/2+threadsPerBlock-1) / threadsPerBlock );
__global__ void dot( int size, float *a, float *b, float *c ) {
__shared__ float cache[threadsPerBlock];
int tid = threadIdx.x + blockIdx.x * blockDim.x;
int cacheIndex = threadIdx.x;
float temp = 0;
while (tid < size) {
temp += a[tid] * b[tid];
tid += blockDim.x * gridDim.x;
}
// set the cache values
cache[cacheIndex] = temp;
// synchronize threads in this block
__syncthreads();
// for reductions, threadsPerBlock must be a power of 2
// because of the following code
int i = blockDim.x/2;
while (i != 0) {
if (cacheIndex < i)
cache[cacheIndex] += cache[cacheIndex + i];
__syncthreads();
i /= 2;
}
if (cacheIndex == 0)
c[blockIdx.x] = cache[0];
}
struct DataStruct {
int deviceID;
int size;
float *a;
float *b;
float returnValue;
};
void* routine( void *pvoidData ) {
DataStruct *data = (DataStruct*)pvoidData;
HANDLE_ERROR( cudaSetDevice( data->deviceID ) ); //為每個線程指定執行的GPU
int size = data->size;
float *a, *b, c, *partial_c;
float *dev_a, *dev_b, *dev_partial_c;
// allocate memory on the CPU side
a = data->a;
b = data->b;
partial_c = (float*)malloc( blocksPerGrid*sizeof(float) );
// allocate the memory on the GPU
HANDLE_ERROR( cudaMalloc( (void**)&dev_a,
size*sizeof(float) ) );
HANDLE_ERROR( cudaMalloc( (void**)&dev_b,
size*sizeof(float) ) );
HANDLE_ERROR( cudaMalloc( (void**)&dev_partial_c,
blocksPerGrid*sizeof(float) ) );
// copy the arrays 'a' and 'b' to the GPU
HANDLE_ERROR( cudaMemcpy( dev_a, a, size*sizeof(float),
cudaMemcpyHostToDevice ) );
HANDLE_ERROR( cudaMemcpy( dev_b, b, size*sizeof(float),
cudaMemcpyHostToDevice ) );
dot<<<blocksPerGrid,threadsPerBlock>>>( size, dev_a, dev_b,
dev_partial_c );
// copy the array 'c' back from the GPU to the CPU
HANDLE_ERROR( cudaMemcpy( partial_c, dev_partial_c,
blocksPerGrid*sizeof(float),
cudaMemcpyDeviceToHost ) );
// finish up on the CPU side
c = 0;
for (int i=0; i<blocksPerGrid; i++) {
c += partial_c[i];
}
HANDLE_ERROR( cudaFree( dev_a ) );
HANDLE_ERROR( cudaFree( dev_b ) );
HANDLE_ERROR( cudaFree( dev_partial_c ) );
// free memory on the CPU side
free( partial_c );
data->returnValue = c;
return 0;
}
int main( void ) {
int deviceCount;
HANDLE_ERROR( cudaGetDeviceCount( &deviceCount ) );
if (deviceCount < 2) {
printf( "We need at least two compute 1.0 or greater "
"devices, but only found %d\n", deviceCount );
return 0;
}
float *a = (float*)malloc( sizeof(float) * N );
HANDLE_NULL( a );
float *b = (float*)malloc( sizeof(float) * N );
HANDLE_NULL( b );
// fill in the host memory with data
for (int i=0; i<N; i++) {
a[i] = i;
b[i] = i*2;
}
// prepare for multithread
DataStruct data[2];
data[0].deviceID = 0;
data[0].size = N/2;
data[0].a = a;
data[0].b = b;
data[1].deviceID = 1;
data[1].size = N/2;
data[1].a = a + N/2;
data[1].b = b + N/2;
CUTThread thread = start_thread( routine, &(data[0]) ); //建立一個新線程，在新線程中調用routine函數
routine( &(data[1]) ); //在主線程中調用routine函數
end_thread( thread ); //等待兩個線程的routine函數執行完成
// free memory on the CPU side
free( a );
free( b );
printf( "Value calculated: %f\n",
data[0].returnValue + data[1].returnValue );
return 0;
}
可移動的頁鎖定記憶體
在使用多個GPU時，cudaHostAlloc配置設定的頁鎖定記憶體隻對于配置設定它們的線程來說是頁鎖定的，而其他線程依然把這塊記憶體當作可分頁記憶體。當其他線程複制這塊記憶體資料時，會當作可分頁記憶體進行複制，降低性能
解決方案：cudaHostAlloc((void **)&a, N * sizeof(float), cudaHostAllocPortable);