CUDA编程入门

rowheighPoint.y
colwidthPoint.x

1 GPU硬件架构及运行机制

文章写的很全面

2 GPU计算基础知识

CUDA编程需要CPU和GPU协同工作，程序中既包含host程序，又包含device程序，他们分别在CPU和GPU上运行，二者间可以进行通信。

host - CPU及其内存 device - GPU及其内存

2.1 kernel 核函数

• kernel是在device上线程中并行执行的函数
• 核函数用__global__符号申明，在调用时需要用<<<grid, block>>>来指定kernel要执行的线程数量
• 在CUDA中，每一个线程都要执行核函数，每个县城分配一个唯一的线程号（thread ID），这个ID值可以通过核函数的内置变量threadldx获得

//Kernel 定义
__global__ void vec_add(double *x, double *y, double *z, int n)
{
   
	int i = get_tid();// 全局变量i；user-defined macro/function
	if(i<n) z[i] = x[i] + y[i];// 向量相加
}
int main()
{
   
	int N = 1000000;// 1M
	int bs = 256;// 每个块有256线程
	int gs = (N + bs -1)/bs;// 网格
	vec_add<<<gs,bs>>>(x, y, z, N);// kernel, call GPU，指定使用的线程数和结构

	}
}

2.2 程序层次结构

• 第一层：一个kernel所启动的所有线程成为一个网格（grid）
  • 同一个网格上的线程，共享相同的全局内存
  • dim3类型
• 第二层：网格可以分为很多线程块（block），一个线程块包含很多线程
  • dim3类型
• 第三层：32线程一组为线程束（warp ）

在CUDA架构中，线程束是指一个包含32个线程的集合，这个线程集合被“编织在一起”并且以“步调一致(Lockstep)”的形式执行。在程序中的每一行，线程束中的每个线程都将在不同的数据上执行相同的指令。

dim3类型 包含三个无符号整数(x, y, z)成员的结构体变量，在定义时，缺省值初始化为1；1-dim、2-dim、3-dim结构均可

global：在device上执行，从host中调用，返回类型必须是void，不支持可变参数，不能成为类成员函数；异步，host会在kernel执行时进行下一步
device：在device上执行，单次可以从device中调用；不可以与__global__同时用
host：在host上执行，仅可以从host上调用，一般省略；不可以与__global__同时用；可以与__device__同时用，此时函数会在devicw和host都编译

2.3 【常用】CUDA内置变量

一个线程需要两个内置的坐标变量（blockldx， threadldx）来唯一标识，它们都是dim3类型变量。
一个线程块上的线程，放在同一个流式多处理器（SM）上
单个SM资源有限，线程块中的线程数有限制

blockldx 指明线程所在grid中的位置：blockldx.x, blockldx.y, blockldx.z
threadldx 指明线程所在block中的位置：threadldx.x ,threadldx.y, threadldx.z

gridDim 网格块各个维度的大小
blockDim 线程块各个维度的大小
对于一个2-dim的block(Dx,Dy)，线程(x,y)的ID值为(x+y∗Dx)，如果是3-dim的block(Dx,Dy,Dz)，线程(x,y,z)的ID值为(x+y∗Dx+z∗Dx∗Dy)。

/* get thread id: 1D block and 2D grid */
# define get_tid()(blockDim.x*(blockldx.x + blockldx.y*gridDim.x) + threadldx.x)

/* get block id: 2D grid */
# define get_bid()(blockldx.x + blockldx.y*gridDim.x)

__global__void vec_add(double *x, double *y, double *z, int n)
{
   
	int i = get_tid();// user-defined function
	if(i<n)z[i] = x[i] + y[i];
}

3 并行编程

3.1 线程并行

#include "cuda_runtime.h"			//CUDA运行时API
#include "device_launch_parameters.h"	
#include <stdio.h>
cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size);
__global__ void addKernel(int *c, const int *a, const int *b)
{
   
    int i = threadIdx.x;
    c[i] = a[i] + b[i];
}
int main()
{
   
    const int arraySize = 5;
    const int a[arraySize] = {
    1, 2, 3, 4, 5 };
    const int b[arraySize] = {
    10, 20, 30, 40, 50 };
    int c[arraySize] = {
    0 };
    // Add vectors in parallel.
    cudaError_t cudaStatus;
	int num = 0;
	cudaDeviceProp prop;
	cudaStatus = cudaGetDeviceCount(&num);
	for(int i = 0;i<num;i++)
	{
   
		cudaGetDeviceProperties(&prop,i);
	}
	cudaStatus = addWithCuda(c, a, b, arraySize);
    if (cudaStatus != cudaSuccess) 
	{
   
        fprintf(stderr, "addWithCuda failed!");
        return 1;
    }
    printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",c[0],c[1],c[2],c[3],c[4]);
    // cudaThreadExit must be called before exiting in order for profiling and
    // tracing tools such as Nsight and Visual Profiler to show complete traces.
    cudaStatus = cudaThreadExit();
    if (cudaStatus != cudaSuccess) 
	{
   
        fprintf(stderr, "cudaThreadExit failed!");
        return 1;
    }
    return 0;
}
// 重点理解这个函数
cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size)
{
   
    int *dev_a = 0;	//GPU设备端数据指针
    int *dev_b = 0;
    int *dev_c = 0;
    cudaError_t cudaStatus;		//状态指示
 
    // Choose which GPU to run on, change this on a multi-GPU system.
    cudaStatus = cudaSetDevice(0);	//选择运行平台
    if (cudaStatus != cudaSuccess) 
	{
   
        fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
        goto Error;
    }
    // 分配GPU设备端内存
    cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
    if (cudaStatus != cudaSuccess) 
	{
   
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }
    cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
    if (cudaStatus != cudaSuccess) 
	{
   
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }
    cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));
    if (cudaStatus != cudaSuccess) 
	{
   
        fprintf(stderr, "cudaMalloc failed!"