CUDA学习笔记一

由于各种需求，所以略微了解学习一下gpu的CUDA怎么操作？CUDA简介CUDA是并行计算的平台和类C编程模型，我们能很容易的实现并行算法，就像写C代码一样。只要配备的NVIDIA GPU，就可以在许多设备上运行你的并行程序，无论是台式机、笔记本抑或平板电脑。熟悉C语言可以帮助你尽快掌握CUDA。CUDA编程CUDA编程允许你的程序执行在异构系统上，即CUP和GPU，二者有各自的存储空间，并由PCI-Express 总线区分开。因此，我们应该先注意二者术语上的区分：Host：CPU and itsmemory (host memory)Device: GPU and its memory (device memory)代码中，一般用h_前缀表示host memory，d_表示device memory。kernel是CUDA编程中的关键，他是跑在GPU的代码，用标示符__global__注明。host可以独立于host进行大部分操作。当一个kernel启动后，控制权会立刻返还给CPU来执行其他额外的任务。所以，CUDA编程是异步的。一个典型的CUDA程序包含由并行代码补足的串行代码，串行代码由host执行，并行代码在device中执行。host端代码是标准C，device是CUDA C代码。我们可以把所有代码放到一个单独的源文件，也可以使用多个文件或库。NVIDIA C编译器（nvcc）可以编译host和device生成可执行程序。这里再次说明下CUDA程序的处理流程：从CPU拷贝数据到GPU。调用kernel来操作存储在GPU的数据。将操作结果从GPU拷贝至CPU。Memory操作cuda程序将系统区分成host和device，二者有各自的memory。kernel可以操作device memory，为了能很好的控制device端内存，CUDA提供了几个内存操作函数： 为了保证和易于学习，CUDA C 的风格跟C很接近，比如：cudaError_t cudaMalloc ( void** devPtr, size_t size )我们主要看看cudaMencpy，其函数原型为：cudaError_t cudaMemcpy ( void* dst, const void* src, size_t count,cudaMemcpyKind kind )其中cudaMemcpykind的可选类型有：cudaMemcpyHostToHostcudaMemcpyHossToDevicecudaMemcpyDeviceToHostcudaMemcpuDeviceToDevice具体含义很好懂，就不多做解释了。对于返回类型cudaError_t，如果正确调用，则返回cudaSuccess，否则返回cudaErrorMemoryAllocation。可以使用char* cudaGetErrorString(cudaError_t error)将其转化为易于理解的格式。组织线程掌握如何组织线程是CUDA编程的重要部分。CUDA线程分成Grid和Block两个层次。 由一个单独的kernel启动的所有线程组成一个grid，grid中所有线程共享global memory。一个grid由许多block组成，block由许多线程组成，grid和block都可以是一维二维或者三维，上图是一个二维grid和二维block。这里介绍几个CUDA内置变量：blockIdx：block的索引，blockIdx.x表示block的x坐标。threadIdx：线程索引，同理blockIdx。blockDim：block维度，上图中blockDim.x=5.gridDim：grid维度，同理blockDim。一般会把grid组织成2D，block为3D。grid和block都使用dim3作为声明，例如：dim3 block(3);

// 后续博文会解释为何这样写grid

dim3 grid((nElem+block.x-1)/block.x);需要注意的是，dim3仅为host端可见，其对应的device端类型为uint3。启动CUDA kernelCUDA kernel的调用格式为：kernel_name<<>>(argument list);其中grid和block即为上文中介绍的类型为dim3的变量。通过这两个变量可以配置一个kernel的线程总和，以及线程的组织形式。例如：kernel_name<<<4, 8>>>(argumentt list);该行代码表明有grid为一维，有4个block，block为一维，每个block有8个线程，故此共有4*8=32个线程。注意，不同于c函数的调用，所有CUDA kernel的启动都是异步的，当CUDA kernel被调用时，控制权会立即返回给CPU。函数类型标示符__device__ 和__host__可以组合使用。kernel的限制：仅能获取device memory 。必须返回void类型。不支持可变数目参数。不支持静态变量。不支持函数指针。异步。代码分析#include

#include

#define CHECK(call) \

{ 　　const cudaError_t error = call; 　　if (error != cudaSuccess) 　　{ 　　　　printf("Error: %s:%d, ", __FILE__, __LINE__); 　　　　printf("code:%d, reason: %s\n", error, cudaGetErrorString(error)); 　　　　exit(1); 　　} }

void checkResult(float *hostRef, float *gpuRef, const int N) {

　　double epsilon = 1.0E-8;

　　bool match = 1;

　　for (int i=0; i

　　　　if (abs(hostRef[i] - gpuRef[i]) > epsilon) {

　　　　　　match = 0;

　　　　　　printf("s do not match!\n");

　　　　　　printf("host %5.2f gpu %5.2f at current %d\n",hostRef[i],gpuRef[i],i);

　　　　　　break;

　　　　}

　　}

　　if (match) printf("s match.\n\n");

}

void initialData(float *ip,int size) {

　　// generate different seed for random number

　　time_t t;

　　srand((unsigned) time(&t));

　　for (int i=0; i

　　　　ip[i] = (float)( rand() & 0xFF )/10.0f;

　　}

}

void sumsOnHost(float *A, float *B, float *C, const int N) {

　　for (int idx=0; idx

　　C[idx] = A[idx] + B[idx];

}

__global__ void sumsOnGPU(float *A, float *B, float *C) {

　　int i = threadIdx.x;

　　C[i] = A[i] + B[i];

}

int main(int argc, char **argv) {

　　printf("%s Starting...\n", argv[0]);

　　// set up device

　　int dev = 0;

　　cudaSetDevice(dev);

　　// set up data size of vectors

　　int nElem = 32;

　　printf("Vector size %d\n", nElem);

　　// malloc host memory

　　size_t nBytes = nElem * sizeof(float);

　　float *h_A, *h_B, *hostRef, *gpuRef;

　　h_A = (float *)malloc(nBytes);

　　h_B = (float *)malloc(nBytes);

　　hostRef = (float *)malloc(nBytes);

　　gpuRef = (float *)malloc(nBytes);

　　// initialize data at host side

　　initialData(h_A, nElem);

　　initialData(h_B, nElem);

　　memset(hostRef, 0, nBytes);

　　memset(gpuRef, 0, nBytes);

　　// malloc device global memory

　　float *d_A, *d_B, *d_C;

　　cudaMalloc((float**)&d_A, nBytes);

　　cudaMalloc((float**)&d_B, nBytes);

　　cudaMalloc((float**)&d_C, nBytes);

　　// transfer data from host to device

　　cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);

　　cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice);

　　// invoke kernel at host side

　　dim3 block (nElem);

　　dim3 grid (nElem/block.x);

　　sumsOnGPU<<< grid, block >>>(d_A, d_B, d_C);

　　printf("Execution configuration <<<%d, %d>>>\n",grid.x,block.x);

　　// copy kernel result back to host side

　　cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);

　　// add vector at host side for result checks

　　sumsOnHost(h_A, h_B, hostRef, nElem);

　　// check device results

　　checkResult(hostRef, gpuRef, nElem);

　　// free device global memory

　　cudaFree(d_A);

　　cudaFree(d_B);

　　cudaFree(d_C);

　　// free host memory

　　free(h_A);

　　free(h_B);

　　free(hostRef);

　　free(gpuRef);

　　return(0);

}编译指令：$nvcc sum.cu -o sum运行： $./sum输出：./sum Starting...

Vector size 32

Execution configuration <<<1, 32>>>

s match.

CUDA学习笔记一