CUDA Visual Profiler

在上180645课程的时候，里面谈到使用CUDA来做矩阵乘法和k均值聚类的加速。在使用n卡的时候，有一个Visual Profiler的东西可以看到GPU的使用信息。

在安装好了CUDA以后，在Ubuntu上登录以后，使用X server。在Ubuntu命令行输入：

ssh -X < your_andrew_id>@ghcXX.ghc.andrew.cmu.edu

然后就登陆了远程服务器，接着呢使用：

computeprof &

如果遇到错误，退出登录再连接就好了。

这样就可以看到了GPU的使用信息了。然后如果是Windows的话，使用Xming或Cygwin。如果是OS X的话，使用XQuartz就可以了。

CUDA编程指导

使用CUDA编程，可以学习CUDA编程指南【1】。接下来我就大概过一遍编程指南。

threadIdx是三维的向量，可以表示为一维、二维、三维的线程索引。如果是二维的话，若尺寸是 $(D_x,D_y)@H_<a href="/tag/301/" target="_blank" class="keywords">301</a>_110@，那么索引的就是 (x+yDx) <script type="math/tex" id="MathJax-Element-2">(x + yD_x)@H_<a href="/tag/301/" target="_blank" class="keywords">301</a>_110@。如果是三维的，索引的就是（x，y，z），那么就是( x+yDx+zDxDy <script type="math/tex" id="MathJax-Element-3">x+yD_x+zD_xD_y@H_<a href="/tag/301/" target="_blank" class="keywords">301</a>_110@) 现在线程块一般是1024个，但是因为有多个线程块。所以总的线程数是每块线程数x线程块数。 通过<a href="/tag/diaoyong/" target="_blank" class="keywords">调用</a>__syncthreads()<a href="/tag/hanshu/" target="_blank" class="keywords">函数</a>进行数据同步。 CUDA的每个线程、线程块等等的内存层次： 除了全局存储之外，还有两种额外的存储：常量和texture memory（这个玩样儿是啥？）。 CUDADeviceReset()的<a href="/tag/diaoyong/" target="_blank" class="keywords">调用</a>使得所有的配置初始化。 CUDA上的存储操作有cudaMalloc(),cudaFree(). cudaMemcpy()。 举一个例子： <pre class="prettyprint">// Device code __global__ void VecAdd(float* A,float* B,float* C,int N) { int i = blockDim.x * blockIdx.x + threadIdx.x; if (i < N) C[i] = A[i] + B[i]; } // Host code int main() { int N = ...; size_t size = N * sizeof(float); // Allocate input vectors h_A and h_B in host memory float* h_A = (float*)malloc(size); float* h_B = (float*)malloc(size); // Initialize input vectors ... // Allocate vectors in device memory float* d_A; cudaMalloc(&d_A,size); float* d_B; cudaMalloc(&d_B,size); float* d_C; cudaMalloc(&d_C,size); // Copy vectors from host memory to device memory cudaMemcpy(d_A,h_A,size,cudaMemcpyHostToDevice); cudaMemcpy(d_B,h_B,cudaMemcpyHostToDevice); // Invoke kernel int threadsPerBlock = 256; int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock; VecAdd<<<blocksPerGrid,threadsPerBlock>>>(d_A,d_B,d_C,N); // Copy result from device memory to host memory // h_C contains the result in host memory cudaMemcpy(h_C,cudaMemcpyDeviceToHost); // Free device memory cudaFree(d_A); cudaFree(d_B); cudaFree(d_C); // Free host memory ... }</pre> cudaMallocPitch(),cudaMalloc3D()可以用来分配内存。另外还有cudaMemcpy2D和cudaMemcpy3D来分配2D和3D的内存。P34行有例子。 <h2 id="shared-memory">shared memory</h2> shared 标识，共享的内存比全局的内存更快。这里举了一个矩阵乘法的例子： P35 在P41页有memory blocking存在，更加快。 <h2 id="page-locked-out-memory">Page locked out memory</h2> 和传统的malloc分配的内存相反的，这种比较固定。 cudaHostAlloc() 和 cudaFreeHost()。 在CUDA里面涉及数据同步和流的东西，这里有<a href="/tag/xianshi/" target="_blank" class="keywords">显示</a>同步和隐式同步。还有更多数据流的东西，比如数据传过去kernel的时候有的已经在执行啦什么的。还有callback<a href="/tag/hanshu/" target="_blank" class="keywords">函数</a>。 P57里面有各种API。 CUDA里面的硬件架构上，有SIMD和多线程。 <h1 id="c">C</h1> 一些CUDA的语法，涉及和C有关的东西。类似于API。 在这里，贴上矩阵的CUDA算法，最基本的，然后需要在上面进行加速： <pre class="prettyprint">#include <cuda.h> #include <cuda_runtime.h> #include "matrix_mul.h" #define TILE_WIDTH 2 namespace cuda { __global__ void matrix_mul_kernel(float *sq_matrix_1,float *sq_matrix_2,float *sq_matrix_result,int sq_dimension) { int tx = threadIdx.x; int ty = threadIdx.y; float sum = 0.0f; for(int k = 0; k < sq_dimension; k++) { sum += sq_matrix_1[ty*sq_dimension + k] * sq_matrix_2[k*sq_dimension + tx]; } sq_matrix_result[ty*sq_dimension + tx] = sum; } void matrix_multiplication(float *sq_matrix_1,unsigned int sq_dimension) { int size = sq_dimension * sq_dimension * sizeof(float); float *sq_matrix_1_d,*sq_matrix_2_d,*sq_matrix_result_d; /*************************************************** 1st Part: Allocation of memory on device memory ****************************************************/ /* copy sq_matrix_1 and sq_matrix_2 to device memory */ cudaMalloc((void**) &sq_matrix_1_d,size); cudaMemcpy(sq_matrix_1_d,sq_matrix_1,size,cudaMemcpyHostToDevice); cudaMalloc((void**) &sq_matrix_2_d,size); cudaMemcpy(sq_matrix_2_d,sq_matrix_2,cudaMemcpyHostToDevice); /*allocate sq_matrix_result on host */ cudaMalloc((void**) &sq_matrix_result_d,size); /*************************************************** 2nd Part: Inovke kernel ****************************************************/ dim3 dimBlock(sq_dimension,sq_dimension); dim3 dimGrid(1,1); matrix_mul_kernel<<<dimGrid,dimBlock,dimBlock.x * dimBlock.x * sizeof(float)>>>(sq_matrix_1_d,sq_matrix_2_d,sq_matrix_result_d,sq_dimension); /*************************************************** 3rd Part: Transfer result from device to host ****************************************************/ cudaMemcpy(sq_matrix_result,cudaMemcpyDeviceToHost); cudaFree(sq_matrix_1_d); cudaFree(sq_matrix_2_d); cudaFree(sq_matrix_result_d); } } // namespace cuda</pre> <h2 id="cuda-调用">CUDA <a href="/tag/diaoyong/" target="_blank" class="keywords">调用</a></h2> 核<a href="/tag/hanshu/" target="_blank" class="keywords">函数</a>是GPU每个thread上运行的程序。必须通过gloabl<a href="/tag/hanshu/" target="_blank" class="keywords">函数</a>类型限定符定义。形式如下： <pre>__global__ void kernel(param list){ }</pre> 核<a href="/tag/hanshu/" target="_blank" class="keywords">函数</a>只能在主机端<a href="/tag/diaoyong/" target="_blank" class="keywords">调用</a>，<a href="/tag/diaoyong/" target="_blank" class="keywords">调用</a>时必须申明执行参数。<a href="/tag/diaoyong/" target="_blank" class="keywords">调用</a>形式如下： <pre>Kernel<<<Dg,Db,Ns,S>>>(param list);</pre> <<<>>>运算符内是核<a href="/tag/hanshu/" target="_blank" class="keywords">函数</a>的执行参数，告诉编译器运行时如何启动核<a href="/tag/hanshu/" target="_blank" class="keywords">函数</a>，用于说明内核<a href="/tag/hanshu/" target="_blank" class="keywords">函数</a>中的线程<a href="/tag/shuliang/" target="_blank" class="keywords">数量</a>，以及线程是如何组织的。 <<<>>>运算符对kernel<a href="/tag/hanshu/" target="_blank" class="keywords">函数</a>完整的执行配置参数形式是<< < Dg,S>>> 【2】 <ul> <li>参数Dg用于定义整个grid的维度和尺寸，即一个grid有多少个block。为dim3类型。Dim3 Dg(Dg.x,Dg.y,1)表示grid中每行有Dg.x个block，每列有Dg.y个block，第三维恒为1(目前一个核<a href="/tag/hanshu/" target="_blank" class="keywords">函数</a>只有一个grid)。整个grid中共有Dg.x*Dg.y个block，其中Dg.x和Dg.y最大值为65535。</li> <li>参数Db用于定义一个block的维度和尺寸，即一个block有多少个thread。为dim3类型。Dim3 Db(Db.x,Db.y,Db.z)表示整个block中每行有Db.x个thread，每列有Db.y个thread，高度为Db.z。Db.x和Db.y最大值为512，Db.z最大值为62。一个block中共有Db.x*Db.y*Db.z个thread。计算能力为1.0,1.1的硬件该乘积的最大值为768，计算能力为1.2,1.3的硬件<a href="/tag/zhichi/" target="_blank" class="keywords">支持</a>的最大值为1024。</li> <li>参数Ns是一个可选参数，用于设置每个block除了静态分配的shared Memory以外，最多能动态分配的shared memory大小，单位为byte。不需要动态分配时该值为0或省略不写。</li> <li>参数S是一个cudaStream_t类型的可选参数，初始值为零，表示该核<a href="/tag/hanshu/" target="_blank" class="keywords">函数</a>处在哪个流之中。</li> </ul> <h2 id="cuda-编程介绍">CUDA 编程介绍</h2> 比如举个计算一个数字每个数字平方和的CUDA实现。 <pre class="prettyprint">#include <st<a href="/tag/dio/" target="_blank" class="keywords">dio</a>.h> __global__ void square(float * d_out,float * d_in) { int idx = threadIdx.x; float f = d_in[idx]; d_out[idx] = f * f; } int main(int argc,char ** argv) { const int ARRAY_SIZE = 64; const int ARRAY_BYTES = ARRAY_SIZE * sizeof(float); // generate the input array on the host float h_in[ARRAY_SIZE]; for (int i = 0; i < ARRAY_SIZE; i++) { h_in[i] = float(i); } float h_out[ARRAY_SIZE]; // declare GPU memory pointers float *d_in; float *d_out; // allocate GPU memory cudaMalloc((void**) &d_in,ARRAY_BYTES); cudaMalloc((void**) &d_out,ARRAY_BYTES); // transfer the array to the GPU cudaMemcpy(d_in,h_in,ARRAY_BYTES,cudaMemcpyHostToDevice); // launch the kernel square<<<1,ARRAY_SIZE>>>(d_out,d_in); // copy back the result array to the <a href="/tag/cpu/" target="_blank" class="keywords">cpu</a> cudaMemcpy(h_out,d_out,cudaMemcpyDeviceToHost); // print out the resulting array for (int i =0; i < ARRAY_SIZE; i++) { printf("%f",h_out[i]); printf(((i % 4) != 3) ? "\t" : "\n"); } cudaFree(d_in); cudaFree(d_out); return 0; }</pre> <h2 id="cuda-数据同步">CUDA 数据同步</h2> 原本有问题的<a href="/tag/daima/" target="_blank" class="keywords">代码</a>： <pre class="prettyprint">__global__ void shift(){ int idx = threadIdx.x; __shared__ int array[128]; array[idx] = threadIdx.x; if (idx < 127) { array[idx] = array[idx + 1]; } }</pre> 设置barrier： <pre class="prettyprint">__global__ void shift(){ int idx = threadIdx.x; __shared__ int array[128]; array[idx] = threadIdx.x; __syncthreads();//执行至此，数组中的每一个元素都被正确的赋值 if (idx < 127) { int temp = array[idx + 1]; __syncthreads();//将一行<a href="/tag/daima/" target="_blank" class="keywords">代码</a>拆分成两行来设置一个barrier,这种技巧非常实用，执行至此，每一个线程都正确的取值 array[idx] = temp; __syncthreads();//确保后续使用array的正确性 } }</pre> 参考资料： 【1】CUDA 编程指南：<a href="javascript:void()" class="outlink" data="/link?url=http://docs.nvidia.com/cuda/pdf/CUDA_C_Programming_Guide.pdf" rel="nofollow" target="_blank">http://docs.nvidia.com/cuda/pdf/CUDA_C_Programming_Guide.pdf</a> 【2】CUDA <a href="/tag/diaoyong/" target="_blank" class="keywords">调用</a>说明：<a href="http://www.jb51.cc/article/p-ducfhdos-zb.html">http://www.jb51.cc/article/p-ducfhdos-zb.html</a> 【3】CUDA 核<a href="/tag/hanshu/" target="_blank" class="keywords">函数</a>的参数解析：<a href="http://www.jb51.cc/article/p-fubmdykp-zd.html">http://www.jb51.cc/article/p-fubmdykp-zd.html</a></div> <div class="topcard-tags"></div> <ul class="list-group"> <li class="list-group-item"><a href="/ubuntu/354438.html" title="Ubuntu安装后常见部署">上一篇：Ubuntu安装后常见部署</a><a href="/ubuntu/354436.html" title="LAMP服务器优化--优化前准备" class="text-muted pull-right">下一篇：LAMP服务器优化--优化前准备</a> </li> </ul> </div> </div> </div>  <div class="row row-sm"> <div class="col-sm-12 col-md-12 col-lg-12"> <div class="card"> <ins class="adsbygoogle" style="display:block" data-ad-format="autorelaxed" data-ad-client="ca-pub-4605373693034661" data-ad-slot="9144498553"></ins> <script> (adsbygoogle = window.adsbygoogle || []).push({});$

CUDA 初体验

CUDA Visual Profiler

CUDA编程指导

猜你在找的Ubuntu相关文章