一般而言,CUDA程序的基本模式是:
分配内存空间和显存空间
初始化内存空间
将要计算的数据从内存上复制到显存上
执行kernel计算
将计算后显存上的数据复制到内存上
处理复制到内存上的数据
// This is the REAL "hello world" for CUDA!
// It takes the string "Hello ", prints it, then passes it to CUDA with an array
// of offsets. Then the offsets are added in parallel to produce the string "World!"
// By Ingemar Ragnemalm 2010
#include <stdio.h>
const int N = 16;
const int blocksize = 16;
__global__ void hello(char *a, int *b)
{
a[threadIdx.x] += b[threadIdx.x];
}
int main()
{
char a[N] = "Hello \0\0\0\0\0\0";
int b[N] = {15, 10, 6, 0, -11, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
char *ad;
int *bd;
const int csize = N*sizeof(char);
const int isize = N*sizeof(int);
printf("%s", a);
cudaMalloc( (void**)&ad, csize );
cudaMalloc( (void**)&bd, isize );
cudaMemcpy( ad, a, csize, cudaMemcpyHostToDevice );
cudaMemcpy( bd, b, isize, cudaMemcpyHostToDevice );
dim3 dimBlock( blocksize, 1 );
dim3 dimGrid( 1, 1 );
hello<<<dimGrid, dimBlock>>>(ad, bd);
cudaMemcpy( a, ad, csize, cudaMemcpyDeviceToHost );
cudaFree( ad );
cudaFree( bd );
printf("%s\n", a);
return EXIT_SUCCESS;
}
注释:
函数类型限定符,用来确定函数是在CPU还是在GPU上执行,以及这个函数是从CPU调用还是从GPU调用。
__device__,__device__表示从GPU上调用,在GPU上执行;
__global__,__global__表示在CPU上调用,在GPU上执行,也就是所谓的内核(kernel)函数;内核主要用来执行多线程调用。
__host__,__host__表明在CPU上调用,在CPU上执行,这是默认时的情况,也就是传统的C函数。CUDA支持__host__和__device__的联用,表示同时为主机和设备编译。此时这个函数不能出现多线程语句。
变量类型限定符,用来规定变量存储什么位置上。在传统的CPU程序上,这个任务由编译器承担。在CUDA中,不仅要使用主机端的内存,还要使用设备端的显存和GPU片上的寄存器、共享存储器和缓存。在CUDA存储器模型中,一共抽象出来了8种不同的存储器。复杂的存储器模型使得必须要使用限定符要说明变量的存储位置。
__device__,__device__表明声明的数据存放在显存中,所有的线程都可以访问,而且主机也可以通过运行时库访问;
__shared__,__shared__表示数据存放在共享存储器在,只有在所在的块内的线程可以访问,其它块内的线程不能访问;
__constant__,__constant__表明数据存放在常量存储器中,可以被所有的线程访问,也可以被主机通过运行时库访问;
执行配置运算符<<< >>>,用来传递内核函数的执行参数。执行配置有四个参数,第一个参数声明网格的大小,第二个参数声明块的大小,第三个参数声明动态分配的共享存储器大小,默认为0,最后一个参数声明执行的流,默认为0。
五个内建变量,用于在运行时获得网格和块的尺寸及线程索引等信息
gridDim, gridDim是一个包含三个元素x,y,z的结构体,分别表示网格在x,y,z三个方向上的尺寸,虽然其有三维,但是目前只能使用二维;
blockDim, blockDim也是一个包含三个元素x,y,z的结构体,分别表示块在x,y,z三个方向上的尺寸,对应于执行配置中的第一个参数,对应于执行配置的第二个参数;
blockIdx, blockIdx也是一个包含三个元素x,y,z的结构体,分别表示当前线程所在块在网格中x,y,z三个方向上的索引;
threadIdx, threadIdx也是一个包含三个元素x,y,z的结构体,分别表示当前线程在其所在块中x,y,z三个方向上的索引;
warpSize,warpSize表明warp的尺寸,在计算能力为1.0的设备中,这个值是24,在1.0以上的设备中,这个值是32。