如何计算哪个SM给定线程正在运行？

Question

我是CUDA初学者.

到目前为止,我了解到,每个SM都有8个块(线程).假设我有简单的工作将数组中的元素乘以2.但是,我的数据少于线程.

不是问题,因为我可以切断线程的"尾部"以使它们闲置.但是,如果我理解正确,这将意味着一些SM将获得100%的工作,而某些部分(甚至没有).

因此,我想计算哪个SM正在运行给定线程并以这种方式进行计算,每个SM具有相同的工作量.

我希望它首先有意义:-)如果是这样,如何计算给定线程运行哪个SM？或者 - 当前SM的索引和它们的总数？换句话说,在SM术语中等效于threadDim/threadIdx.

更新

评论太久了.

罗伯特,谢谢你的回答.当我试图消化所有,这里是我做的-我有一个"大"数组,我只是有乘以值*2和它(存储阵列输出作为热身,顺便说一句我做所有的计算,数学上是正确的. ).所以首先我在1个块,1个线程中运行它.精细.接下来,我尝试以这样的方式拆分工作,即每个乘法只由一个线程完成一次.结果我的程序运行速度慢了大约6倍.我甚至感觉到为什么-小罚获取有关GPU的信息,然后计算有多少块和线程我应该使用,则每个线程,而不是单一的乘法中,现在我身边有10多只的乘法计算阵列中的偏移量一个线程.一方面,我试图找出如何改变这种不受欢迎的行为,另一方面,我想在SM中均匀地传播线程的"尾部".

我改写 - 也许我错了,但我想解决这个问题.我有1G小工作(*2就是全部) - 我应该用1K线程创建1K块,或者用1个线程创建1M块,用1M线程创建1块,依此类推.到目前为止,我读取了GPU属性,除法,除法,并盲目地使用网格/块的每个维度的最大值(如果没有要计算的数据,则使用所需的值).

代码

size是输入和输出数组的大小.一般来说:

output_array[i] = input_array[i]*2;

计算我需要多少块/线程.

size_t total_threads = props.maxThreadsPerMultiProcessor
                       * props.multiProcessorCount;
if (size<total_threads)
    total_threads = size;

size_t total_blocks = 1+(total_threads-1)/props.maxThreadsPerBlock;

size_t threads_per_block = 1+(total_threads-1)/total_blocks;  

拥有props.maxGridSize和props.maxThreadsDim我以类似的方式计算块和线程的维度 - 来自total_blocks和threads_per_block.

然后是杀手部分,计算线程的偏移量(线程内部):

size_t offset = threadIdx.z;
size_t dim = blockDim.x;
offset += threadIdx.y*dim;
dim *= blockDim.y;
offset += threadIdx.z*dim;
dim *= blockDim.z;
offset += blockIdx.x*dim;
dim *= gridDim.x;
offset += blockIdx.y*dim;
dim *= gridDim.y;

size_t chunk = 1+(size-1)/dim;

所以现在我有当前线程的起始偏移量,以及用于乘法的数组(块)中的数据量.我没有使用grimDim.z上面,因为AFAIK总是1,对吧？

Answer 1

尝试这样做是不寻常的事情.鉴于你是一名CUDA初学者,在我看来这样的问题表明试图不正当地解决问题.你想解决的问题是什么？如果您在SM X与SM Y上执行特定线程,它如何帮助您解决问题？如果您希望从机器中获得最大性能,请以一种所有线程处理器和SM都可以处于活动状态的方式构建您的工作,事实上,所有人都有"足够多的工作".GPU依赖于超额订阅资源来隐藏延迟.

作为CUDA初学者,您的目标应该是:

• 在块和线程中创建足够的工作
• 有效地访问内存(这主要与合并有关 - 你可以阅读)

确保"每个SM都有相同的工作量"没有任何好处.如果您在网格中创建足够的块,每个SM 都会有工作的大致等量.这是调度程序的工作,您应该让调度程序执行此操作.如果你没有创建足够的块,你的第一个目标应该是创建或找到更多的工作要做,而不是为每个块提出花哨的工作细分而不会产生任何好处.

费米GPU中的每个SM(例如)具有32个线程处理器.为了使这些处理器即使在由于存储器访问等而存在不可避免的机器停顿的情况下也保持忙碌,该机器被设计成当发生停顿时通过交换另一个线程的扭曲(32)来隐藏等待时间,从而处理可以继续.为了实现这一点,您应该尝试每个SM拥有大量可用的warp.通过以下方式促进了这一点:

• 网格中的许多线程块(至少是GPU中SM数量的6倍)
• 每个threadblock有多个warp(可能至少有4到8个warp,所以每个块有128到256个线程)

由于(Fermi)SM总是一次执行32个线程,如果我在任何时刻拥有的线程少于我的GPU中SM数量的32倍,那么我的机器利用率就会低下.如果我的整个问题仅由20个线程组成,那么它根本就没有很好地利用任何GPU,并且将这20个线程分成多个SM /线程块也不太可能有任何明显的好处.

编辑:由于您不想发布您的代码,我会再提出一些建议或意见.

1. 你试图修改一些代码,发现它运行得慢,然后跳到(我认为)错误的结论.
2. 您应该熟悉一个像vector vector这样的简单代码示例.它不是每个元素的倍增,但结构是接近的.使用单个线程执行此向量添加无法实际运行得更快.我想如果你研究这个例子,你会找到一种直接的方法来扩展它来做数组元素乘以2.
3. 没有人按照你概述的方式计算每个块的线程数.首先,每个块的线程数应该是32的倍数.其次,习惯上选择每个块的线程作为起点,并从中构建其他启动参数,而不是相反.对于一个大问题,只需从每个块的256或512个线程开始,并省略其计算.
4. 根据您选择的线程块大小构建其他启动参数(网格大小).您的问题本质上是1D,因此1D网格块的1D网格是一个很好的起点.如果此计算超出了x维中最大块的机器限制,那么您可以让每个线程循环处理多个元素,或者扩展到2D网格(1D线程块).
5. 您的偏移计算不必要地复杂.请参阅向量添加示例,了解如何使用相对简单的偏移计算来创建线程网格以处理数组.