CUDA笔记
写在前面:
这是cuda官网上cuda programming basice的笔记,cuda环境的搭建在win下只需要安装vs2013+cuda 工具包就可以了,linux下安装工具包后需要再安装gcc等编译环境并配置环境变量,mac类似linux,只不过编译环境可以直接使用xcode。
一、编译
物理运行方式:
发布模式:nvcc <文件名>.cu [-o 输出文件名]
调试模式(只能调试主机代码,不能调试gpu代码):nvcc –g <文件名>.cu
模拟运行方式:
无符号模式(所有代码运行于主机):nvcc–deviceemu <文件名>.cu
有符号模式(所有代码运行于主机):nvcc–deviceemu -g <文件名>.cu
二、内存管理
首先分配内存,然后从主机内存向gpu内存拷贝,然后进行计算,将结果拷贝回主机内存,最后释放不需要的内存。
1.分配内存与释放内存
1)函数原型
(1) cudaMalloc(void ** 指针,size_t 字节数)
(2) cudaMemset(void * 指针,int 值,size_t设置数目)
(3) cudaFree(void * 指针)
2)用例
int n=1024;
int nbytes =1024*sizeof(int);
int *d_a=0;
cudaMalloc((void **) &d_a, nbytes);
cudaMemset(d_a,0,nbytes);
cudaFree(d_a);
2.数据复制
1)函数原型
cudaMemcpy(void * dst,void * src,size_t nbytes,enum cudaMemcpyKind direction)
disection是用于标示复制设备的值,无论是多少,此api均将把src指向的空间内的数据复制到dst指向的空间中,调用此api的cpu线程将在复制完成后返回,此api不能在GPU中调用。
(1) enum cudaMemcpyKind解释
<1>cudaMemcpyHostToDevice表示从主机复制到设备,实际上是从内存复制到显存
<2>cudaMemcpyDeviceToHost表示从设备复制到主机,实际上是从显存复制到内存
<3>cudaMemcpyDeviceToDevice表示从设备复制到设备,实际上用于两个显卡间的显存复制或两个显存空间间的复制
2)用例
int main(void)
{
float *a_h,*b_h;//主机数据
float *a_d,*b_d;//设备数据
int N=14,nBytes,I;
nBytes=N*sizeof(float);
a_h=(float *)malloc(nBytes);
b_h=(float *)malloc(nBytes);
cudaMalloc((void**)&a_d,nBytes);
cudaMalloc((void**)&b_d,nBytes);
for(i=0,i<N;i++)
{
a_h[i]=100.0f+I;
}
cudaMemcpy(a_d,a_h,nBytes,cudaMemcpyHostToDevice);
cudaMemcpy(b_d,a_d,nBytes,cudaMemcpyDeviceToDevice);
cudaMemcpy(b_h,b_d,nBytes,cudaMemcpyDeviceToHost);
for(i=0; i<N;i++)
{
assert(a_h[i]==b_h[i]);
}
free(a_h);
free(b_h);
cudaFree(a_d);
cudaFree(b_d);
return 0;
}
三、核函数(内核函数)
1.内核函数是一种有一些限制的c函数
1)不能访问主机内存
2)返回类型只能是void
3)不允许可变参数
4)不允许递归
5)不允许静态变量
2函数参数将自动地从主机复制到设备
3.函数修饰语
1)__global__
(1)此种函数从主机调用,在设备上执行
(2)只允许void返回
2)__device__
(1)此种函数从设备上调用,在设备上执行
(2)不允许被主机代码调用
3)__host__
(1)此种函数从主机调用,在主机上执行(如不写修饰语,默认函数为__host__修饰)
4.核函数的调用
1)原型
函数名<<<dim3dG,dim3 dB>>(参数);
2)说明
(1)dG表明grid(网格)中blocks(块)的数目
<1>只有两个维度有效,x和y,z始终只能为1(当前限制,不知道以后会不会解除)
<2>块的总量:dG.x*dG.y
(2)dB表明块中threads(线程)的数目
<1>三个维度都有效,x、y和z
<2>线程的总量:dB.x*dB.y*dB.z
(3)dim3中未被定义的值默认为1
3)用例
(1)
dim3 grid,block;
grid.x=2;
grid.y=4;
block.x=8;
block.y=16;
kernel<<<grid,block>>>(…);
(2)
dim3 grid(2,4),block(8,16);
kernel<<<grid,block>>>(…);
(3)
kernel<<<32,512>>>(…);
4)设备变量管理
(1)所有设备上运行的函数都能自由访问它们的局部变量
(2)特殊变量
<1>dim3gridDim:网格中块的数目(最多二维)
<2>dim3blockDim:块中线程数目
<3>dim3blockIdx:在网格中的块引索(运行时自动生成的唯一值)
<4>dim3threadIdx:块中的线程引索(运行时自动生成的唯一值)
5)主机同步
(1)所有核函数的运行都是异步的
<1>主机调用后会立即返回
<2>内核函数将在cuda之前的calls完成后执行
(2)cudaMemcpy()是个同步过程
<1>将会在复制完成后返回主机
<2>复制在cuda之前的calls完成后执行
<3>使用方式:先复制数据到设备,调用核函数,执行其他主机代码,从设备复制结果数据
(3)cudaThreadSynchronize()
<1>执行到此处后等待块中所有线程均运行到此处后继续执行
6)变量修饰语
(1)__device__
<1>在设备全局内存中,容量大,所有线程均可访问,无缓存
<2>使用cudaMalloc进行分配(隐含__device__修饰)
<3>存在于整个程序生存周期内(你不释放或程序结束都在那)
(2)__shared__
<1>块内线程可访问,容量小,速度快
<2>存在于块线程存在时(块线程结束自动释放)
(3)未修饰变量
<1>标量和运行时建立的矢量在寄存器中
<2>无法安置在寄存器中的变量保存在local(本地)内存中
7)设备同步
(1)void __syncthreads();
<1>设备上调用
<2>只有块中所有线程均执行到此处时才可继续向下执行
(2)分支操作
<1>只有处于一个分支的所有线程执行完毕才会执行另一个分支的线程
8)GPU原子操作
此处需要查询书
9)变量类型(向量)
(1)可以被用于gpu和cpu代码中的变量类型
<1>[u]char[1..4],[u]short[1..4],[u]int[1..4],[u]long[1..4],float[1..4],double[1..4]
可以使用下标:x,y,z,w
例子:
float4 xx;
float yy=xx.w;
<2>dim3
基于unit3,初始值为(1,1,1)
10)cuda错误反馈(主机端)
(1)除调用核函数外的所有cuda函数调用都会返回错误代码,类型为cudaError_t,未发生错误返回cudaSuccess。
(2)cudaError_t cudaGetLastError(void)
<1>返回最后一个错误的错误代码(如果没有错误也会返回一个值,值为cudaSuccess)
<2>可以用于获取核函数调用时发生的错误的错误代码
(3)char* cudaGetErrorString(cudaError_t)
<1>用于将错误代码转化为文本信息