CUDA中文教程02之心得体会.doc

CUDA中文教程02之心得体会第二集主要是讲了CUDA编程模型，真正的介绍了如何编写一个CUDA程序。首先是介绍了一些基本的概念及数据类型。作为CPU的协处理器，GPU有自己的存储器，可以并行的进行多线程计算，是一种并行的处理设备。接下来是介绍了一些关键字以及应用程序接口等。CUDA语言跟C语言相似。一个CUDA kernel 是用一个序列数组来执行的，所有的线程都运行一样的代码，每一个线程都有自己的ID，负责被指定的任务。前面提过，一个block包含了512个threads，在同一个block内的thread可以通过shared memory 同步合作，但是不同的block里的thread 是不能进行合作的。每个线程有自己ID去决定处理什么数据。block也有自己的ID，现在的block ID只包含两维，e.g. Block（0，0），Block （1，0）。thread ID包含三维，e.g. Thread （0，0，0）。在CPU（host）和GPU（device）之间进行读写数据操作要通过 Global memory。它的内容对于所有线程来说是可见的。但是每一次访问Global memory 有很长的延迟。CUDA的应用程序接口（API）是C语言的扩展。CUDA的存储分配函数cudaMalloc（），函数将对象分配到Global memory 中，这个函数需要两个参数，1、指向这个对象的指针地址，2、分配对象的大小。当完成任务后，用cudaFree（）释放Global memory里面的对象。从host到Global memory 的数据传输，用到的函数是cudaMemcpy（），这个函数需要四个参数，1、指向目的地的指针，2、指向源头的指针，3、需要复制的字节数，4、传输的类型（host to host / host to device / device to host / device to device）。 下面是一个例子：/从host到device的传输，cudaMemcpyHostToDevice是一个象征常量 cudaMemcpy（Md.elements,M.elements,size,cudaMemcpyHostToDevice）; /从device到host的传输，cudaMemcpyDeviceToHost是一个象征常量 cudaMemcpy（M.elements,Md.elements,size,cudaMemcpyDeviceToHost）;CUDA关键字。1、当函数是在device上获取且执行的，可以在函数前冠以_device_:e.g. _device_ float DeviceFunc()2、当函数是在host上获取但是在device上执行的，可以在函数前冠以_global_:&M t8nQl g 0e.g. _global_ void KernelFunc() /must return voidITPUB个人空间*SW:Hf6 w3、当函数是在host上获取且执行的，可以在函数前冠以_host_:ITPUB个人空间 LW-*d.*e6|e.g. _host_ float HostFunc()一般情况下，_host_可以省略，是一种缺省状态,只有当定义在device上的函数也想要在host上使用，就可以同时将_device_和_host_同时冠以函数前。 注意，在device上执行的函数没有递归，没有静态变量，没有可变参数。 矩阵相乘例子：void MatrixMulOnDevice(float* M, float* N, float* P, int Width)ITPUB个人空间D4gb-_)mUe6Ze(d&k)8Z0 int size = Width * Width * sizeof(float);ITPUB个人空间panR2uPZy%D float* Md, Nd, Pd;)M s7C76A0 7_!i1xwv M?&U01. / Allocate and Load M, N to device memoryITPUB个人空间(b0X.wx5e,5gQ.C!V cudaMalloc(&Md, size);5a%!R?:LTWh(p+o0 cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice);NqW8qd6J p*R0 cudaMalloc(&Nd, size);ITPUB个人空间+6mi q/T3i,Lp&S cudaMemcpy(Nd, N, size, cudaMemcpyHostToDevice);ITPUB个人空间8t(aYh6D / Allocate P on the device(e&F8UBjz/a l0 cudaMalloc(&Pd, size);P;w0j.l Y*sx1X0 ITPUB个人空间/N.eaP8Uy5a2. / Kernel invocation code ITPUB个人空间w7WZlU I6z6t;v eJXIe/w4J0/ Matrix multiplication kernel per thread codeITPUB个人空间%E21W V(a$H(Gq_global_ void MatrixMulKernel(float* Md, float* Nd, float* Pd, int Width)EeZH4rZ0r| iL(mw S50 / 2D Thread IDm8m8Q5sW;0 int tx = threadIdx.x;ITPUB个人空间8a1P0ya_5lud int ty = threadIdx.y;ITPUB个人空间3Zm5Z U.mb0o*z0 | / Pvalue is used to store the element of the matrixT&xR1 v2E)Q0 / that is computed by the thread|#GN-BH D$l0 float Pvalue = 0; )C _ j7 g h3B0 for (int k = 0; k Width; +k)G7S?1o euC0 vp-Bg.yIT|0 float Melement = Mdty * Width + k;ITPUB个人空间mF;l%tw1km(a+Gui float Nelement = Ndk * Width + tx;7e(nH/Ss$t1a9D0 Pvalue += Melement * Nelement; b+|&sJ iQ O0 e7ZR;$S1mTZ0 / Write the matrix to device memory;ITPUB个人空间kjZ7Na:B / each thread writes one element&|U&HC4;syL0 Pdty * Width + tx = Pvalue;ITPUB个人空间7L*:o*q)A Q3sWlbEI$6p.w&QB0 8MI(w0d0 / Setup the execution configurationh fM.|m0 dim3 dimBlock(Width, Width);ITPUB个人空间3|6Y IQPN2w!R dim3 dimGrid(1, 1);a- I:mX3c9E0m0 / Launch the device computation threads!ITPUB个人空间fC1v.T!j-E#tk MatrixMulKernel(Md, Nd, Pd);9Ag_x!p*c4?T/P M0 ITPUB个人空间$|7W H8GG9p2P3. / Read P from the deviceP*WzN_5y6O0