cuda学习1-初始庐山真面目
cuda作为gpu计算中的代表，拥有着超级高的计算效率，其原因是gpu实际相当与一台超级并行机组，使用过MPI做并行计算的人们可能知道，所谓的并行计算，简单讲就是用多个U（计算单元）来完成一个U的计算任务，MPI中将其叫做核，我们知道一个cpu有一个或2，4，8个核，超级厉害的也就16个吧，原来人们为了做大规模的并行计算，将一大堆cpu装在柜子里，组成计算集群，但是那种设备大的吓人，而且又有多少人会用呢。gpu则不同，一个小小的芯片上就存在着成千上万的线程，是由分为grid，block，thread三级结构实现的。（本人初学者，学到哪写到哪，借鉴需谨慎）
所谓的并行计算，就是用多个计算单元共同完成一个计算任务，那这有什么难度呢，这在生活中很常见啊，似乎没什么可学的，可以想象我们造一座房子，我们要找一堆人过来，然后分配下任务，这里面有人做墙，有人做地板，屋顶等等吧，然后一声令下开始吧，大家各司其职，把自己那部分做好，房子就做好了。但是事情没有那么简单，这个过程中存在很多问题会影响效率，甚至结果。比如说做墙的人要用锤子，但是锤子在被其他人占用着，怎么办；做地板的把地做好了，但是又被做墙的人给不小心砸坏了；下面的结构还没做好，就有人来装屋顶了等等。这时就需要一个精确的管理方案，而这个管理方案就是所有并行API需要做的事。例如后面会学到的共享内存，同步等知识。目前看到书就是cuda by example了，很适合入门，看的很愉快。下面就开始了，什么hello world就不写了，直接学干货了。
1 /* 2 * Copyright 1993-2010 NVIDIA Corporation. All rights reserved. 3 * 4 * NVIDIA Corporation and its licensors retain all intellectual property and 5 * proprietary rights in and to this software and related documentation. 6 * Any use, reproduction, disclosure, or distribution of this software 7 * and related documentation without an express license agreement from 8 * NVIDIA Corporation is strictly prohibited. 9 * 10 * Please refer to the applicable NVIDIA end user license agreement (EULA) 11 * associated with this source code for terms and conditions that govern 12 * your use of this NVIDIA software. 13 * 14 */ 15 16 17 #include "../common/book.h" 18 19 #define N 10 20 21 __global__ void add( int *a, int *b, int *c ) { 22 int tid = blockIdx.x; // this thread handles the data at its thread id 23 if (tid < N) 24 c[tid] = a[tid] + b[tid]; 25 } 26 27 int main( void ) { 28 int a[N], b[N], c[N]; 29 int *dev_a, *dev_b, *dev_c; 30 31 // allocate the memory on the GPU 32 HANDLE_ERROR( cudaMalloc( (void**)&dev_a, N * sizeof(int) ) ); 33 HANDLE_ERROR( cudaMalloc( (void**)&dev_b, N * sizeof(int) ) ); 34 HANDLE_ERROR( cudaMalloc( (void**)&dev_c, N * sizeof(int) ) ); 35 36 // fill the arrays 'a' and 'b' on the CPU 37 for (int i=0; i<N; i++) { 38 a[i] = -i; 39 b[i] = i * i; 40 } 41 42 // copy the arrays 'a' and 'b' to the GPU 43 HANDLE_ERROR( cudaMemcpy( dev_a, a, N * sizeof(int), 44 cudaMemcpyHostToDevice ) ); 45 HANDLE_ERROR( cudaMemcpy( dev_b, b, N * sizeof(int), 46 cudaMemcpyHostToDevice ) ); 47 48 add<<<N,1>>>( dev_a, dev_b, dev_c ); 49 50 // copy the array 'c' back from the GPU to the CPU 51 HANDLE_ERROR( cudaMemcpy( c, dev_c, N * sizeof(int), 52 cudaMemcpyDeviceToHost ) ); 53 54 // display the results 55 for (int i=0; i<N; i++) { 56 printf( "%d + %d = %d\n", a[i], b[i], c[i] ); 57 } 58 59 // free the memory allocated on the GPU 60 HANDLE_ERROR( cudaFree( dev_a ) ); 61 HANDLE_ERROR( cudaFree( dev_b ) ); 62 HANDLE_ERROR( cudaFree( dev_c ) ); 63 64 return 0; 65 }
这段代码讲述了如何将两个长度是10的向量相加，属于gpu计算中基础中的基础，我们借助这段毫无难度的代码熟悉一下cuda中的一些基本规则。
首先，每一个cuda代码中必有kernel函数，也就是前面标有__global__的函数，如下：
1 __global__ void add( int *a, int *b, int *c ) { 2 int tid = blockIdx.x; // this thread handles the data at its thread id 3 if (tid < N) 4 c[tid] = a[tid] + b[tid]; 5 }
kernel函数的用意就是gpu中的每一个thread都会执行kernel，从而达到并行的目的。a，b，c三个参数传入所有thread，每一个thread完成加法操作，为了使每一个thread的加法是对应数组中不同的元素，所以变量tid就意义重大，blockIdx.x是runtime中提供的变量，通常来讲，一个gpu有1个grid，1个grid有多个block，这些block以一维或二维或三维数组的形式排列，blockIdx.x就是每个block在x方向上的索引值（就是序号），而每一个block又可以分为多个thread，thread按照一维或二维的方式排列。网上摘图一个，以供理解
这个图还说明了一个kernel就有一个grid，多个kernel有多个grid，涨姿势。
接着说 tid = blockIdx.x，就是说要把每一个block的序号赋值给tid，block的序号是0，1，2，。。。这样排列的，数组a，b，c的索引值tid也是0，1，2.。。。这样的，这不就说明每个block都会计算他自己的那个数组分量了，这不是巧了吗。
cudaMalloc( (void**)&dev_a, N * sizeof(int) )函数可以在device上申请存储空间，这里有一个原则，就是host中的代码不能操作用cudaMalloc申请的空间，因此想要释放空间，用cudaFree。kernel才是操作cudaMalloc申请的变量的函数，kernel中用到的其他函数需要以__device__标明。
为变量申请空间后用函数cudaMemcpy( dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice )把host中的变量a赋值进device中即dev_a。然后device自己计算自己那份任务，即kernel中的计算。计算后用函数cudaMemcpy( c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost )将结果dev_c传会host。整体就是这么个流程。
需要注意的是对计算发出命令的add<<<N,1>>>( dev_a, dev_b, dev_c );kernel函数在调用时需要两个参数，参数告诉runtime如何launch the kernel，就是怎么把kernel复制给那些计算单元的意思吧N代表N个一维block，1代表每个block里面包含1个thread。
