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__:

e.g. __global__ void KernelFunc() //must return void

3、当函数是在host上获取且执行的，可以在函数前冠以__host__:

e.g.  __host__ float HostFunc()

一般情况下，__host__可以省略，是一种缺省状态,只有当定义在device上的函数也想要在host上使用，就可以同时将——__device__和__host__同时冠以函数前。

 

注意，在device上执行的函数没有递归，没有静态变量，没有可变参数。

 

矩阵相乘例子：

void MatrixMulOnDevice(float* M, float* N, float* P, int Width)

{

   int size = Width * Width * sizeof(float);

    float* Md, Nd, Pd;

   …

1. // Allocate and Load M, N to device memory

    cudaMalloc(&Md, size);

    cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice);

     cudaMalloc(&Nd, size);

     cudaMemcpy(Nd, N, size, cudaMemcpyHostToDevice);

     // Allocate P on the device

    cudaMalloc(&Pd, size);

 

2.  // Kernel invocation code    

   

// Matrix multiplication kernel – per thread code

__global__ void MatrixMulKernel(float* Md, float* Nd, float* Pd, int Width)

{

    // 2D Thread ID

    int tx = threadIdx.x;

    int ty = threadIdx.y;

    // Pvalue is used to store the element of the matrix

    // that is computed by the thread

    float Pvalue = 0;   

    for (int k = 0; k < Width; ++k)

    {

         float Melement = Md[ty * Width + k];

         float Nelement = Nd[k * Width + tx];

         Pvalue += Melement * Nelement;

    }

    // Write the matrix to device memory;

    // each thread writes one element

    Pd[ty * Width + tx] = Pvalue;

}

  

    // Setup the execution configuration

    dim3 dimBlock(Width, Width);

    dim3 dimGrid(1, 1);

    // Launch the device computation threads!

    MatrixMulKernel<<<dimGrid, dimBlock>>>(Md, Nd, Pd);

 

3.    // Read P from the device

      cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost);

       // Free device matrices

      cudaFree(Md); cudaFree(Nd); cudaFree (Pd);

}