cuda编程笔记

一个Grid中含有多个Block，一个Block中含有多个thread

gridDim.x表示网格的块数量
blockIdx.x表示当前块的索引
blockDim.x表示一个块中的线程数量
threadIdx.x表示当前块中线程的索引
<<>>
启动核函数时，核函数代码由每个已配置的线程块中的每个线程执行；

#include 

// cuda函数专用标志符，且不能有返回值
__global__ void gpu(int* a, int N) {
	// 计算当前的线程标志（线程会从0开到总的可用线程数）
	int i = blockIdx.x * blockDim.x + threadIdx.x;
	// 如果还没开够N个线程，就继续开线程运算
	if (i < N) {
		a[i]=i*2;
	}
}

void check(int* a, int N) {
	for (int i = 0; i < N; i++) {
		if (a[i] != i * 2) {
			printf("errorn");
			return;
		}
	}
	printf("okn");
}


int main() {
	// 数组的长度
	const int N = 100;
	// 需要的内存空间
	size_t size = N * sizeof(int);
	int* a;
	// 分配内存空间，cpu和gpu都可以使用
	cudaMallocManaged(&a, size);
	
	// 每个block都有256个线程
	size_t threads = 256;
	// 计算需要多少个block，向上取整
	size_t blocks = (N + threads - 1) / threads;
	// 使用cuda计算
	gpu<<>>(a, N);
	// 等待同步
	cudaDeviceSynchronize();
	// 检查结果是否正确
	check(a, N);
	// 释放内存
	cudaFree(a);
	return 1;
}

当数据量大于总的分配线程数时，需要将数据切分成块来运行，也就是跨步循环运行；

#include 

// cuda函数专用标志符，且不能有返回值
__global__ void gpu(int* a, int N) {
	// 记录当前线程标记（线程会从0开到总的可用线程数）
	int threadi = blockIdx.x * blockDim.x + threadIdx.x;
	// 计算总的可用线程数
	int stride = gridDim.x * blockDim.x;
	// 跨步循环，一个线程跨步处理多个数据
	for (size_t i = threadi; i < N; i += stride) {
		a[i] = i * 2;
	}
}

void check(int* a, int N) {
	for (int i = 0; i < N; i++) {
		if (a[i] != i * 2) {
			printf("errorn");
			return;
		}
	}
	printf("okn");
}


int main() {
	// 数组的长度
	const int N = 1000;
	// 需要的内存空间
	size_t size = N * sizeof(int);
	int* a;
	// 分配内存空间，cpu和gpu都可以使用
	cudaMallocManaged(&a, size);
	
	// 每个block都有256个线程
	size_t threads = 256;
	// 只开辟一个block
	size_t blocks = 1;
	// 使用cuda计算
	gpu<<>>(a, N);
	// 等待同步
	cudaDeviceSynchronize();
	// 检查结果是否正确
	check(a, N);
	// 释放内存
	cudaFree(a);
	return 1;
}

用二维矩阵的方式来定义blocks和threads，在cpu上用完数据，如果不提前开辟显存空间，会发生缺页错误，因此在cpu用完数据之后，在gpu使用之前，提前 开辟显存，将数据预先存储到显存中。

#include 
#include 


#define N 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];
		}
	}
}

void check(int* c_cpu, int* c_gpu) {
	for (int r = 0; r < N; r++) {
		for (int c = 0; c < N; c++) {
			if (c_cpu[r * N + c] != c_gpu[r * N + c]) {
				printf("error");
				return ;
			}
		}
	}
	printf("ok");
}

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);

	// 定义线程的维度（为了简洁）
	dim3 threads(16, 16, 1);
	// 计算需要开多少个blocks
	dim3 blocks((N + threads.x - 1) / threads.x, (N + threads.x - 1) / threads.x, 1);
	cpu(a, b, c_cpu);
	gpu << > > (a, b, c_gpu);
	cudaDeviceSynchronize();
	check(c_cpu, c_gpu);

	cudaFree(a);
	cudaFree(b);
	cudaFree(c_cpu);
	cudaFree(c_gpu);

	return 1;
}