cuda基础学习笔记

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>用于将错误代码转化为文本信息