C++ 使用 __global__
声明说明符定义内核,并使用新的 <<<...>>>
执行配置语法指定内核调用的 CUDA 线程数(请参阅 C++ 语言扩展)。 每个执行内核的线程都有一个唯一的线程 ID,可以通过内置变量在内核中访问。
示例代码使用内置变量 threadIdx
将两个大小为 N 的向量 A 和 B 相加,并将结果存储到向量 C 中:
#include <cuda_runtime_api.h>
#include <iostream>
#define RANDOM(x) (rand() % x)
#define MAX 10
// single block multiple threads
__global__ void vector_add_gpu_2(int *d_a, int *d_b, int *d_c, int n){
int tid = threadIdx.x;
const int t_n = blockDim.x;
while(tid < n){
d_c[tid] = d_a[tid] + d_b[tid];
tid+=t_n;
}
}
int main(){
/***向量相加的实现***/
int n = 5;
int *a = (int *)malloc(sizeof(int)*n);
int *b = (int *)malloc(sizeof(int)*n);
int *c = (int *)malloc(sizeof(int)*n);
for (size_t i = 0; i<n; i++){
a[i] = RANDOM(MAX);
b[i] = RANDOM(MAX);
std::cout << a[i] << " " << b[i] << std::endl;
}
cudaError_t cudaStatus;
// GPU memory allocate
int *d_a, *d_b, *d_c;
cudaMalloc((void **)&d_a, sizeof(int)*n);
cudaMalloc((void **)&d_b, sizeof(int)*n);
cudaMalloc((void **)&d_c, sizeof(int)*n);
// data a and b copy to GPU
cudaStatus = cudaMemcpy(d_a, a, sizeof(int)*n, cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
std::cout << ("Memory copy failed! error code: %s", cudaGetErrorString(cudaStatus)) << std::endl;
}
cudaStatus = cudaMemcpy(d_b, b, sizeof(int)*n, cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
std::cout << ("Memory copy failed! error code: %s", cudaGetErrorString(cudaStatus)) << std::endl;
}
vector_add_gpu_2<<<1, 3>>>(d_a, d_b, d_c, n);
// result copy back to CPU
cudaMemcpy(c, d_c, sizeof(int)*n, cudaMemcpyDeviceToHost);
std::cout << "the result of add is: " << std::endl;
for (size_t i = 0; i<n; i++){
std::cout << c[i] << " ";
}
std::cout << std::endl;
// GPU memory free
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
free(a);
free(b);
free(c);
return 0;
}
思考:
1 VecAdd 被执行几次?
2 怎么被调用的?
3 有多少个线程在执行?线程间数据相互影响吗?