CUDA8.0矩陣乘法例子解釋（matrixMul.cpp）
通過學習英偉達自帶的例子matrixMul學CUDA庫的使用。
簡略部分垃圾。只說核心程式碼。
這個例子是實現 C=A*B的矩陣相乘
// Use a larger block size for Fermi and above int block_size = 32; //original: dim3 dimsA(5*2*block_size, 5*2*block_size, 1); dim3 dimsB(5*4*block_size, 5*2*block_size, 1); // reduce sizes to avoid running out of memory //dim3 dimsA(32,32, 1); //dim3 dimsB(32,32,1);
定義了矩陣的大小。 dim3是一個三維結構。xyz。分別代表長寬高。是cuda的內建結構
struct __device_builtin__ dim3
{
unsigned int x, y, z;
};矩陣A列5*2*32 （320） 行 5*2*32 （320）
矩陣B列5*4*32 （640）行 5*2*32（320）
下面說為了避免記憶體用完。可以減小一些寬和高。
隨後呼叫
int matrix_result = matrixMultiply(argc, argv, block_size, dimsA, dimsB);進入matrixMultiply以後的程式碼如下
unsigned int size_A = dimsA.x * dimsA.y; unsigned int mem_size_A = sizeof(float) * size_A; float *h_A = (float *)malloc(mem_size_A); unsigned int size_B = dimsB.x * dimsB.y; unsigned int mem_size_B = sizeof(float) * size_B; float *h_B = (float *)malloc(mem_size_B); // Initialize host memory const float valB = 0.01f; constantInit(h_A, size_A, 1.0f); constantInit(h_B, size_B, valB);
size_A和size_B分別是AB矩陣元素的個數。
mem_size_A和mem_size_B分別是矩陣所需要的記憶體的大小，這裡每個元素都是浮點數。
h_A和h_B則是分配好的記憶體的起始地址。
隨後對A B的資料進行初始化。A每個元素賦值為1.0f。B每個元素賦值為0.01f (constantInit就是實現賦值的過程)
// Allocate device memory
CUdeviceptr d_A, d_B, d_C;
char *ptx, *kernel_file;
size_t ptxSize;定義一些後續需要的變數
kernel_file = sdkFindFilePath("matrixMul_kernel.cu", argv[0]);
compileFileToPTX(kernel_file, 0, NULL, &ptx, &ptxSize);
CUmodule module = loadPTX(ptx, argc, argv);
找到cu檔案的位置。cu檔案是C語言語法的，就是字尾不同，這裡面主要是實現演算法。隨後呼叫
compile將cu檔案編譯成GPU可以理解執行的程式碼，然後通loadPTX執行載入函式。
就是將cu檔案編譯成GPU可以理解的東西。相當於翻譯的過程。然後載入進來。
這裡需要argc和argv是可能在argv中指定使用某個特殊的裝置。比如我有幾張顯示卡。可能要選擇這樣。否則就按照預設配置來選擇。
其次關於 compileFileToptX和loadPtx。是拿底層的SDK進行了基礎的封裝。你也可以呼叫底層SDK實現。具體的函式在nvrtc_helper.h檔案裡面。
// Allocate host matrix C
dim3 dimsC(dimsB.x, dimsA.y, 1);
unsigned int mem_size_C = dimsC.x * dimsC.y * sizeof(float);
float *h_C = (float *) malloc(mem_size_C);
定義計算後儲存的結果，不細說了。
checkCudaErrors(cuMemAlloc(&d_A, mem_size_A));
checkCudaErrors(cuMemAlloc(&d_B, mem_size_B));
checkCudaErrors(cuMemAlloc(&d_C, mem_size_C));
// copy host memory to device
checkCudaErrors(cuMemcpyHtoD(d_A, h_A, mem_size_A));
checkCudaErrors(cuMemcpyHtoD(d_B, h_B, mem_size_B));
cuMemAlloc這次是在視訊記憶體中分配記憶體了。分配了矩陣A B 和結果C。所以這裡要注意。千萬不要記憶體不夠了。因此儘量關閉不需要用到的視訊記憶體。
接下來吧矩陣A B 的記憶體資料（h_A h_B）拷貝到視訊記憶體(d_A d_B)當中去。
// Setup execution parameters
dim3 threads(block_size, block_size);
dim3 grid(dimsB.x / threads.x, dimsA.y / threads.y);
定義了執行引數3維 x y z
threads (x=block_size ，y=block_size ，z=1) threads(32,32,1)
grid (x=矩陣B的x/threads.x，y=矩陣A的y/thread.Y，z=1) grid (640/32=20 , 320/32=10,1 )=（20,10,1）
CUfunction kernel_addr;
if (block_size == 16)
{
checkCudaErrors(cuModuleGetFunction(&kernel_addr, module, "matrixMulCUDA_block16"));
}
else
{
checkCudaErrors(cuModuleGetFunction(&kernel_addr, module, "matrixMulCUDA_block32"));
}
通過cuModuleGetFunction得到cu模組中函式的地址，放在kernel_addr中。
void matrixMulCUDA_block32(float *C, float *A, float *B, int wA, int wB)
這個就是cu檔案中實際的執行函數了，這裡先看一下函式，有個映像。
void *arr[] = { (void *)&d_C, (void *)&d_A, (void *)&d_B, (void *)&dimsA.x, (void *)&dimsB.x };
接著定義了函式需要的引數。
int nIter = 300;
for (int j = 0; j < nIter; j++)
{
checkCudaErrors(cuLaunchKernel(kernel_addr,
grid.x, grid.y, grid.z, /* grid dim */
threads.x, threads.y, threads.z, /* block dim */
0,0, /* shared mem, stream */
&arr[0], /* arguments */
0));
checkCudaErrors(cuCtxSynchronize());
}
執行計算操作。
這裡重點介紹cuLaunchKernel函式
CUresult CUDAAPI cuLaunchKernel(CUfunction f,
unsigned int gridDimX,
unsigned int gridDimY,
unsigned int gridDimZ,
unsigned int blockDimX,
unsigned int blockDimY,
unsigned int blockDimZ,
unsigned int sharedMemBytes,
CUstream hStream,
void **kernelParams,
void **extra);
呼叫核心執行f函式。處理gridDimX gridDimY gridDimZ大小的資料塊。每個資料塊包含 blockDimX blockDImY blockDimZ執行緒
sharedMemBytes指定了每個資料塊可以共享的動態記憶體。
f函式的引數可以有兩種形式
1 通過kernelParams引數指定。如果f有N個引數。那麼kernelParams就是一個N大小的引數陣列的指標。從kernelParams[0]到
kernelParams[N-1]，每個引數必須指向核心將要拷貝的一塊記憶體，這裡意思是核心需要的是地址，而不是值。
比如f(int x) 需要的是一個int的x。 那麼kernelParams[0]=&x; 而不能直接指定為x。這點要特別注意。
核心引數的數量，大小和偏移不需要指定，那些都是直接從核心的image裡面直接得到，（這句暫未理解什麼意思）。
2引數也可以通過程式打包到一個單獨的buffer通過extra引數傳遞過去。這個就需要用對其等來處理好每個引數的大小等等。
extra的存在主要是允許 culaunchKernel函式拿到一些不通用的引數。extra指定了這些引數的名字和對應的值。必須以NULL或者CU_LAUNCH_PARAM_END結束
比如
void *config[] = {
CU_LAUNCH_PARAM_BUFFER_POINTER, argBuffer,
CU_LAUNCH_PARAM_BUFFER_SIZE, &argBufferSize,
CU_LAUNCH_PARAM_END
};
status = cuLaunchKernel(f, gx, gy, gz, bx, by, bz, sh, s, NULL, config);
並且 兩種方式只能選用一種。兩個值都指定的話會導致函式執行錯誤。
該cuLaunchKernel函式等同的呼叫函式順序如下
* Calling ::cuLaunchKernel() sets persistent function state that is
* the same as function state set through the following deprecated APIs:
* ::cuFuncSetBlockShape(),
* ::cuFuncSetSharedSize(),
* ::cuParamSetSize(),
* ::cuParamSeti(),
* ::cuParamSetf(),
* ::cuParamSetv().
通過cuLaunchKernel呼叫後，會覆蓋前面設定的塊大小。引數資訊。共享大小等等。
* \param f - Kernel to launch
* \param gridDimX - Width of grid in blocks
* \param gridDimY - Height of grid in blocks
* \param gridDimZ - Depth of grid in blocks
* \param blockDimX - X dimension of each thread block
* \param blockDimY - Y dimension of each thread block
* \param blockDimZ - Z dimension of each thread block
* \param sharedMemBytes - Dynamic shared-memory size per thread block in bytes
* \param hStream - Stream identifier
* \param kernelParams - Array of pointers to kernel parameters
* \param extra - Extra options
// Copy result from device to host
checkCudaErrors(cuMemcpyDtoH(h_C, d_C, mem_size_C));
將結果從視訊記憶體直接拷貝回記憶體
free(h_A);
free(h_B);
free(h_C);
checkCudaErrors(cuMemFree(d_A));
checkCudaErrors(cuMemFree(d_B));
checkCudaErrors(cuMemFree(d_C));
最後就是資源的釋放