第二集主要是讲了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);
}