关于线程，线程块，线程格概念的理解请教大神

加速计算 CUDA
1

线程是GPU的基本运行单位假设有以下代码

global void addcuda(int * a, int * b, int * c)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
while (tid < N)
{
c[tid] = a[tid] + b[tid];
tid += blockDim.x * gridDim.x;
}
return;
}

GPU计算时，同时并行有N个线程，threadID分别是0~N-1

然后线程块是对多个并行线程的划分，一个线程块中有多少线程取决于kernel函数启动时的执行配置和硬件上的限制

一个线程块内的所有线程都是在同一个流处理器内执行的

上面代码中红色部分就是我所不理解的线程格的概念

顺便问个问题，关于归约算法，用以下代码做例子

int i = blockDim.x/2;

while (i!=0 && cacheIndex<i)
{
	cache[cacheIndex] += cache[cacheIndex+i];
	__syncthreads();
	i /= 2;
}

当这段代码并行计算结束后，归约的结果不是已经存放在cache[cacheIndex]中了吗？

新手有很多不明白，请大神指点！

2

这些都是基本问题，建议好好看看书，你这根本就没认为看完手册的第一章。

关于你的2个问题：
（1）每次tid += blockDim.x * gridDim.x; 是在干嘛：
这个很好理解，只需要你明白blockDim.x * gridDim.x是什么。
你应该看出（此信息你没给出），这是个1维的启动配置，blockDim.x * gridDim.x是当前grid中所在的所有线程的总数。而每个线程的线程编号每次循环增加上“这个总数”，是为了能在线程数目不够问题的元素规模（N）的时候，能够处理完毕所有的元素。

这个问题的关键在于1维度的启动线程形状，以及，启动的线程数目不够处理问题的规模上。
但可惜您的帖子正好遗漏了这2个要点，从而导致您无法理解。

（2）关于第二点的规约中的cacheIndex是什么：
这里您要知道，此代码同时被多个线程执行的，每个线程都有自己的cacheIndex，是不同的。
所以对于您的block中的多个线程，cacheIndex的值可能从A到B(A和B是2个您自己代码的计算cacheIndex的值）每个线程的值都不同，范围是从A到B.

而最终cacheIndex在哪里，一般是在您[A]里。您需要这样理解。而不是理解成[cacheIndex]里，否则这些cacheIndex有多种指定，您这等于没说。

此问题的重点在于cacheIndex的计算，但同样很可惜，您的帖子依然遗忘了此最重要的要点。

以及，楼主如果不是故意来捣乱的话（我相信您不是），
我强烈建议您仔细阅读下手册，那怕英文的您看不懂也没事，论坛有中文版（在资源板块）可以下载，您不妨看看。学会的知识是自己的，稍微花一点时间也是有好处的。
您觉得呢？

感谢来访。

3

非常感谢大神的回复！
第一个问题是我没问好，我想问的是线程块映射到一个流处理器，那么线程格映射的物理结构是什么？因为我之前理解为线程格映射的物理结构为整个GPU
第二个问题，关于归约，我想问的是那段代码执行完以后归约是否已经完成了？cacheindex我知道是什么
因为我手头缺少cuda的文献，只靠两本官方的书籍。。估计也不是给初学者看的。。所以很多基础概念没理解清楚。。谢谢大神告诉我论坛里还有这样的资料，之前一直不知道，btw我不是捣乱的哈哈

4

楼主您好，

(1)block和SP之前并无任何固定映射关系。SP可以为多个不同的blocks来执行代码。

(2)您可以认为grid映射在GPU上，但具体点说，一般情况下GPU上同时只有grid中的部分线程的映射，当这部分线程执行完毕后，grid中的下一批线程继续被放到GPU上来执行。（因为我们的线程数目一般要大于GPU能同时执行的线程数目）。

(3)那要看你的问题，如果您只是需要将blockDim.x个数据进行求和，自然是完成的。反之，如果每个block只是在对一大批数据中的部分在处理，那么你还需要将这些和进一步求和才能完成最终的和的计算。您觉得呢？

(4)官方的2本书已经足够，不妨看下。

感谢来访，欢迎再来。

5

线程和SP里面的CUDA核心有映射关系吗？

6

楼主您好，

无固定的映射关系的，一个SP可以执行多个不同线程里的代码的（当然不是同一时刻）。

感谢来访。

7

谢谢回答，我的意思是，线程块内的所有线程都在同一个SP上执行
然后，SP里面的CUDA核心是不是就是一个CUDA核心对应一个线程的执行？

8

LZ您好：

您这里的概念有问题，SP（stream processor）就是CUDA core的，两者是同一个东西。

所以您的描述就无法理解了。

9

…

您的7#我拒绝回答，建议您看点书，稍微对CUDA理解点基本概念再来提问。
否则这是对您的时间的浪费。

感谢来访。

10

谢谢回复，但是官方书《高性能CUDA应用设计与开发》里P73到P75的描述和您的描述貌似有点不相符

11

LZ您好：

1：您提到的这本书不是NVIDIA的官方资料，以及您也没有提供该资料的具体表述。因此我不评价您10#的说法。

2：我在8#指出的“SP（stream processor）就是CUDA core”观点，接受任何人对任何NVIDIA官方资料的查证。

12

谢谢回复，但是官方书《高性能CUDA应用设计与开发》里P73到P75的描述让我产生疑惑[attach]3390[/attach]这张图里说明的是一个block对内的thread在一个SP里运行，和我目前的理解一样
而书本《高性能CUDA应用设计与开发》P73到P75页的描述也和这个一致，但是
在P75的图中引进了CUDA核心这个没有描述过的概念

感谢您的回答，一般能解决的问题我都会自己解决的，但是有一些文献给出了不同的描述
让我感到比较困惑，所以上论坛询问，而这些问题正正是一些概念上的理解问题，所以请您谅解

13

那请问官方资料是不是以NV官网上的为准？因为CUDA的文献实在少，而且大家的说法貌似都有点不一样，所以想搞清楚概念和基础以方便后来的编程，请谅解

14

您好，另一位版主告诉我SP单元即为CUDA核心，非常感谢您耐心回答！

15

LZ您好：

本版一般不对第三方资料进行回应，不过我可以就您给出的这个图简单说一下：

1：此图没有任何表明“一个block内的threads在一个SP中执行”的含义。该图只是示意性地用箭头表明可以有两个blocks同时resident在一个SM（stream multi-processor）上。

2：“书本《高性能CUDA应用设计与开发》P73到P75页的描述也和这个一致”，此书不是NV官方资料，而且您没有提供该书的原始描述，而只是提出了您自己的理解。在此不对该观点评论。

3：“P75的图中引进了CUDA核心这个没有描述过的概念”，请您联系该书的作者或者译者以获得“CUDA核心”在书中的具体含义。
按照NVIDIA资料中的习惯CUDA core等同于SP，SP是较早期的叫法，CUDA core是现在的叫法。

16

LZ您好：

请以NVIDIA提供的官方文档为准（虽然也有错误）。
该文档可以在 CUDA Toolkit Documentation 12.3 Update 2 这里获取在线版本的最新版。
以及，随CUDA Toolkit免费发放。

最常用的是CUDA C Programming Guide和CUDA Reference Manual（CUDA API Reference Manual）这两本。请以英文原版为准，中文译版为辅。

祝您好运~