CUDA 4.0 編程實踐

拿到CUDA 4.0版本與之前的CUDA 2.3有很大的不同。至少，cubin格式變成了ELF檔案，不再是decuda的輸入檔案。用GPU-Z測試了一下我的顯卡，GT218 支援OpenCL，CUDA，DirectCompute4.1。好了，一切都安裝完畢，包括vs2008。

下面是一個簡單的CUDA式的Hello World。

/*************************************************************************  [!output PROJECT_NAME].cu *  This is a example of the CUDA program. ************************************************************************/#include <stdio.h>#include <stdlib.h>#include <cuda_runtime.h>/************************************************************************//* Example                                                              *//************************************************************************/__global__ static void HelloCUDA(char* result, int num, clock_t* time){    int i = 0;      char p_HelloCUDA[] = "Hello CUDA!";    clock_t start = clock();    for(i = 0; i < num; i++) {        result[i] = p_HelloCUDA[i];    }      *time = clock() - start;}int main(int argc, char** argv){    char        *device_result  = 0;    clock_t     *time           = 0;    char        host_result[12] ={0};    clock_t     time_used       = 0;    int deviceCount;    int device;        cudaGetDeviceCount(&deviceCount);    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);    }    cudaMalloc((void**) &device_result, sizeof(char) * 11);    cudaMalloc((void**) &time, sizeof(clock_t));    HelloCUDA<<<1, 1, 0>>>(device_result, 11 , time);    cudaMemcpy(&host_result, device_result, sizeof(char) * 11, cudaMemcpyDeviceToHost);    cudaMemcpy(&time_used, time, sizeof(clock_t), cudaMemcpyDeviceToHost);    cudaFree(device_result);    cudaFree(time);    printf("%s,%d\n", host_result, time_used);    return 0;}

用命令列來編譯，相對簡單許多，nvcc.exe --help >nvcc.txt得到協助文檔，方便查看。對應上面的程式，批次檔如下(儲存為make.bat雙擊運行)：

@echo off  set myFun=sample   call "%VS90COMNTOOLS%vsvars32.bat"set include=%CUDA_INC_PATH%;%include%set lib=%CUDA_LIB_PATH%;%lib%set path=%CUDA_BIN_PATH%;%path%echo ------------------===By GoldenSpider 2011-10-8===------------------ nvcc %myFun%.cu -c  -Xcompiler "/MD " -o "%myFun%.obj"  link /OUT:"%myFun%.exe" /SUBSYSTEM:console /nologo %myFun%.obj cudart.lib kernel32.lib msvcrt.libecho -------------------------------------------------------------------echo Good Job, Compiler Success!! Run EXE(Y/?)  pause  %myFun%.exe    pause  

效果：(好像不支援上傳圖片，就複製一下cmd下的結果吧^_^)

Setting environment for using Microsoft Visual Studio 2008 x86 tools.
------------------===By GoldenSpider 2011-10-8===------------------
sample.cu
tmpxft_00000cf0_00000000-3_sample.cudafe1.gpu
tmpxft_00000cf0_00000000-8_sample.cudafe2.gpu
sample.cu
tmpxft_00000cf0_00000000-3_sample.cudafe1.cpp
tmpxft_00000cf0_00000000-14_sample.ii
-------------------------------------------------------------------
Good Job, Compiler Success!! Run EXE(Y/?)
請按任意鍵繼續. . .
Device 0 has compute capability 1.2 .
Hello CUDA!,8876
請按任意鍵繼續. . .

上面是基本入門，如果想用vc6.0編譯怎麼辦呢，要是用彙編該怎麼寫呢，思路也很簡單，就是用CUDA Driver API.裝置碼交給nvcc編譯，得到ptx或cubin。主機碼交給vc編譯或彙編器來編譯。ptx、cubin僅僅作為資料。實質上也是這麼做的。具體的可以參考vectorAddDrv這個執行個體。你能：

call "%VS90COMNTOOLS%vsvars32.bat"set include=%CUDA_INC_PATH%;%include%set lib=%CUDA_LIB_PATH%;%lib%set path=%CUDA_BIN_PATH%;%path%nvcc -ptx  VecAdd.cu

再:

@echo off call "E:\Microsoft Visual Studio\VC98\Bin\vcvars32.bat"set include=%CUDA_INC_PATH%;%include%set lib=%CUDA_LIB_PATH%;%lib%set myHost=maincl /c /MD %myHost%.cpplink  /SUBSYSTEM:console /nologo %myHost%.obj cuda.lib kernel32.lib msvcrt.lib%myHost%.exepause

執行效果：

cuDeviceGet returns: 0cuCtxCreate returns: 0cuModuleLoad returns: 0allocating d_a returns: 0copy data for a returns: 0getting the function handle returns: 0kernel launch returns: 0copy from device to host returns: 02.1000  ....

查看其匯入庫:

匯入表所處的節： .rdata
----------------------------------------------------------
匯入庫： nvcuda.dll
----------------------------------------------------------
OriginalFirstThunk 000020FC
TimeDateStamp 00000000
ForwarderChain 00000000
FirstThunk 00002044
----------------------------------------------------------
匯入序號 匯入函數名稱
----------------------------------------------------------
00000084 cuInit
00000059 cuDeviceGetCount
00000057 cuDeviceGet
0000000D cuCtxCreate_v2
000000E0 cuModuleLoad
0000008E cuMemAlloc_v2
000000C6 cuMemcpyHtoD_v2
000000DB cuModuleGetFunction
00000088 cuLaunchKernel
000000BE cuMemcpyDtoH_v2

----------------------------------------------------------
匯入庫： MSVCRT.dll
----------------------------------------------------------
OriginalFirstThunk 000020B8
TimeDateStamp 00000000
ForwarderChain 00000000
FirstThunk 00002000

運行庫已經不再是MSVCR90.dll ，呵呵。上面的VecAdd.cu代碼如下：

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