gpu block 划分

转载

mob64ca14144dde 2025-01-01 14:16:46

文章标签 gpu block 划分 cuda CUDA #include #define 文章分类 游戏开发

仅作个人记录

过程中碰到问题，还是建议查看官方API文档
CUDA Runtime API :: CUDA Toolkit Documentation

经过并行 + 连续存取的优化后，速度已经很快了，下面还可以进一步加速。

CUDA不仅提供Thread，还提供了Grid和Block以及Share Memory这些非常重要的机制，1050ti的Thread最大是1024，但通过block和Grid，线程的数量还能成倍增长，甚至几万个线程。

gpu block 划分_gpu block 划分

CUDA 架构下，gpu执行最小单位是thread，多个 thread 可以组成一个block。一个 block 中的 thread 能存取同一块共享内存，而且可以快速进行同步的动作。

每一个 block 能包含的 thread 数目有限。不过，执行相同程序的 block，又可以组成grid（套娃）。不同 block 中的 thread 无法存取同一块共享内存，因此无法直接互通或同步。因此，不同 block 中的 thread 能合作的程度是比较低的。不过也因为如此，可以让程序不用担心gpu实际上能够同时执行的 thread 数目限制。例如，一个执行单元很少的gpu，可能会把各个 block 中的 thread 顺序执行，而非同时执行。不同的 grid 则可以执行不同的程序（即 kernel）。

通过多个block使用更多的线程

修改过程：CUDA编程（六）进一步并行_MingChao_Sun-CSDN博客

#include <iostream>
#include <stdlib.h>
#include <time.h> //用于计时
#include <cuda_runtime.h>
#include <device_launch_parameters.h>

using namespace std;

//32 个 block，每个 block 有 256个 thread，共有 32*256= 8192个thread
#define BLOCK_NUM 32
#define THREAD_NUM 256
#define DATA_SIZE 1048576
int data[DATA_SIZE];

//产生随机数
void generateNum(int *data, int size){
	for (int i = 0; i < size; ++i)	data[i] = rand() % 10;
}

void printDeviceProp(const cudaDeviceProp &prop){
	cout << "Device Name: " << prop.name << endl;
	cout << "totalGlobalMem: " << prop.totalGlobalMem << endl;
	cout << "sharedMemPerBlock: " << prop.sharedMemPerBlock << endl;
	cout << "regsPerBlock: " << prop.regsPerBlock << endl;
	cout << "warpSize: " << prop.warpSize << endl;
	cout << "memPitch: " << prop.memPitch << endl;
	cout << "maxThreadsPerBlock:" << prop.maxThreadsPerBlock << endl;
	cout << "maxThreadsDim[0 - 2]: " << prop.maxThreadsDim[0] << " " << prop.maxThreadsDim[1] << " " << prop.maxThreadsDim[2] << endl;
	cout << "maxGridSize[0 - 2]: " << prop.maxGridSize[0] << " " << prop.maxGridSize[1] << " " << prop.maxGridSize[2] << endl;
	cout << "totalConstMem:" << prop.totalConstMem << endl;
	cout << "major.minor:" << prop.major << " " << prop.minor << endl;
	cout << "clockRate:" << prop.clockRate << endl;
	cout << "textureAlignment:" << prop.textureAlignment << endl;
	cout << "deviceOverlap:" << prop.deviceOverlap << endl;
	cout << "multiProcessorCount:" << prop.multiProcessorCount << endl;
}

//cuda初始化
bool InitCuda(){
	int count;
	cudaGetDeviceCount(&count);//获取能够使用的gpu数量，编号从0开始
	if (count == 0)	return false;//没有支持cuda的gpu
	int device = 0;
	for (; device < count; ++device){
		cudaDeviceProp prop;
		if (cudaGetDeviceProperties(&prop, device) == cudaSuccess){
			printDeviceProp(prop);
			break;//寻找一个可用的gpu
		}
	}
	cudaSetDevice(device);//决定使用编号为device的gpu
	return true;
}

//__global__函数（GPU上执行），计算立方和
__global__ void sum_Squares(int *num, int *result, clock_t *time){
	const int thread_id = threadIdx.x;//当前的线程编号（0开始）
	const int block_id = blockIdx.x;//当前的 thread 属于第几个 block（0 开始）
	const int size = DATA_SIZE / THREAD_NUM;//分配给每个线程的量
	clock_t start;
	if (thread_id == 0)	time[block_id] = clock();//计算时间，只在 threadid ==0 时进行，每个 block 都会记录开始时间及结束时间
	int sum = 0;
	for (int i =block_id*THREAD_NUM + thread_id; i < DATA_SIZE; i += BLOCK_NUM*THREAD_NUM)	sum += num[i] * num[i] * num[i];
	result[block_id*THREAD_NUM + thread_id] = sum;
	if (thread_id == 0) time[block_id+BLOCK_NUM] = clock();
}

int main(){
	if (!InitCuda())	return 0;
	//生成随机数
	generateNum(data, DATA_SIZE);

	int *gpudata, *result;
	clock_t *time;
	//gpu上开内存空间存储数组以及计算结果
	cudaMalloc((void **)&gpudata, sizeof(int)*DATA_SIZE);//第一个参数是指针的指针
	cudaMalloc((void **)&result, sizeof(int)*THREAD_NUM*BLOCK_NUM);//thread,block增多
	cudaMalloc((void **)&time, sizeof(clock_t)*BLOCK_NUM*2);

	//数据从cpu搬运到gpu
	cudaMemcpy(gpudata, data, sizeof(int)*DATA_SIZE, cudaMemcpyHostToDevice);

	//CUDA 中执行函数 语法：函数名称<<<block数目, thread数目, shared memory大小>>>(args...)
	sum_Squares <<<BLOCK_NUM, THREAD_NUM, 0 >>>(gpudata, result, time);//512个线程进行运算

	//运算结果又从gpu搬运回cpu
	int sum[THREAD_NUM*BLOCK_NUM];//进行修改
	clock_t time_cost[BLOCK_NUM*2];
	cudaMemcpy(&sum, result, sizeof(int)*THREAD_NUM*BLOCK_NUM, cudaMemcpyDeviceToHost);
	cudaMemcpy(&time_cost, time, sizeof(clock_t)*BLOCK_NUM * 2, cudaMemcpyDeviceToHost);

	//释放gpu上面开的内存
	cudaFree(gpudata);
	cudaFree(result);
	cudaFree(time);

	int all_sum = 0;//cpu端进行加和
	for (int i = 0; i < THREAD_NUM*BLOCK_NUM; ++i)	all_sum += sum[i];


	//新的计时策略 把每个 block 最早开始时间和最晚结束时间之差，为总运行时间
	clock_t min_start = time_cost[0], max_end = time_cost[BLOCK_NUM];

	for (int i = 1; i < BLOCK_NUM; i++) {
		if (min_start > time_cost[i])
			min_start = time_cost[i];
		if (max_end < time_cost[i + BLOCK_NUM])
			max_end = time_cost[i + BLOCK_NUM];
	}
	cout << "GPU_sum: " << all_sum << " time cost: " << max_end - min_start << endl;

	all_sum = 0;//cpu上面也计算一次进行验证
	for (int i = 0; i < DATA_SIZE; ++i)	all_sum += data[i] * data[i] * data[i];
	cout << "CPU_sum: " << all_sum << endl;

	return 0;
}

运行结果：

gpu block 划分_cuda_02