CUDA中block和thread的合理划分配置

最新推荐文章于 2024-09-18 14:48:41 发布
最新推荐文章于 2024-09-18 14:48:41 发布
阅读量3.3w 收藏 65
点赞数 4
分类专栏： CUDA 文章标签： CUDA OpenCV Grid Block Thread
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/dcrmg/article/details/54850766
版权

CUDA并行编程的基本思路是把一个很大的任务划分成N个简单重复的操作，创建N个线程分别执行执行，每个网格（Grid）可以最多创建65535个线程块，每个线程块（Block）一般最多可以创建512个并行线程，在第一个CUDA程序中对核函数的调用是：


addKernel<<<1, size>>>(dev_c, dev_a, dev_b);


这里的<<<>>>运算符内是核函数的执行参数，告诉编译器运行时如何启动核函数，用于说明内核函数中的线程数量，以及线程是如何组织的。


<<<>>>运算符完整的执行配置参数形式是<<<Dg, Db, Ns, S>>>

  • 参数Dg用于定义整个grid的维度和尺寸，即一个grid有多少个block。为dim3类型。Dim3 Dg(Dg.x, Dg.y, 1)表示grid中每行有Dg.x个block，每列有Dg.y个block，第三维恒为1(目前一个核函数只有一个grid)。整个grid中共有Dg.x*Dg.y个block，其中Dg.x和Dg.y最大值为65535。
  • 参数Db用于定义一个block的维度和尺寸，即一个block有多少个thread。为dim3类型。Dim3 Db(Db.x, Db.y, Db.z)表示整个block中每行有Db.x个thread，每列有Db.y个thread，高度为Db.z。Db.x和Db.y最大值为512，Db.z最大值为62。 一个block中共有Db.x*Db.y*Db.z个thread。计算能力为1.0,1.1的硬件该乘积的最大值为768，计算能力为1.2,1.3的硬件支持的最大值为1024。
  • 参数Ns是一个可选参数，用于设置每个block除了静态分配的shared Memory以外，最多能动态分配的shared memory大小，单位为byte。不需要动态分配时该值为0或省略不写。
  • 参数S是一个cudaStream_t类型的可选参数，初始值为零，表示该核函数处在哪个流之中。
在第一个CUDA程序中使用了1个线程块，每个线程块包含size个并行线程，每个线程的索引是threadIdx.x。
也可以选择创建size个线程块，每个线程块包含1个线程，核函数的调用更改为：

addKernel<<<size, 1>>>(dev_c, dev_a, dev_b);

线程的索引更改为blockIdx.x。完整程序如下： 

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size);

__global__ void addKernel(int *c, const int *a, const int *b)
{
	int i = blockIdx.x;
	c[i] = a[i] + b[i];
}

int main()
{
	const int arraySize = 5;
	const int a[arraySize] = { 1, 2, 3, 4, 5 };
	const int b[arraySize] = { 10, 20, 30, 40, 50 };
	int c[arraySize] = { 0 };

	// Add vectors in parallel.
	cudaError_t cudaStatus = addWithCuda(c, a, b, arraySize);
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "addWithCuda failed!");
		return 1;
	}

	printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",
		c[0], c[1], c[2], c[3], c[4]);

	// cudaDeviceReset must be called before exiting in order for profiling and
	// tracing tools such as Nsight and Visual Profiler to show complete traces.
	cudaStatus = cudaDeviceReset();
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaDeviceReset failed!");
		return 1;
	}
	getchar();
	return 0;
}

// Helper function for using CUDA to add vectors in parallel.
cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size)
{
	int *dev_a = 0;
	int *dev_b = 0;
	int *dev_c = 0;
	cudaError_t cudaStatus;

	// Choose which GPU to run on, change this on a multi-GPU system.
	cudaStatus = cudaSetDevice(0);
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
		goto Error;
	}

	// Allocate GPU buffers for three vectors (two input, one output)    .
	cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaMalloc failed!");
		goto Error;
	}

	cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaMalloc failed!");
		goto Error;
	}

	cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaMalloc failed!");
		goto Error;
	}

	// Copy input vectors from host memory to GPU buffers.
	cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaMemcpy failed!");
		goto Error;
	}

	cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaMemcpy failed!");
		goto Error;
	}

	// Launch a kernel on the GPU with one thread for each element.
	//addKernel<<<1, size>>>(dev_c, dev_a, dev_b);
	addKernel << <size, 1 >> > (dev_c, dev_a, dev_b);

	// Check for any errors launching the kernel
	cudaStatus = cudaGetLastError();
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
		goto Error;
	}

	// cudaDeviceSynchronize waits for the kernel to finish, and returns
	// any errors encountered during the launch.
	cudaStatus = cudaDeviceSynchronize();
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
		goto Error;
	}

	// Copy output vector from GPU buffer to host memory.
	cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaMemcpy failed!");
		goto Error;
	}

Error:
	cudaFree(dev_c);
	cudaFree(dev_a);
	cudaFree(dev_b);

	return cudaStatus;
}

执行结果一致：



更普遍的情况是需要创建多个线程块，每个线程块包含多个并行线程，这种情况下线程索引的计算为：

int tid=threadIdx.x+blockIdx.x*blockDim.x;

blockIdx代表线程块在网格中的索引值，blockDim代表线程块的尺寸大小，另外还有gridDim代表网格的尺寸大小。

如果有N个并行的任务，我们希望每个线程块固定包含6个并行的线程，则可以使用以下的核函数调用：

addKernel<<<(N+5)/6, 6>>>(dev_c, dev_a, dev_b);

把第一个CUDA程序的向量个数增加到15个，修改成以上调用方式： 

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size);

__global__ void addKernel(int *c, const int *a, const int *b)
{
	int i = threadIdx.x + blockIdx.x*blockDim.x;
	if (i < 15)
		c[i] = a[i] + b[i];
}

int main()
{
	const int arraySize = 15;
	const int a[arraySize] = { 1, 2, 3, 4, 5,6,7,8,9,10,11,12,13,14,15 };
	const int b[arraySize] = { 10, 20, 30, 40, 50,60,70,80,90,100,110,120,130,140,150 };
	int c[arraySize] = { 0 };

	// Add vectors in parallel.
	cudaError_t cudaStatus = addWithCuda(c, a, b, arraySize);
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "addWithCuda failed!");
		return 1;
	}
	printf("{ 1, 2, 3, 4, 5,6,7,8,9,10,11,12,13,14,15}+\n{ 10, 20, 30, 40, 50,60,70,80,90,100,110,120,130,140,150}=\n{%d,%d,%d,%d,%d,%d,%d,%d,%d,%d,%d,%d,%d,%d,%d}\n",
		c[0], c[1], c[2], c[3], c[4], c[5], c[6], c[7], c[8], c[9], c[10], c[11], c[12], c[13], c[14]);

	// cudaDeviceReset must be called before exiting in order for profiling and
	// tracing tools such as Nsight and Visual Profiler to show complete traces.
	cudaStatus = cudaDeviceReset();
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaDeviceReset failed!");
		return 1;
	}
	getchar();
	return 0;
}

// Helper function for using CUDA to add vectors in parallel.
cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size)
{
	int *dev_a = 0;
	int *dev_b = 0;
	int *dev_c = 0;
	cudaError_t cudaStatus;

	// Choose which GPU to run on, change this on a multi-GPU system.
	cudaStatus = cudaSetDevice(0);
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
		goto Error;
	}

	// Allocate GPU buffers for three vectors (two input, one output)    .
	cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaMalloc failed!");
		goto Error;
	}

	cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaMalloc failed!");
		goto Error;
	}

	cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaMalloc failed!");
		goto Error;
	}

	// Copy input vectors from host memory to GPU buffers.
	cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaMemcpy failed!");
		goto Error;
	}

	cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaMemcpy failed!");
		goto Error;
	}

	// Launch a kernel on the GPU with one thread for each element.
	//addKernel<<<1, size>>>(dev_c, dev_a, dev_b);
	addKernel << <(size + 5) / 6, 6 >> > (dev_c, dev_a, dev_b);

	// Check for any errors launching the kernel
	cudaStatus = cudaGetLastError();
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
		goto Error;
	}

	// cudaDeviceSynchronize waits for the kernel to finish, and returns
	// any errors encountered during the launch.
	cudaStatus = cudaDeviceSynchronize();
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
		goto Error;
	}

	// Copy output vector from GPU buffer to host memory.
	cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, "cudaMemcpy failed!");
		goto Error;
	}

Error:
	cudaFree(dev_c);
	cudaFree(dev_a);
	cudaFree(dev_b);

	return cudaStatus;
}

执行结果：


以下CUDA和OpenCV混合编程，对一幅图像上每个像素点的颜色执行一次运算，生成一幅规则的图形。
新建了一个 dim3类型的变量grid(DIM, DIM)，代表一个二维的网格，尺寸大小是DIM*DIM个线程块: 

#include "cuda_runtime.h"  
#include <highgui.hpp>  

using namespace cv;

#define DIM 600   //图像长宽  

__global__ void kernel(unsigned char *ptr)
{
	// map from blockIdx to pixel position  
	int x = blockIdx.x;
	int y = blockIdx.y;
	int offset = x + y * gridDim.x;

	//BGR设置
	ptr[offset * 3 + 0] = 999 * x*y % 255;
	ptr[offset * 3 + 1] = 99 * x*x*y*y % 255;
	ptr[offset * 3 + 2] = 9 * offset*offset % 255;
}

// globals needed by the update routine  
struct DataBlock
{
	unsigned char   *dev_bitmap;
};

int main(void)
{
	DataBlock   data;
	cudaError_t error;

	Mat image = Mat(DIM, DIM, CV_8UC3, Scalar::all(0));

	data.dev_bitmap = image.data;
	unsigned char    *dev_bitmap;

	error = cudaMalloc((void**)&dev_bitmap, 3 * image.cols*image.rows);
	data.dev_bitmap = dev_bitmap;

	dim3    grid(DIM, DIM);
	//DIM*DIM个线程块
	kernel <<<grid, 1 >>> (dev_bitmap);

	error = cudaMemcpy(image.data, dev_bitmap,
		3 * image.cols*image.rows,
		cudaMemcpyDeviceToHost);

	error = cudaFree(dev_bitmap);

	imshow("CUDA Grid/Block/Thread)", image);
	waitKey();
}

执行效果：


-牧野-
关注
专栏目录
CUDA中Block大小的选择
学以致用 知行合一
02-20 1422
这里介绍两种比较官方的做法，这两种做法的原理是一致的，都是基于SM占用率最大化的原则。1、API函数--cudaOccupancyMaxPotentialBlockSize。注意：对于二维数据，根据计算出的blocksize，自己调整为二维格式，例如blocksize=1024，则二维dim3 blocksize(32,32);它可以计算出在SM占用率最大时blocksize的大小，可以根据此值配置Block。这样就可以不去关心各种硬件资源的限制，那确实是一件很麻烦的事情。2、占用率计算器的方法。
CUDA 编程(1):使用Grid 和 Block分配线程
@bangbang的博客
06-09 545
核函数以线程为单位进行计算的函数，cuda编程会涉及到大量的线程(thread),几千个到几万个thread同时并行计算，所有的thread其实都是在执行同一个核函数。。之所以这么划分，比如将Grid划分为多维的Block，Block划分为多维thread，它其实是为了帮助你快速索引对应的thread。Grid 和Block是逻辑上意义的概念，它并不是在你的GPU硬件上实际存在的。
CUDA执行模型
A11en3的博客
04-28 1008
GPU的架构是围绕一个流式多处理器（SM）的可扩展阵列搭建的。可以通过复制这种架构的构建块来实现GPU的硬件并行。SM是GPU中的一个重要组件，负责执行CUDA核函数中的线程块（Thread Block）。每个SM包含多个CUDA核心，可以同时执行多个线程块中的线程，以实现并行计算。每个CUDA核心有自己的缓存器，也可以访问同一个SM中的共享内存。每个CUDA核心可以并行多个线程。简单来说，一个GPU可以并行的最大线程数=CUDA核心数*一个CUDA核心可以并行的线程数。
CUDA结构之grid block thread SM SP warp 随笔
最新发布
weixin_52173987的博客
09-18 2094
CUDA的线程模型
CUDA 关于 BLOCK数目与Thread数目设置
weixin_30650859的博客
08-30 1908
GPU的计算核心是以一定数量的StreamingProcessor(SP)组成的处理器阵列，NV称之为TextureProcessingClusters(TPC)，每个TPC中又包含一定数量的StreamingMulti-Processor(SM)，每个SM包含8个SP。SP的主要结构为一个ALU（逻辑运算单元），一个FPU（浮点运算单元）以及一个Regi...
cuda thread和block值怎么配置（个人经验）
蜗蜗的专栏
06-19 280
详见CUDA_C_Programming_Guide.pdf，Chapter 19. Compute Capabilities。1)一个block里的thread的总数是32的倍数，是1536的约数；2）一个block里的thread的总数 能够被 待处理数据量整除；3）block的数量受限于寄存器分配和共享内存；1.查询自己的GPU算力；2.对应算力查找资源限制；4.thread限制。
How do CUDA blocks/warps/threads map onto CUDA cores?
十年磨一剑
10-31 365
https://stackoverflow.com/questions/10460742/how-do-cuda-blocks-warps-threads-map-onto-cuda-cores http://www.nvidia.com/content/PDF/fermi_white_papers/NVIDIA_Fermi_Compute_Architecture_Whitepaper
【CUDA】grid、block、thread的关系及thread索引的计算
热门推荐
hujingshuang
11-09 3万+
CUDA中grid、block、thread的关系及thread索引的计算
CUDA如何选取Blocks和Threads
李宏宗的博客
11-11 1446
Blocks 早期的卡，最多支持65535个blocks, 后面的卡最多可以支持231−1=21474836472^{31}-1=2147483647231−1=2147483647 可以理解为blocks可以随便造，主要是要把threads的大小选好 Threads 一般来讲每一个block不能超过512（Compute Capability 1.x）/ 1024（2.x和之后版本）个线程 你电脑不旧的话，一般都是不超过1024 每个块的线程数应该是warp size大小的整数倍，在所有
使用二维网格和二维块对矩阵求和_cuda矩阵求和_cuda_矩阵求和_cuda并行_
10-03
在CUDA编程中，矩阵求和是一个常见的操作，特别是在高性能计算和大数据分析领域。CUDA（Compute Unified Device Architecture）是由NVIDIA提出的GPU并行计算框架，它允许程序员直接利用GPU的并行处理能力，以加速...
cuda7.0：主机多线程流实现kernel并行
06-29
cuda7.0：新功能实现主机多线程多流之间 实现kernel并行。主机的每个线程分配一个流 采用这个新特性可以实现多流之间的kernel并发执行
GPU中基于CUDA计算SM中active block和active warp及普遍优化流程
GabrielGu的博客
10-23 3851
CUDA优化第一要义是设计block和thread的dimension，这时active block和active warp就必须要考虑了。（一维的同学请回避谢谢） 首先通俗地讲一些CUDA优化的常识吧 accurancy：只在关键步骤使用双精度，其他部分仍然使用单精度浮点以获得指令吞吐量和精度的平衡；建议将双精度放在CPU里计算。异构并行会给项目加分哦 latency：降低延迟需要考虑缓冲
高性能计算平台建设——浅谈 OpenMP 和 CUDA 的异同点
AI天才研究院
08-27 543
随着计算机系统的不断发展、应用领域的扩展、硬件的普及，越来越多的人们期盼着能够实现更快的运算能力，更大的计算容量。同时，由于科技的进步，对算法设计要求也越来越高。所以，如何更好地充分利用现代计算机硬件资源、提升并行编程能力成为热门话题。最近两年来，以 NVIDIA CUDA 为代表的并行编程模型和框架已经成为非常热门的研究方向。它是一种面向GPU的并行编程模型，可以用来进行复杂的数据并行计算任务，尤其适合于高性能计算领域。
GPU 架构与 CUDA 关系 并行计算平台和编程模型 CUDA 线程层次结构 GPU 的算力是如何计算的 算力峰值
万有文的博客
04-28 1387
本文主要包含 NVIDIA GPU 硬件的基础概念、CUDA（Compute Unified Device Architecture）并行计算平台和编程模型，详细讲解 CUDA 线程层次结构，最后将讲解 GPU 的算力是如何计算的，这将有助于计算大模型的算力峰值和算力利用率。GPU 硬件基础概念GPU 架构与 CUDA 的关系紧密相连，两者共同构成了 NVIDIA 在并行计算领域的核心技术。
【cuda编程】CUDA的运行方式以及grid、block结构关系
非晚非晚的博客
10-26 4033
从GPU至线程的关系依次为：显卡（GPU）->网格（grid）->线程块（block）->线程（thread） 。从网格开始最大为3维，当然也可以1维了。 网格(grid)：一个内核函数(kernel)就是一个网格，里面所有线程都在这个网格范围内，里面的线程共享全局内存空间 线程块(block)：一个网格可以包含很多个block，block之间可以通过“同步”和“共享内存”进行协作，block之间的区分通过“blockIdx” 线程(thread)：一个线程块可以包含很多个thread，thread之间区
cuda 线程调度
doublewei的专栏
03-28 2403
硬件描述 从硬件上看，一块显卡的最小单元是GPU核(或者叫做Stream Processor),所有核心平均分配在多个SM中，而多个SM共同构成整块显卡的核心。比如RTX2070有36个Streaming Multiprocessors，而每个SM有64个CUDA Cores，RTX2070具有36*64=2304个CUDA Cores。 软件描述 thread，block，grid，war...
cuda block grid等介绍
MakeYouClimax的博客
06-02 652
比如，我们可以定义一个2D线程块，其中每个线程块由多个2D线程组成，每个线程处理一个小块或者一个小行/列。对于单个块的情况，可以使用共享内存来提高访问速度，以便线程块中的所有线程都可以快速、有效地访问它。同时，在使用共享内存时，需要确保线程块中的线程都能够正常读写，以避免竞争条件的发生。在使用 CUDA 进行矩阵乘法等大规模矩阵计算时，通常需要将输入矩阵按照块的形式划分为多个子矩阵，然后分配到不同的线程块中去计算，以充分利用GPU并行计算的优势。申请内存时是线性的内存，需要知道是按行还是按列排列。
cuda中grid，block、thread
03-16
在CUDA中，grid、block和thread是三个重要的概念。 Grid是一个由多个block组成的二维数组，用于描述并行计算的总体规模。每个block都可以包含多个thread。 Block是一个由多个thread组成的三维数组，用于描述并行计算的一个小规模。每个block中的thread可以协同工作，共同完成一个任务。 Thread是CUDA中最小的执行单元，每个thread都可以独立执行一些计算任务。在一个block中，所有的thread都可以访问同一块共享内存，并且可以通过同步机制来协同工作。 通过合理的设置grid、block和thread的数量，可以充分利用GPU的并行计算能力，提高程序的执行效率。
评论 6
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值