CUDA - 如何让线程和内存对应

最新推荐文章于 2025-03-18 09:54:50 发布
原创 最新推荐文章于 2025-03-18 09:54:50 发布 · 466 阅读 · 5 ·
CC 4.0 BY-SA版权
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
文章标签：

#CUDA

前提：

本文的目的就是设置的程序中，每个线程可以负责一个单独的计算任务。帮助学习和理解线程是如何组织的。

本文处理一个二维数据的加法。

数据在内存中的存储

以线性、行为主的方式存储。

例如，一个16*8的一维数组，在内存中是一段连续的128个地址存储该数据

如下图，每个小格子表示一行数据

想要GPU充分发挥他的优点就是每个线程处理不同的数据，避免同一个线程处理多个数据，或者避免线程没有组织的胡乱访问内存。

组织线程模型

二维网格二维线程块 2D grid \ 2D block

如图，不同颜色的方块表示的是一个线程块。因为数组大小是16*8 =128，先定义每个块的维度是（4，4），所以可以计算得到网格的维度是（4，2）。

定义： 

gridsize(4,2)

blocksize(4,4)

目的是为了让线程和数组内存中的分布一一对应。

线程和二维矩阵映射关系

ix = threadIdx.x + blockIdx.x * blockDim.x;
iy = threadIdx.y + blockIdx.y * blockDim.y;

如下图

线程和二维矩阵映射关系

idx = iy * gridDim.x * blockDim.x + ix;

 编写代码如下：实现二维网格二维线程块进行二维数组的加法

#include "cuda_runtime_api.h"
#include <device_launch_parameters.h>
#include <iostream>

static void CheckCudaErrorAux(const char*, unsigned, const char*, cudaError_t);
#define CUDA_CHECK_RETURN(value) CheckCudaErrorAux(__FILE__,__LINE__, #value, value)

__global__ void addMatrix(int* A, int* B, int* C, const int nx, const int ny)
{
	int ix = threadIdx.x + blockIdx.x * blockDim.x;
	int iy = threadIdx.y + blockIdx.y * blockDim.y;
	int idx = iy * gridDim.x * blockDim.x + ix;

	if (ix < nx && iy < ny)
	{
		C[idx] = A[idx] + B[idx];
	}
}

int main()
{
	const int nx = 16;
	const int ny = 8;
	const int nxy = nx * ny;
	size_t stBytesCount = nxy * sizeof(int);
	int* ipHost_A = new int[nxy];
	int* ipHost_B = new int[nxy];
	int* ipHost_C = new int[nxy];

	for (size_t i = 0; i < nxy; i++)
	{
		ipHost_A[i] = i;
		ipHost_B[i] = i + 1;
	}
	memset(ipHost_C, 0, stBytesCount);
	
	int* ipDevice_A, * ipDevice_B, * ipDevice_C;
	CUDA_CHECK_RETURN(cudaMalloc((void**)&ipDevice_A, stBytesCount));
	CUDA_CHECK_RETURN(cudaMalloc((void**)&ipDevice_B, stBytesCount));
	CUDA_CHECK_RETURN(cudaMalloc((void**)&ipDevice_C, stBytesCount));

	CUDA_CHECK_RETURN(cudaMemcpy(ipDevice_A, ipHost_A, stBytesCount, cudaMemcpyHostToDevice));
	CUDA_CHECK_RETURN(cudaMemcpy(ipDevice_B, ipHost_B, stBytesCount, cudaMemcpyHostToDevice));

	dim3 block(4,4);
	dim3 grid((nx + block.x - 1) / block.x, (ny + block.y - 1) / block.y);

	addMatrix <<<grid, block >>> (ipDevice_A, ipDevice_B, ipDevice_C, nx, ny);

	CUDA_CHECK_RETURN(cudaMemcpy(ipHost_C, ipDevice_C, stBytesCount, cudaMemcpyDeviceToHost));

	for (size_t i = 0; i < nxy; i++)
	{
		if (i % 4 == 0 && i)
			std::cout << std::endl;

		std::cout << ipHost_A[i] << " + " << ipHost_B[i] << " = " << ipHost_C[i] << "\t";

	}

	cudaFree(ipDevice_A);
	cudaFree(ipDevice_B);
	cudaFree(ipDevice_C);

	delete []ipHost_A;
	delete []ipHost_B;
	delete []ipHost_C;

	ipHost_A =nullptr;
	ipHost_B =nullptr;
	ipHost_C =nullptr;


	return 0;
}

static void CheckCudaErrorAux(const char* file, unsigned line, const char* statement, cudaError_t err)
{
	if (err == cudaSuccess)
		return;
	std::cerr << statement << " returned: " << cudaGetErrorName(err) << "  \t : " << cudaGetErrorString(err) << "(" << err << ") at " << file << ":" << line << std::endl;
	exit(1);
}

二维网格一维线程块 2D grid \ 1D block

如图，不同颜色的方块表示的是一个线程块。因为数组大小是16*8 =128，先定义每个块的维度是（4，1），所以可以计算得到网格的维度是（4，8）。

线程和二维矩阵映射关系

这里定义的网格是一维的，所以blockDim.y = 1, threadIdx.y 始终是0

ix = threadIdx.x + blockIdx.x * blockDim.x;
iy = threadIdx.y + blockIdx.y * blockDim.y;

如下图

线程和二维矩阵映射关系

idx = iy * gridDim.x * blockDim.x + ix;

 编写代码如下：实现二维网格一维线程块进行二维数组的加法

#include "cuda_runtime_api.h"
#include <device_launch_parameters.h>
#include <iostream>

static void CheckCudaErrorAux(const char*, unsigned, const char*, cudaError_t);
#define CUDA_CHECK_RETURN(value) CheckCudaErrorAux(__FILE__,__LINE__, #value, value)

__global__ void kernel_addMatrix(int* A, int* B, int* C, const int nx, const int ny)
{
	int ix = threadIdx.x + blockIdx.x * blockDim.x;
	int iy = threadIdx.y + blockIdx.y * blockDim.y;  // 因为block是一维的，所以threadIdx.y始终是0
	int idx = iy * gridDim.x * blockDim.x + ix;

	if (ix < nx && iy < ny)
	{
		C[idx] = A[idx] + B[idx];
	}
}

int main()
{
	const int nx = 16;
	const int ny = 8;
	const int nxy = nx * ny;
	size_t stBytesCount = nxy * sizeof(int);
	int* ipHost_A = new int[nxy];
	int* ipHost_B = new int[nxy];
	int* ipHost_C = new int[nxy];

	for (size_t i = 0; i < nxy; i++)
	{
		ipHost_A[i] = i;
		ipHost_B[i] = i + 1;
	}
	memset(ipHost_C, 0, stBytesCount);

	int* ipDevice_A, * ipDevice_B, * ipDevice_C;
	CUDA_CHECK_RETURN(cudaMalloc((void**)&ipDevice_A, stBytesCount));
	CUDA_CHECK_RETURN(cudaMalloc((void**)&ipDevice_B, stBytesCount));
	CUDA_CHECK_RETURN(cudaMalloc((void**)&ipDevice_C, stBytesCount));

	CUDA_CHECK_RETURN(cudaMemcpy(ipDevice_A, ipHost_A, stBytesCount, cudaMemcpyHostToDevice));
	CUDA_CHECK_RETURN(cudaMemcpy(ipDevice_B, ipHost_B, stBytesCount, cudaMemcpyHostToDevice));

	dim3 block(4, 1);
	dim3 grid((nx + block.x - 1) / block.x, ny);

	kernel_addMatrix <<<grid, block >>> (ipDevice_A, ipDevice_B, ipDevice_C, nx, ny);	

	cudaThreadSynchronize();

	CUDA_CHECK_RETURN(cudaMemcpy(ipHost_C, ipDevice_C, stBytesCount, cudaMemcpyDeviceToHost));

	for (size_t i = 0; i < nxy; i++)
	{
		if (i % 4 == 0 && i)
			std::cout << std::endl;

		std::cout << ipHost_A[i] << " + " << ipHost_B[i] << " = " << ipHost_C[i] << "\t";

	}

	cudaFree(ipDevice_A);
	cudaFree(ipDevice_B);
	cudaFree(ipDevice_C);

	delete[]ipHost_A;
	delete[]ipHost_B;
	delete[]ipHost_C;

	ipHost_A = nullptr;
	ipHost_B = nullptr;
	ipHost_C = nullptr;


	return 0;
}

static void CheckCudaErrorAux(const char* file, unsigned line, const char* statement, cudaError_t err)
{
	if (err == cudaSuccess)
		return;
	std::cerr << statement << " returned: " << cudaGetErrorName(err) << "  \t : " << cudaGetErrorString(err) << "(" << err << ") at " << file << ":" << line << std::endl;
	exit(1);
}

示例结果

一维网格一维线程块 1D grid \ 1D block

之前的GPU线程数和数组的大小是相等的，如果说不相等的情况下，GPU每个线程处理的就不是一个运算，而是多个数据的运算。也就是说在核函数中需要使用循环进行处理。

针对本文示例，网格块定义为（4，1），线程块定义为（4，1）。也就是说每个线程处理的分布如下图：

这个例子中，每个线程需要处理的是一列的数据。

线程和二维矩阵映射关系

ix = threadIdx.x + blockIdx.x * blockDim.x;
iy = threadIdx.y + blockIdx.y * blockDim.y;

这里因为grid\block都是一维的，所以threadIdx.y、blockIdx.y都始终是0.

如下

编码如下：

#include "cuda_runtime_api.h"
#include <device_launch_parameters.h>
#include <iostream>

static void CheckCudaErrorAux(const char*, unsigned, const char*, cudaError_t);
#define CUDA_CHECK_RETURN(value) CheckCudaErrorAux(__FILE__,__LINE__, #value, value)

__global__ void _addMatrix(int* A, int* B, int* C, const int nx, const int ny)
{
	int ix = threadIdx.x + blockIdx.x * blockDim.x;

	int iy = 0; 
	int offset = gridDim.x * blockDim.x;

	if (ix < nx)
	{
		for (size_t i = 0; i < ny; i++)
		{
			int idx = i * offset + ix;
			if (idx < nx*ny)
			{
				C[idx] = A[idx] + B[idx];
			}
		}
	}
}

int main()
{
	const int nx = 16;
	const int ny = 8;
	const int nxy = nx * ny;
	size_t stBytesCount = nxy * sizeof(int);
	int* ipHost_A = new int[nxy];
	int* ipHost_B = new int[nxy];
	int* ipHost_C = new int[nxy];

	for (size_t i = 0; i < nxy; i++)
	{
		ipHost_A[i] = i;
		ipHost_B[i] = i + 1;
	}
	memset(ipHost_C, 0, stBytesCount);

	int* ipDevice_A, * ipDevice_B, * ipDevice_C;
	CUDA_CHECK_RETURN(cudaMalloc((void**)&ipDevice_A, stBytesCount));
	CUDA_CHECK_RETURN(cudaMalloc((void**)&ipDevice_B, stBytesCount));
	CUDA_CHECK_RETURN(cudaMalloc((void**)&ipDevice_C, stBytesCount));

	CUDA_CHECK_RETURN(cudaMemcpy(ipDevice_A, ipHost_A, stBytesCount, cudaMemcpyHostToDevice));
	CUDA_CHECK_RETURN(cudaMemcpy(ipDevice_B, ipHost_B, stBytesCount, cudaMemcpyHostToDevice));

	dim3 block(4, 1);
	dim3 grid(4, 1);

	_addMatrix << <grid, block >> > (ipDevice_A, ipDevice_B, ipDevice_C, nx, ny);

	cudaThreadSynchronize();

	CUDA_CHECK_RETURN(cudaMemcpy(ipHost_C, ipDevice_C, stBytesCount, cudaMemcpyDeviceToHost));

	for (size_t i = 0; i < nxy; i++)
	{
		if (i % 4 == 0 && i)
			std::cout << std::endl;

		std::cout << ipHost_A[i] << " + " << ipHost_B[i] << " = " << ipHost_C[i] << "\t";

	}

	cudaFree(ipDevice_A);
	cudaFree(ipDevice_B);
	cudaFree(ipDevice_C);

	delete[]ipHost_A;
	delete[]ipHost_B;
	delete[]ipHost_C;

	ipHost_A = nullptr;
	ipHost_B = nullptr;
	ipHost_C = nullptr;


	return 0;
}

static void CheckCudaErrorAux(const char* file, unsigned line, const char* statement, cudaError_t err)
{
	if (err == cudaSuccess)
		return;
	std::cerr << statement << " returned: " << cudaGetErrorName(err) << "  \t : " << cudaGetErrorString(err) << "(" << err << ") at " << file << ":" << line << std::endl;
	exit(1);
}

cuda中使用二维矩阵
youzjuer的博客
09-19 367
经过上面的操作后，就可以像操作二维数组了t=O83At=O83ACUDA学习之一：二维矩阵加法 - 冷豆东 - 博客园 (cnblogs.com)https://www.cnblogs.com/jugg1024/p/4349243.htmlCUDA 中的 cudaMalloc使用二重指针（void**)的一些理解_cuda申请二重指针指针-CSDN博客https://blog.csdn.net/lingyunxianhe/article/details/92001270。
CUDA编程- - GPU线程的理解 thread,block,grid - 学习记录
点云算法工程师
01-27 3030
3、同一个线程束中的线程存储访问合并在一起了，长方形布局只需要一次访问操作就可以获得连续的内存数据 // 正方形要两次访问。在 blockIdx.x = 1 中， blockDim.x 表示所要求每个线程块启动的线程数量，在这 = 128。3、在核函数中编写单个线程的使用代码，再调用核函数，便可简单的达到 cpu 中 openmp 的多线程方式。在 cpu 中，用 openmp 并行计算，有限的线程数对 128 进行分组运算。2、线程数量的多少是由线程块，线程，线程网格等决定的，
CUDA（10）之深入理解threadIdx
热门推荐
林微的博客
06-21 1万+
摘要 本文主要讲述CUDA的threadIdx。   1. Grid，Block和Thread三者的关系 其中，一个grid包含多个blocks，这些blocks的组织方式可以是一维，二维或者三维。任何一个block包含有多个Threads，这些Threads的组织方式也可以是一维，二维或者三维。举例来讲：比如上图中，任何一个block中有10个Thread，那么，Block(0,0)...
CUDA-MODE 第四课: 计算和内存基础（基于PMPP 书的第4-5章）
I good vegetable a!
08-14 1056
我的课程笔记，欢迎关注：https://github.com/BBuf/how-to-optim-algorithm-in-cuda/tree/master/cuda-mode。
CUDA编程---全局内存
qq_40268672的博客
04-14 1478
既可以表明标量也可以表示数组。+只在计算能力2.x的设备上进行缓存。
深入理解CUDA内置变量—线程索引，内存位置，硬件特性
u012276729的博客
03-06 1434
作为C/C++编程中作为内置变量，不是直接对应硬件指令，而是CUDA编程模型中的抽象概念，用于标识线程在网格（grid）和块（block）中的位置，帮助开发者组织和标识线程。此外，线程索引通常不会直接保存在某种类型的“内存”上，而是作为CUDA内核函数的参数传递，并在函数内部使用。线程索引通常保存在GPU的内存上，具体来说，它们可能存储在全局内存、常量内存或纹理内存中，这取决于你如何使用它们。值，用于计算线程的全局索引，从而访问全局内存中的数据，或者用于执行特定的任务。
CUDA中线程与数据的对应关系
gigglesun的专栏
03-28 1562
引子 由于NVIDIA GPU采用的是SIMT的运行模式，CUDA编程中线程数量与数据的对应关系是什么呢？首先，我们来看一个经典的例子： #include <stdio.h> #define N (2048*2048) #define THREADS_PER_BLOCK 512 __global__ void add(int *a, int *b, int *c) { int index = threadIdx.x + blockIdx.x * blockDim.x; c[index]
【CUDA编程笔记】如何使用CUDA统一内存来优化多进程多线程程序的性能？
联系方式：927632640
08-26 1004
要使用CUDA统一内存优化多进程多线程程序的性能，可以采取以下步骤。
cuda-bitonic-sorting:Bitonic Sorting Cuda C++ 和一些额外的
06-14
CUDA编程主要通过定义kernel函数来实现，这些函数可以在GPU的多个线程块和线程上并行执行。Bitonic排序在CUDA中的实现通常涉及到以下关键步骤： 1. **初始化**: 在CUDA程序开始时，需要将数据从主机内存复制到设备...
cuda-使用cuda并行加速实现之matrixMul.zip
04-24
2. **CUDA线程和线程块**：CUDA将计算任务划分为线程和线程块。线程是执行计算的基本单元，而线程块是一组协同工作的线程。线程块内的线程可以共享数据，且在GPU硬件中可以一起调度，这有助于优化内存访问和计算效率...
CUDA-MODE课程笔记 第8课: CUDA性能检查清单
I good vegetable a!
08-19 1612
这节课相当于对Lecture 1更系统的补充，重点介绍GPU kernel优化的一些实用技术和分析工具ncu。合并全局内存访问： 确保线程访问连续的内存位置，以最大限度地提高内存带宽利用率。最大化占用率： 优化kernel启动参数，以充分利用 GPU 的处理能力。理解内存与计算限制： 分析kernel特征以确定限制因素（内存带宽或计算能力），并应用适当的优化技术。最小化线程的分化： 避免导致warp内的线程采用不同执行路径的条件语句，从而导致性能下降。
【CUDA编程】线程ID与全局ID的映射关系
weixin_46566156的博客
05-04 813
blockDim.方向 指的是该方向上thread的总数 girdDim.方向 指的是该方向上block的总数 1. 1D grid 1D block threadId = blockIdx.x * blockDim.x + threadIdx.x; 2. 1D grid 2D block threadIdx = blockIdx.x * blockDim.x * blockDim.y + // 确定在哪一个block threadIdx.y.
CUDA并行执行模型与多维数据映射
最新发布
weixin_35949153的博客
03-18 416
本文深入探讨了CUDA中数据并行执行模型的核心概念，特别是如何将线程映射到多维数据。通过分析二维网格处理图片的例子，我们了解到CUDA中的线程组织通常基于数据的性质，并且使用threadIdx、blockIdx、blockDim和gridDim变量来控制线程映射。此外，我们还学习了内存空间在现代计算机中的简化视图，以及线性代数函数在科学和工程应用中的重要性。
CUDA下在Host端分配的几种内存模式
bit262426738的专栏
04-21 6285
平时我们使用的内存都是Pageable(交换页)的，而另一个模式就是Pinned(Page-locked)，实质是强制让系统在物理内存中完成内存申请和释放的工作，不参与页交换，从而提高系统效率，需要使用cudaHostAlloc和cudaFreeHost来分配和释放。 　　优点 　　1、带宽更高 　　2、内核处理和内存拷贝可同时进行 　　3、可以内存映射(mapped)
CUDA学习笔记（十二）共享内存
云中寻雾的博客
07-21 1万+
CUDA SHARED MEMORY shared memory在之前的博文有些介绍，这部分会专门讲解其内容。在global Memory部分，数据对齐和连续是很重要的话题，当使用L1的时候，对齐问题可以忽略，但是非连续的获取内存依然会降低性能。依赖于算法本质，某些情况下，非连续访问是不可避免的。使用shared memory是另一种提高性能的方式。 GPU上的memory有两种： · On...
CUDA统一内存、零复制内存、锁页内存
葫芦与瓢的博客
12-15 1万+
CUDA 6.0以后，出现了：统一内存，unified memory，其具有以下特性： 1 其声明有以下两种方式： a 使用关键字managed，可选的和device,一起使用，比如：device managed int ret[1000]; b 使用函数cudaMallocManaged(): int *ret; cudaMallocManaged(&ret,1000*s
CUDA 编程--线程层次
caisense的专栏
04-26 3864
1.内核CUDA通过允许程序员定义称为内核的C函数扩展了C，内核调用时会被N个CUDA线程执行N次（译者注：这句话要好好理解，其实每个线程只执 行了一次），这和普通的C函数只执行一次不同。 内核使用 global 声明符定义，使用一种新<<< … >>>执行配置语法指定执行某一指定内核的线程数（参看下面代码）。每个执行内核的线程拥有一个独一无二的线程ID，可以通过内置的threadIdx变量在内
CUDA内存拷贝
weixin_30530339的博客
04-12 406
1、cudaMemcpy()<--> cudaMalloc() //线性内存拷贝 1 //线性内存拷贝 2 cudaMalloc((void**)&dev_A, data_size); 3 cudaMemcpy(dev_A, host_A, data_size, cudaMemcpyHostToDevice); 2、cudaMemcpy2D()&lt...
CUDA-By-Example: Julia 修复包详解与应用
CUDA提供多种内存类型，包括全局内存、共享内存、常量内存和纹理内存。全局内存是最容易访问的内存类型，但它相对较慢。共享内存用于同一块内线程之间的通信，比全局内存快得多。常量和纹理内存是只读内存，常用于...
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

抵扣说明：

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

余额充值