cuda 函数传地址_CUDA编程实战-向量加法

1 较短矢量的求和

#include <iostream>

#define N 10

__global__ void ArrayAdd(int* a, int* b, int* c) {
    // thread id, GPU中的线程ID，不同线程处理不同数据
    // blockIdx是cuda运行时定义的内置变量，用于标识线程ID
    int tid = blockIdx.x;
    if (tid < N) c[tid] = a[tid] + b[tid];
}

int main() {
    // 分配CPU内存
    int a[N] = {}, b[N] = {}; c[N] = {};
    // 位于CPU内存上的指向GPU内存地址的指针
    int *dev_a = nullptr, *dev_b = nullptr, *dev_c = nullptr;
    
    // 生成输入数据
    for (int idx = 0; idx < N; ++idx) {
        a[idx] = 3 * idx;
        b[idx] = idx * idx;
    }
        
    // 分配GPU内存
    cudaMalloc((void**)&dev_a, N * sizeof(int));
    cudaMalloc((void**)&dev_b, N * sizeof(int));
    cudaMalloc((void**)&dev_c, N * sizeof(int));

    // 从CPU内存到GPU内存拷贝输入数据
    cudaMemcpy(dev_a, a, N * sizeof(int), cudaMemcpyHostToDevide);
    cudaMemcpy(dev_b, b, N * sizeof(int), cudaMemcpyHostToDevide);
    
    // 执行GPU程序
    ArrayAdd<<<N, 1>>>(dev_a, dev_b, dev_c);
    
    // 从GPU内存到CPU内存拷贝输出数据
    cudaMemcpy(c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost);
    
    for (int idx = 0; idx < N; ++idx) {
        printf("%d + %d = %d", a[idx], b[idx], c[idx];
    }
    
    // 释放分配的GPU内存
    cudaFree(dev_a); dev_a = nullptr;
    cudaFree(dev_b); dev_b = nullptr;
    cudaFree(dev_c); dev_c = nullptr;
    
    return 0;
}

• kernel<<<grid dim, block dim>>>
  • grid dim是二维的block网格
  • block dim是三维的thread数组
• 核函数的索引计算方法
  • blockIdx: grid中block的id，分为x, y坐标
  • threadIdx: block中thread的id，分为x, y, z坐标
  • tid = theadIdx.x + blockIdx.x * blockDim.x （一维grid和一维block的情形）

2 较长矢量的求和

• CUDA编程通常要结合block并行和thread并行
  • block数量受硬件限制，不超过65535
  • 设备属性maxTreadsPerBlock（通常限制每个线程块512个线程）

// 修改核函数调用，保证每个线程块完全利用其中的128个线程，且不启动多余的线程块
ArrayAdd<<<(N + 127) / 128, 128>>>(dev_a, dev_b, dev_c);

3 任意长度矢量的求和

• 按位置拆分数据，形成多个计算任务，每个thead中循环执行分配的计算
• 数据分块，GPU顺序处理多块数据，GPU threads并行处理单块数据

__global__ void ArrayAdd(int* a, int * b, int* c) {
	int tid = blockIdx.x * blockDim.x + threadIdx.x;
	while (tid < N) {
		c[tid] = a[tid] + b[tid];
		tid += blockDim.x * gridDim.x;  // tid递增步长是线程数目
	}
}

// CPU code
ArrayAdd<<<(128, 128>>>(dev_a, dev_b, dev_c);