CUDA 初体验

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是三维的向量，可以表示为一维、二维、三维的线程索引。如果是二维的话，若尺寸是(Dx,Dy)，那么索引的就是(x+yDx)。如果是三维的，索引的就是（x，y，z），那么就是(x+yDx+zDxDy)

现在线程块一般是1024个，但是因为有多个线程块。所以总的线程数是每块线程数x线程块数。

通过调用__syncthreads()函数进行数据同步。

CUDA的每个线程、线程块等等的内存层次：

除了全局存储之外，还有两种额外的存储：常量和texture memory（这个玩样儿是啥？）。

CUDADeviceReset()的调用使得所有的配置初始化。

CUDA上的存储操作有cudaMalloc(), cudaFree(). cudaMemcpy()。

举一个例子：

// 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 codeint main(){int N = ...;size_t size = N * sizeof(float);// Allocate input vectors h_A and h_B in host memoryfloat* h_A = (float*)malloc(size);float* h_B = (float*)malloc(size);// Initialize input vectors...// Allocate vectors in device memoryfloat* 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 memorycudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);// Invoke kernelint 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 memorycudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);// Free device memorycudaFree(d_A);cudaFree(d_B);cudaFree(d_C);// Free host memory...}

cudaMallocPitch(), cudaMalloc3D()可以用来分配内存。另外还有cudaMemcpy2D和cudaMemcpy3D来分配2D和3D的内存。P34行有例子。

shared memory

shared 标识，共享的内存比全局的内存更快。这里举了一个矩阵乘法的例子： P35

在P41页有memory blocking存在，更加快。

Page locked out memory

和传统的malloc分配的内存相反的，这种比较固定。

cudaHostAlloc() 和 cudaFreeHost()。

在CUDA里面涉及数据同步和流的东西，这里有显示同步和隐式同步。还有更多数据流的东西，比如数据传过去kernel的时候有的已经在执行啦什么的。还有callback函数。

P57里面有各种API。

CUDA里面的硬件架构上，有SIMD和多线程。

C

一些CUDA的语法，涉及和C有关的东西。类似于API。

在这里，贴上矩阵的CUDA算法，最基本的，然后需要在上面进行加速：

#include <cuda.h>#include <cuda_runtime.h>#include "matrix_mul.h"#define TILE_WIDTH 2namespace 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, float *sq_matrix_2, float *sq_matrix_result, 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, size, 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, sq_matrix_result_d, size, cudaMemcpyDeviceToHost);    cudaFree(sq_matrix_1_d);    cudaFree(sq_matrix_2_d);    cudaFree(sq_matrix_result_d);  }} // namespace cuda

CUDA 调用

核函数是GPU每个thread上运行的程序。必须通过gloabl函数类型限定符定义。形式如下：

            __global__ void kernel(param list){  }

核函数只能在主机端调用，调用时必须申明执行参数。调用形式如下：

            Kernel<<<Dg,Db, Ns, S>>>(param list);

<<<>>>运算符内是核函数的执行参数，告诉编译器运行时如何启动核函数，用于说明内核函数中的线程数量，以及线程是如何组织的。

<<<>>>运算符对kernel函数完整的执行配置参数形式是<< < Dg, Db, Ns, S>>> 【2】

• 参数Dg用于定义整个grid的维度和尺寸，即一个grid有多少个block。为dim3类型。Dim3 Dg(Dg.x, Dg.y, 1)表示grid中每行有Dg.x个block，每列有Dg.y个block，第三维恒为1(目前一个核函数只有一个grid)。整个grid中共有Dg.x*Dg.y个block，其中Dg.x和Dg.y最大值为65535。
• 参数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的硬件支持的最大值为1024。
• 参数Ns是一个可选参数，用于设置每个block除了静态分配的shared Memory以外，最多能动态分配的shared memory大小，单位为byte。不需要动态分配时该值为0或省略不写。
• 参数S是一个cudaStream_t类型的可选参数，初始值为零，表示该核函数处在哪个流之中。

CUDA 编程介绍

比如举个计算一个数字每个数字平方和的CUDA实现。

#include <stdio.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 CPU       cudaMemcpy(h_out, d_out, ARRAY_BYTES, 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;  }  

CUDA 数据同步

原本有问题的代码：

__global__ void shift(){      int idx = threadIdx.x;      __shared__ int array[128];      array[idx] = threadIdx.x;      if (idx < 127) {          array[idx] = array[idx + 1];      }  }  

设置barrier：

__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();//将一行代码拆分成两行来设置一个barrier,这种技巧非常实用，执行至此，每一个线程都正确的取值        array[idx] = temp;        __syncthreads();//确保后续使用array的正确性    }}

参考资料：
【1】CUDA 编程指南：http://docs.nvidia.com/cuda/pdf/CUDA_C_Programming_Guide.pdf
【2】CUDA 调用说明：http://blog.csdn.net/augusdi/article/details/12204121
【3】CUDA 核函数的参数解析：http://blog.csdn.net/a925907195/article/details/39500915