cuda编程入门笔记

1. Hello Cuda

一个基本的hello cuda程序包含以下三个部分：

1. GPU函数前加 __global__ 前缀，且核函数必须为void类型
2. 调用GPU函数时指定资源: <<<griDim, blockDim>>>
3. 使用同步函数

#include <stdio.h>

 void cpu(){
    printf("hello cpu\n");
}

// 1. 加上 __global__ 前缀
__global__ void gpu(){
    printf("hello gpu\n");
}

int main() {
    cpu();
    // 2. 加上<<<griDim, blockDim>>>  线程块的个数与线程的个数
    gpu<<<1,1>>>();
    // 3. 加上cudaDeviceSynchronize()， 等待GPU代码执行完成，才能在cpu上恢复执行
    cudaDeviceSynchronize();
    return 0;
}

2. GPU基本知识

2.1 基本资源：网格（grid）、块（block）、线程（thread）

  一块GPU拥有多个Grid，一个Grid含有多个Block，一个Block中含有多个Thread，每个Thread都可以并行处理一个程序，下图黑框表示一个Grid， 蓝色框表示Block，白色框表示 Thread

gridDim.x：表示一个Grid里面Block块数量（图示为2）
blockIdx.x：表示当前Block在Grid中的索引
blockDim.x：表示一个Block中x方向上Thread数量
threadIdx.x：表示当前Thread在对应Block中的索引
<<<NUMBER_OF_BLOCK, NUMBER_OF_THREAD_PER_BLOCK>>>
启动核函数时，核函数代码在每个已配置的线程块中的每个线程中都执行；

2.2 创建多线程

  设置 function<<<2, 2>>> 则表示创建两个block，每个block有2个线程来同时执行function函数：

#include <stdio.h>

 void cpu(){
    printf("hello cpu\n");
}

// 1. 加上 __global__ 前缀
__global__ void gpu(){
    printf("hello gpu\n");
}

int main() {

    // 2. 加上<<<griDim, blockDim>>>  线程块的个数与线程的个数
    gpu<<<2,2>>>();	//创建两个Block，每个block有2个thread来执行该函数
    // 3. 加上cudaDeviceSynchronize()， 等待GPU代码执行完成，才能在cpu上恢复执行
    cudaDeviceSynchronize();
    return 0;
}

同时4个线程执行该函数，得到4个输出

---

  若只需要在block 0 的 thread 0 来进行一些操作：就需要使用 blockIdx与 threadIdx，若只在第一个block的第一个thread输出，只需要将函数修改为：

__global__ void gpu(){
	// 在第一个block的第一个thread输出
    if(blockIdx.x ==0 && threadIdx.x==0) {
        printf("hello gpu\n");
    }
}

---

2.3 获取Grid中线程的唯一索引

// 一个block中的线程数量 * 当前block的Id + 当前线程在当前block中的id (将二维索引转换为1维索引)
int threadi = blockDim.x * blockIdx.x + threadIdx.x;	

---

2.4 cuda显存分配

2.4.1 统一内存（UM）分配

cudaMallocManaged(&ptr, size)：
  该函数既可以为cpu预分配内存，也可以为gpu预分配内存，分配UM时，内存尚未驻留在主机或设备上。当主机或设备尝试访问内存时会发生 页错误，此时设备会将数据进行迁移。同理，当GPU尝试访问尚未驻留在显存上的数据时会触发缺页中断并将数据进行迁移，这个过程是耗时的。

#include <stdio.h>
#include <stdlib.h>

int main() {
    // 定义预分配内存的大小
    const int N = 100;
    size_t size = N * sizeof(int);
    
    // 在cpu上预分配内存
    int *a = (int*)malloc(size);
    free(a);
    
    // 在Gpu上预分配统一内存
    int *b;
    cudaMallocManaged(&b, size);    // 可以给cpu预分配内存，也可以给gpu预分配内存
    cudaFree(b);
    
    return 0;
}

数据迁移cudaMemPrefetchAsync(&ptr, size, DeviceId)：
  在即将要在gpu上使用某个数据时提前开辟Gpu显存空间，避免缺页中断导致的效率低问题，同理在Gpu运算完成后，后面通过cpu处理该数据前需要在cpu预分配内存；

#include <stdio.h>
#include <stdlib.h>

int main() {
	// 获取设备id
	int id;
	cudaGetDevice(&id);
	
    // 定义预分配内存的大小
    const int N = 100;
    size_t size = N * sizeof(int);
    
    // 统一分配内存（cpu、gpu都可以使用）
    int *a;
   	cudaMallocManaged(&a, size); 
	
	// cpu上使用a
	// 一系列操作
		
	// 准备在gpu上使用a时，提前在显存上开辟空间（避免缺页中断，导致效率低下）
	cudaMemPrefetchAsync(a, size, id);
	// 在gpu上使用a
	// 一系列操作
	
	// 即将cpu使用a之前，将数据迁移到cpu，防止发生缺页中断
	cudaMemPrefetchAsync(a, size, cudaCpuDeviceId);	
	// 在cpu上使用a
	// 一系列操作

	// 释放空间
    cudaFree(a);
    
    return 0;
}

---

例程

将数组中的每个元素扩大两倍

#include <stdio.h>
#include <iostream>


void init_a(int *a, int len){
    for(int i=0;i<len;++i){
        a[i] = i;
    }
}


// cuda函数专用标志符，且不能有返回值
__global__ void gpu(int* a, int N) {
    // 计算当前的线程索引
    int threadi = blockIdx.x * blockDim.x + threadIdx.x;
    // 每个线程处理一个位置
    if (i < N) {
        a[i] = i*2;
    }
}


int main() {
	// 获取设备id
	int id;
	cudaGetDevice(&id);
	
    // 数组的长度
    const int N = 5;
    // 需要的内存空间
    size_t size = N * sizeof(int);
    int* a;
    // 分配内存空间，cpu和gpu都可以使用
    cudaMallocManaged(&a, size);
    // 初始化数组
    init_a(a, N);
	
	// 提前开辟显存空间
	cudaMemPrefetchAsync(a, size, id);
    // 每个block都有256个线程
    size_t threads_num = 256;
    // 计算需要多少个block，向上取整，让数组的每一位拥有一个独立的线程
    size_t blocks_num = (N + threads_num - 1) / threads_num;
    // 使用cuda计算
    gpu<<<blocks_num, threads_num>>>(a, N);	
	// 将数据迁移到cpu
	cudaMemPrefetchAsync(a, size, cudaCpuDeviceId);	
    // 等待同步
    cudaDeviceSynchronize();
    // 释放内存
    cudaFree(a);
    return 1;
}

输出：

2.4.2 手动内存分配

• cudaMalloc：将直接处于活动状态的 GPU 分配显存，可以防止出现所有的 GPU 分页错误，而代价是主机代码无法访问该命令返回的指针；
• cudaMallocHost: 将直接为 CPU 分配内存，该命令可 “固定” 内存（pinned memory） 或 “锁页”内存（page-lock memory），次举允许将内存异步拷贝至 GPU 或从 GPU 异步拷贝至内存。固定内存过多会干扰 cpu的性能。释放固定内存使用 cudaFreeHost 命令
• 无论是从主机到设备还是从设备到主机，cudaMemcpy命令均可以拷贝(非传输) 内存。

// 在cpu和gpu都注册一块内存
int *host_a, *device_a;
cudaMalloc(&device_a, size);	// 在gpu上分配显存
initializeOnHost(host_a, N);	// 将cpu内存空间初始化

// 将cpu内存上的值host_a拷贝到显存device_a上
cudaMemcpy(device_a, host_a, size, cudaMemcpyHostToDevice);

// 在gpu上执行核函数
kenel_function<<<blocks_num, thread_num, 0, someStream>>>(device_a, N);

// 将显存中的值拷device_a贝到cpu内存host_a中
cudaMemcpy(host_a, device_a, size, cudaMemcpyDeviceToHost);

verifyOnHost(host_a, N);
// 释放内存和显存
cudaFree(device_a);
cudaFreeHost(host_a);

3. 跨步循环

当数据量大于总的分配线程数时，需要将数据切分成块来运行，即以跨步循环的方式运行；如：当数组长度为32而注册的总线程数为8时，需要分4次来 循环处理数组中的元素。

#include <stdio.h>
#include <iostream>


void init_a(int *a, int len){
    for(int i=0;i<len;++i){
        a[i] = i;
    }
}


// cuda函数专用标志符，且不能有返回值
__global__ void gpu(int* a, int N) {
    // 计算当前的线程索引
    int threadi = blockIdx.x * blockDim.x + threadIdx.x;
    // 计算总的线程数量，每次都要跨这么多步（线程0处理完a[0]位置后处理a8]，然后处理a[16], 然后处理a[24]）
    int total_thread = gridDim.x * blockDim.x;
    // 每个线程处理一个位置
    for(int i=threadi;i<N;i+=total_thread){
        a[i]*=2;
    }
}


int main() {	
    // 数组的长度
    const int N = 32;
    // 需要的内存空间
    size_t size = N * sizeof(int);
    int* a_cpu, a_gpu;
    cudaMallocHost(&a_cpu, size); // cpu 上分配内存
    cudaMall(&a_gpu, size);		// gpu 上分配显存
    // cpu上初始化数组
    init_a(a_cpu, N);
	// 将cpu数据copy到gpu中
	cudaMemcpy(a_gpu, a_cpu, size, cudaMemcpyHostToDevice);
	// gpu操作数据
    size_t blocks_num = 2;
    size_t threads_num = 4;
    gpu<<<blocks_num, threads_num>>>(a_gpu, N);
	// gpu数据copy到cpu
	cudaMemcpy(a_cpu, a_gpu, size, cudaMemcpyDeviceToHost);
	// 同步
    cudaDeviceSynchronize();
    // 释放内存
    cudaFree(a_gpu);
    cudaFreeHost(a_cpu);
    return 1;
}

4. 异常处理

#include <stdio.h>
#include <iostream>
#include <assert.h>


inline cudaError_t checkCuda(cudaError_t result){
    if(result != cudaSuccess){
        fprintf(stderr, "CUDA runtime error: %s\n", cudaGetErrorString(result));
        assert(result == cudaSuccess);
    }
    return result;
}

int main() {
    // 数组的长度
    const int N = 32;
    // 需要的内存空间
    size_t size = N * sizeof(int);
    int* a;
    // 检查cuda函数是否出错
    checkCuda(cudaMallocManaged(&a, size));
	return 0;
}

---

5. 矩阵加法

#include <stdio.h>
#include <cuda_runtime.h>


#define N 64	// 矩阵大小为 64*64

__global__ void gpu(int *a, int *b, int *c_gpu) {
	// 找出线程对应的坐标位置
	int r = blockDim.x * blockIdx.x + threadIdx.x;
	int c = blockDim.y * blockIdx.y + threadIdx.y;
	if (r < N && c < N) {
		c_gpu[r * N + c] = a[r * N + c] + b[r * N + c];
	}
}

void cpu(int* a, int* b, int* c_cpu) {
	for (int r = 0; r < N; r++) {
		for (int c = 0; c < N; c++) {
			c_cpu[r * N + c] = a[r * N + c] + b[r * N + c];
		}
	}
}



int main() {
	// 定义两个数组和结果存储数组
	int* a, * b, * c_cpu, * c_gpu;
	size_t size = N * N * sizeof(int);
	cudaMallocManaged(&a, size);
	cudaMallocManaged(&b, size);
	cudaMallocManaged(&c_cpu, size);
	cudaMallocManaged(&c_gpu, size);
	// 数组初始化
	for (int i = 0; i < N; ++i) {
		for (int j = 0; j < N; ++j) {
			a[i * N + j] = 1;
			b[i * N + j] = 2;
			c_cpu[i * N + j] = 0;
			c_gpu[i * N + j] = 0;
		}
	}
	
	// 获取设备id
	int id;
	cudaGetDevice(&id);
	// 提前在显存上开辟空间
	cudaMemPrefetchAsync(a, size, id);
	cudaMemPrefetchAsync(b, size, id);

	// 定义block中每个方向的线程的数量（分为两个方向，为了方便计算矩阵）
	dim3 threads(16, 16, 1);
	// 计算需要每个方向上要开多少个blocks，才能得到64*64个线程来处理两个矩阵相加
	dim3 blocks((N + threads.x - 1) / threads.x, (N + threads.x - 1) / threads.x, 1);
	// 使用cpu计算两个矩阵加法
	cpu(a, b, c_cpu);
	// 使用gpu计算矩阵加法
	gpu << <blocks, threads >> > (a, b, c_gpu);
	// 同步
	cudaDeviceSynchronize();
	// 释放空间
	cudaFree(a);
	cudaFree(b);
	cudaFree(c_cpu);
	cudaFree(c_gpu);

	return 1;
}

6. 流

cuda中的流遵循以下规则：

• 指定流中的所有操作会按序执行；
• 非默认流之间的核函数的执行顺序是无法保证的；
• 默认流具有阻断能力，会等待其他流完成当前的操作后开始执行和函数，期间其他流无法执行和函数，直到默认流的该核函数执行完毕；
• 

sudaStream_t stream;	// 创建流
cudaStreamCreate(&stream);	
// 执行核函数
kernel_function<<<blocks_num, threads_num, 0, stream>>>();

// 销毁流
cudaStreamDestroy(stream);