第二集主要是讲了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<<>>(Md, Nd, Pd);
 
3.    // Read P from the device
      cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost);
       // Free device matrices
      cudaFree(Md); cudaFree(Nd); cudaFree (Pd);
}

源自：xieyan blog