线程网格、线程块以及线程

第五章

写完本篇，做一个小总结，写在前面（限于笔者水平有限，仅提出自己的思考，以供参考）：
推荐理解的逻辑是：硬件 --> 软件 --> 编程 --> 解决实际问题

  • GPU从硬件上，有若干个SM，每个SM有若干个SP，SP是实际运行GPU线程的硬件，一个线程对应一个SP。由SM的数量和每个SM拥有SP的数量，可以得到该GPU理论上最大的并行线程数量。但由于SM内的寄存器大小等，可能会有影响。
  • 想要调用GPU上的线程，因此CUDA实现了一层抽象，相当于对于CUDA的编程人员而言，引出了线程束Warp、线程块block、线程网格grid等来组织GPU上的线程，方便编程人员。必须要了解硬件底层的设计，才能充分利用其性能。
  • 又因为GPU的硬件上有很多类型的存储器，比如：全局内存，纹理内存、常量内存、共享内存、寄存器等等，因此想要写一个高性能的CUDA程序，必须要利用好这些GPU硬件。

5.1 简介

  • 英伟达为它的硬件调度方式选择了一种比较有趣的模型，即SPMD（Single Program，Multiple Data），属于SIMD的一种变体。

根据弗林分类法，计算机的结构类型有：
● SIMD ---- 单指令，多数据
● MIMD ---- 多指令，多数据 ---- 相当于今天的双核或4核PC
● SISD ---- 单指令，单数据 ---- 相当于一个单核CPU在一个时刻只能执行一个任务
● MISD ---- 多指令，单数据

  • 并行编程的核心是线程的概念，一个线程就是程序中的一个单一的执行流，一个个线程组合在一起就形成了并行程序。
  • CUDA 的编程模型将线程组合在一起形成了线程束(warp)、线程块(block)以及线程网格(grid)。

5.2 线程

CPU与GPU的不同

GPU 和 CPU 设备的架构是迥异的：
● CPU 的设计是用来运行少量比较复杂的任务
● GPU的设计是用来运行大量比较简单的任务

CPU 与 GPU支持线程的方式不同。CPU的每个核只有少量的寄存器，每个寄存器都将在执行任何已分配的任务中被用到。为了能执行不同的任务，CPU将在任务与任务之间进行快速的上下文切换。
● 从时间的角度来看，CPU上下文切换的代价是非常昂贵的，因为每一次上下文切换都要将寄存器组里的数据保存到RAM中，等到重新执行这个任务时，又从RAM中恢复；
● 相比之下，GPU同样用到上下文切换这个概念，但它拥有多个寄存器组而不是单个寄存器组。因此，一次上下文切换调度者只需要设置一个寄存器组，用于将当前寄存器组里的内容换入、换出，它的速度比将数据保存到RAM中要快好几个数量级。

CPU 和 GPU 的一个主要差别就是每台设备上处理器数量的巨大差异；
GPU为每个SM提供了唯一并且高速的存储器，即共享内存，它为设备提供了在标准寄存器文件之外的本地工作区。自此，程序员可以安心地将数据留在内存中，不必担心由于上下文切换操作需要将数据移出去。另外，共享内存也为线程之间的通信提供了重要机制。

GPU线程

void some_func(void) {
    int i;
    for (i = 0; i < 128; i++) {
        a[i] = b[i] * c[i];
    } 
}

在CUDA中，可以创建一个内核函数的方式将上述的循环并行化。所谓的内核函数，就是一个只能在GPU上执行而不能直接在CPU上执行的函数。从概念上看，GPU的内核函数和循环体是一样，只不过将循环的结构移除了。

__global__ void some_kernel_func(int* const a, const int* const b, const int* const c) {
    a[i] = b[i] * c[i];
}

__global__ 前缀是告诉编译器在编译这个函数的时候生成的是GPU代码而不是CPU代码，并且这段GPU代码在CPU上是全局可见的。
调用内核函数时必须按照以下语法：kernel_function<<<num_blocks,num_threads>>>(param1,param2,...)
num_blocks: 表示线程块数
num_threads: 表示执行内核函数的线程数量
该内核函数总共会被调用 num_blocks*num_threads次
内核调用的下一部分就是参数的传递，我们可以通过寄存器或常量内存进行参数传递，具体是哪一种则视编译器而定。如果使用寄存器传参，每个线程用一个寄存器来传递一个参数。上面的例子中，现在有128个线程，每个线程传递3个参数，那么需要384个寄存器，听起来很多，但其实在每个SM中至少有8192个寄存器，而且随着后续硬件的发展，可能会更多。

__device__ 限定词声明一个函数是：1. 在设备上执行的; 2. 仅可从设备调用。
__global__ 限定词声明一个函数作为一个存在的kernel。这样的一个函数是：1. 在设备上执行的; 2. 仅可从主机调用。
__host__ 限定词声明的函数是：1. 在主机上执行的 2. 仅可从主机调用。

  • CPU 和 GPU有各自独立的内存空间，因此在GPU代码中，不可以直接访问CPU端的参数，反过来在CPU代码中，也不可以直接访问GPU端的参数。（稍后，我们将介绍一种特殊的方法来解决这个问题）因此，我们之前申明的全局数组a,b,c全是在CPU端的内存中，GPU端的代码是无法直接访问的，所以我们必须在GPU端的内存中也声明这几个数组，然后将数据从CPU端复制到GPU端，以GPU内存指针的方式传递给GPU的内存空间进行读写操作，在计算完毕之后，再将计算的结果复制回CPU端。
  • i 不再是循环控制变量，而是用来标识当前所运行的线程的一个变量。 CUDA提供了一个特殊的变量，它在每个线程中的值都不一样，使得它可以标识每一个线程。这就是线程的索引，即线程ID。我们可以直接将这个线程标号用作数组的下标对数组进行访问。线程的信息是由一个结构体存储的，这个例子中只用到了这个结构体的一个元素。
__global__ void some_kernel_func(int* const a, const int* const b, const int* const c) {
    const unsigned int thread_idx = threadIdx.x; 
    a[thread_idx] = b[thread_idx] * c[thread_idx];
}
/*
每个线程都进行了两次读内存操作，一次乘法操作，一次存储操作
注意到，每个线程执行的代码是一样的，但是数据却不相同，这就是CUDA的核心--SPMD模型
*/
  • 与CPU很相似，GPU的每个线程组被送到SM中，然后N个SP开始执行代码。在上面这个例子中，在得到每个线程的标号之后的第一件事就是从数组b和数组c各取一个数然后进行乘法操作，但这不是立即发生的，实际上，当从存储子系统取得所需要的数之后，已经过了400~600个GPU时钟周期。
  • 线程都是以32个为一组（这些线程组叫做线程束），当所有32个线程都在等待诸如内存读取这样的操作时，它们就会被挂起。

5.3 线程块

image
引入线程块后，则thread_idx 参数的计算变得更加复杂，而thread_idx通常又用来表示数组的下标位置。因此，我们之前简单的内核就要稍作调整(先只看上图左侧)

__global__ void some_kernel_func(int* const a, const int* const b, const int* const c) {
    const unsigned int thread_idx = (blockIdx.x * blockDim.x) + threadIdx.x; 
    a[thread_idx] = b[thread_idx] * c[thread_idx];
}
// blockIdx: The block index within the grid
// blockDim：The dimensions of the block
// threadIdx: The thread index within the block

还是对上面128个线程例子而言，我们启动两个线程块，每个线程块启动64个线程，则调用some_kernel_func<<<2,64>>>(a,b,c)

5.4 线程网格

一个线程网格是由若干线程块组成的。引入线程网格后，我们再观察线程的索引怎么计算：

const unsigned int idx = (blockIdx.x * blockDim.x) + threadIdx.x;
const unsigned int idy = (blockIdx.y * blockDim.y) + threadIdx.y;
const unsigned int thread_idx = ((gridDim.x * blockDim.x) * idy) + idx;
/*
gridDim.x : 线程网格X维度上线程块的数量
gridDim.y : 线程网格Y维度上线程块的数量
gridDim.z : 线程网格Z维度上线程块的数量

blockDim.x : 一个线程块X维度上的线程数量
blockDim.y : 一个线程块Y维度上的线程数量
blockDim.z : 一个线程块Z维度上的线程数量

threadIdx.x : 线程块X维度上的线程索引
threadIdx.y : 线程块Y维度上的线程索引
threadIdx.z : 线程块Z维度上的线程索引
*/

5.5 线程束

线程束是GPU的基本执行单位。每个线程束中的线程同时执行，在理想情况下，获得当前指令只需要一次访存，然后将指令广播到这个线程束所占用的所有SP中。当前，GPU上的一个线程束的大小为32，英伟达公司保留着对这个参数修改的权力，因此，他们提供了一个固有变量wrapSize，我们可以通过这个变量来获取当前硬件所支持的线程束的大小。

分支

我们之所以如此关注线程束的大小，一个很重要的原因就是分支。一个线程束是一个单独的执行单元，使用分支(例如，使用if、else、for、while、do、switch等语句)可以产生不同的执行流。
image

上图列举8个线程同步执行的情况，当该Warp存在if-else分支时，执行底层会有个变量exec数组，存储这每个线程当前的mask状态，如果分支为false不执行，那么就会标记该线程masked out，但仍旧需要跟着其他为true的线程同步“执行”,浪费时间。最差的情况就是32个线程中，只有一个为true，但是其他31个都需要同步等待，此时是利用率仅有1/32。同理，for循环的次数如果不一样也存在着相同的问题，当某些线程的循环次数较少，或者提前break，即使已经完成了也需要等待其他线程。

  • CPU有分支预测，它根据之前的运行情况来预测下一次执行到底要执行哪一块代码。在CPU上，指令流通常都会被预取，然后放入CPU的指令管线中。假设预测是准确的，那么CPU就避免了一次失速事件，如果预测错误，CPU则需要重新执行预测指令，然后获取另一个分支的指令，再将其添入管线之中。
  • 相比之下，GPU对分支的处理就没有这么复杂。GPU在执行完分支结构的一个分支后会接着执行另一个分支。对不满足分支条件的线程，GPU在执行这块代码的时候会将他们设置成未激活状态。当这块代码执行完毕之后，GPU继续执行另一个分支，这时，刚刚不满足分支条件的线程如果满足当前的分支条件，那么他们将被激活，然后执行这一段代码。最后所有的线程聚合，继续向下执行。
__global__ some_func(void) {
  if (thread_idx % 2) {
    action_a();
  } else {
    action_b();
  }
}

在上面的代码中，由于硬件每次只能为一个线程束获取一条指令，线程束中一半的线程要执行条件为真的代码段，一半线程执行条件为false的代码段，因此，这时会有一半的线程会被阻塞，而另一半线程会执行满足条件的那个分支。如此，硬件的利用率只达到了50%。

事实上，在指令执行层，硬件的调度是基于半个线程束，而不是整个线程束（笔者猜测，此处可能是指Fermi架构，因为Fermi架构的一个SM中有两个Warp schedule，公32个Core，故一个Warp只能运行16个线程，因此此处说硬件的调度是基于半个线程束的）。这意味着，只要我们能将半个线程束中连续的16个线程划分到同一分支指令中，那么硬件就能同时执行分支结构的两个不同条件的分支块，例如，使用下面的if-else的分支结构，这时，硬件的利用率就可以达到100%

__global__ some_func(void) {
  if ((thread_idx % 32) < 16) {
    action_a();
  } else {
    action_b();
  }
}

image
对于Fermi架构，有16个SM，每个SM中有2个Warp Schedule，共有32个Cuda Core，则在该设备中，SM每次可以有32个线程并行，并且这32个线程分成两个线程束，即该架构下，SM每次可以同时执行2个Warp，而每个Warp只能运行16个线程，后面的架构会变成32个，且可以同时执行更多的Warp。

grid、block以及warp关系

block是CUDA给我们划分的单位，warp是硬件层面中SM对应执行线程的单位。

kernel(指的是在GPU上执行的函数)在执行时会以一个grid为整体，划分若干个block，然后将block分配给SM进行运算。block中的线程以32个为一组，称为warp，进行分组计算。block会以连续的方式划分warp。例如，如果一个block由64个thread，则分为2组warp。0-31为warp0，32-63为warp1.如果block不是32的倍数，则多余的thread独立分成一组warp。例如block有65个thread，则最后一个thread单独为一个warp，那么此时这个warp中的其他thread处于非活动状态。

  • 在一个block内的warp次序是未定义的，但通过协调全局或者共享内存的存取，它们可以同步的执行。如果一个通过warp 线程执行的指令写入全局或共享内存的同一位置，写的次序是未定义的。
  • 在一个grid内的block次序是未定义的，并且在block之间不存在同步机制，因此来自同一个grid的二个不同block的线程不能通过全局内存彼此安全地通讯。
    image

image
当某个块中的warp在存取数据时，会切换到同一个块中的其他warp执行。所以我们在划分block时要确保有足够多的warp可共SM切换。
image

Grid，Block，thread都是线程的组织形式，最小的逻辑单位是一个thread，最小的硬件执行单位是thread warp，若干个thread组成一个block，block被加载到SM上运行，多个block组成一个Grid。
总而言之，一个kernel对应一个Grid，该Grid又包含若干个Block，Block内包含若干个thread。Grid跑GPU的时候，可能是独占一个GPU，也可能是多个kernel并发占用一个GPU。block是常驻在SM上的，一个SM可能有一个或者多个Block，具体根据资源占用分析。

5.6 线程块的调度

在线程块调度者为每个SM初始化分配了线程块之后，就会处于闲置状态，直到有线程块执行完毕。当线程块执行完毕之后就会从SM中撤出，并释放其占用的资源。由于线程块都是相同的大小，因此一个线程块从SM中撤出后另一个在等待队列中的线程块就会被调度执行。所有的线程块的执行顺序是随机、不确定的。因此，当我们在编写一个程序解决一个问题的时候，不要假定线程块的执行顺序，因为线程块根本就不会按照我们所想的顺序去执行。
保证在每个GPU中，线程块的数目都是SM数目的整数倍，以此提高设备的利用率。
其实从负载均衡的角度来看，这个问题还有待优化。因此，之后的CUDA运行时库中支持重叠的内核以及在同一块CUDA设备上可以运行多个单独的内核。通过这种方式，我们就可以维持吞吐量，使GPU集群不止有一个任务源可以调度。一旦设备出现闲置，他就会从内核流中选择另一个内核进行执行。

5.7 一个实例———统计直方图

在这个例子中，我们来统计数组array中数字0~255出现的次数，并将结果保存在bin数组中。array数组里面的每个值只占一个字节。
用串行算法解决非常简单，如下所示

for (unsigned int i = 0; i < max; i++) {
    bin[array[i]]++;
}
  • 当对这段串行代码转化成并行的时候，会有一个问题：如果我们用256个线程来执行，若多个线程同时对同一个bin进行++操作，将出现竞争的情况。这个问题并不少见，CUDA针对此也提出了一个较为简单的方法: atomicAdd(&value) 这个操作保证了对value这个值进行加法操作在所有线程之间是串行执行的。

既然问题得到了解决，那么我们就来考虑改用什么样的方式将之前串行代码转化为并行代码。主要有两种方式：

  • 一种是基于任务分解的模型
  • 一种是基于数据分解的模型

基于任务分解的模型

基于任务分解的模型主要是将输入数组中的元素分配到每一个线程中，然后在进行原子加法操作。这种方法对于程序来说是最简单的解决方案，但同时也有一定缺点，因为有些数据是共享资源。

如果直方图中所统计的信息是平均分布的，那么每个结果数组的每个位置(bin[i])产生的竞争的次数就等于输入数组中array数组中数字i出现的次数。 假设输入数组的大小为512MB（一共524 288个元素），一共出现了四种不同的数字，则每个bin就会有131 072个线程进行竞争。最糟糕的是，如果输入数组中所有的元素都对应同一个数字，由于通过原子操作对这个bin进行内存读写，因此我们的并行程序也就完全变成了一个串行程序。

下面是多线程采用任务分解模型的具体实现。以下是GPU程序代码：

/* Each thread writes to one block of 256 elements of global memory and contends for write access
                  每个线程写入一个包含 256 个元素的全局内存块并竞争写访问权 */
__global__ void myhistogram256Kernel_01(const unsigned char const* d_hist_data,unsigned int* const d_bin_data)
{
  /*Work out out thread id*/
  const unsigned int idx = (blockIdx.x * blockDim.x) + threadIdx.x;
  const unsigned int idy = (blockIdx.y * blockDim.y) + threadIdx.y;
  const unsigned int tid = idx + idy * gridDim.x * blockDim.x;

  /*Fetch the data value*/
  const unsigned char value = d_hist_data[tid];
  atomicAdd(&(d_bin_data[value]),1);
}

我们用一块GTX460显卡对这种方式进行测试，测得其处理速度为1025MB/s。有趣的是，该方法的处理速度并不随输入数组的大小改变而加快或降低。无论输入数组的大小是多少，都会得到一个固定的速度值，而且这个速度非常慢。对于显存为1GB的GTX460显卡而言，他的存储带宽为115GB/s，而我们测得的处理速度为1025MB/s，可见，这种方案的性能极低。
在该程序中，其中最可能影响性能的因素就是存储带宽。在程序中，一个有两次对内存的读/写操作，先是每次从输入数组中获取N个数，然后又压缩至N次写回一块1K的内存区(256个元素 X 每个数int占4字节)。

首先讨论一下读操作

  • 每个线程从输入数组中读得一个字节的元素，但由于线程束的读操作合并到了一起，即每半个线程束（16个线程）同时做一次读操作（不理解为什么是半个线程束，因为上文中的“事实上，在指令执行层，硬件的调度是基于半个线程束的，而不是整个线程束”？ 难道是因为当时Fermi架构下一个warp只能运行16个线程？后面的架构会变成32个。），由于最少的传输大小为32字节，但现在只读了16个字节，浪费了50%的存储带宽。
  • 在最好的情况下，每半个线程束最多能读128个字节，如果是这样，每个线程将从内存中获得4个字节，即可以通过处理4个数字，而不是先前的一个数字。
    一下是修改后的内核代码：
/* Each read is 4 bytes, not one, 32 X 4 = 128 bytes reads */
__global__ void myhistogram256Kernel_02(const unsigned int const* d_hist_data, unsigned int* const d_bin_data)
{
  /*Work out out thread id*/
  const unsigned int idx = (blockIdx.x * blockDim.x) + threadIdx.x;
  const unsigned int idy = (blockIdx.y * blockDim.y) + threadIdx.y;
  const unsigned int tid = idx + idy * gridDim.x * blockDim.x;

  /*Fetch the data value as 32 bit*/
  const unsigned int value_u32 = d_hist_data[tid];
  atomicAdd(&(d_bin_data[(value_u32 & 0x000000FF)]),1);
  atomicAdd(&(d_bin_data[(value_u32 & 0x0000FF00) >> 8]),1);
  atomicAdd(&(d_bin_data[(value_u32 & 0x00FF0000) >> 16]),1);
  atomicAdd(&(d_bin_data[(value_u32 & 0xFF000000) >> 24]),1);
}

但在运行内核时我们会发现，我们所做的努力只实现了零加速。事实上，这种现象在我们优化程序的时候会经常发生，原因主要是我们还没真正弄清楚引起性能瓶颈的因素。

  • 为什么没有优化呢？原因主要是内核是在计算能力为2.x的硬件上运行的。半个线程束合并读取内存数据的方式对计算能力为2.x的硬件来说并不会产生很大影响，因为它已经将整个线程束的内存读取合并到了一起。也就是说，在测试设备GTX460上（计算能力为2.1的硬件），来自一个线程束的32次单字节读取已经被合并成一次32字节的读取操作了。
  • 很明显，相对于存储带宽带来的微小影响，原子写操作才可能是性能瓶颈。为此，我们将采用另外一种方案来编写内核，即基于数据分解的模型

基于数据分解的模型

通过观察我们会发现，内核中有一些数据会被再次用到，而我们可以将这些被再次利用的数据放入能高效处理共享数据的存储区中，例如二级缓存和共享内存，以此来提高程序的性能。
image

  • 我们知道最初引起程序性能低下的原因就是对256个bin所产生的竞争。多个SM中多个线程块都要将它们的计算结果协会内存，然后硬件对每个处理器的缓存中的bin数组进行同步。分开来看，就是程序从内存中获取数据，然后做加法操作，最后将计算得到的新值写回内存。而这其中有好几步都可以一直在二级缓存中进行。在Fermi架构的硬件上，二级缓存中的数据在SM之间是共享的。相比之下，如果在计算能力1.x的硬件上，如果将数据都放在全局内存上进行读写操作，程序的性能会降低好几个数量级。
    image
    但即使在Fermi架构的硬件上使用了二级缓存，我们仍然需要对所有的SM进行同步。另一种方案就是让每一个SM都计算出一个统计直方图，最后再将所有的直方图汇总到一块主存上。无论是CPU还是GPU编程，我们都尽量实现这种方案，因为利用的资源越接近处理器（例如SM），程序运行就越快。
  • 之前我们提到使用共享内存。共享内存是一块比较特殊的内存，因为它存在于芯片上并且它的存取比全局内存更快。image
    我们可以在共享内存上创建一个包含256个bin的局部统计直方图，最后将所有共享内存上计算得到的统计直方图通过原子操作汇总到全局内存。假设每个线程块处理一个统计直方图（256字节），而对全局内存读写的操作次数也不会因此而减少，但写回内存的操作却因此可以合并起来。以下是这种方法的内核代码：
/* Each read is 4 bytes, not one, 32 X 4 = 128 bytes reads */
__shared__ unsigned int d_bin_data_shared[256];
__global__ void myhistogram256Kernel_03(const unsigned int const* d_hist_data, unsigned int* const d_bin_data)
{
  /*Work out out thread id*/
  const unsigned int idx = (blockIdx.x * blockDim.x) + threadIdx.x;
  const unsigned int idy = (blockIdx.y * blockDim.y) + threadIdx.y;
  const unsigned int tid = idx + idy * gridDim.x * blockDim.x;

  /* Clear shared memory */
  d_bin_data_shared[threadIdx.x] = 0; // 内核必须针对共享内存做一次额外的清楚操作，一面之前执行的内核随机残留一些数据

  /*Fetch the data value as 32 bit*/
  const unsigned int value_u32 = d_hist_data[tid];

  /* Wait for all threads to update shared memory */
  __syncthreads(); // 在将下一个线程块中所有线程的数据更新到共享内存中之前，必须等待前一个线程块的所有线程完成共享内存的清除操作
  // 意思应该是要等待一个SM中所有的线程块中的线程都将自己对应的区域清除之后，才能开始对共享内存进行原子写操作

  atomicAdd(&(d_bin_data_shared[(value_u32 & 0x000000FF)]),1);
  atomicAdd(&(d_bin_data_shared[(value_u32 & 0x0000FF00) >> 8]),1);
  atomicAdd(&(d_bin_data_shared[(value_u32 & 0x00FF0000) >> 16]),1);
  atomicAdd(&(d_bin_data_shared[(value_u32 & 0xFF000000) >> 24]),1);

  /* Wait for all threads to update shared memory */
  __syncthreads(); // 在将结果写回全局内存之气那，需要同步操作，以保证所有的线程都完成了计算

  /* The write the accumulated data back to global memory in blocks, not scattered */
  atomicAdd(&(d_bin_data[threadIdx.x]),d_bin_data_shared[threadIdx.x]);
}

此时，通过将连续的写操作合并，我们程序的性能提升了6倍，处理速度达到了6800MB/s。将连续的写操作合并起来之后，我们需要考虑如何减少全局内存的阻塞，我们已经对读数据的大小进行了优化，每次从源数据中读出一个值，而且每个值只需要读一次，因此，我们只需要考虑减少对全局内存写操作的次数。假设每个线程块处理的直方图不是一个而是N个，那么我们对全局内存的写操作的带宽就会减少N倍。

/* Each read is 4 bytes, not one, 32 X 4 = 128 bytes reads */
/* Accumulate into shared memory N times */
__shared__ unsigned int d_bin_data_shared[256];
__global__ void myhistogram256Kernel_04(const unsigned int const* d_hist_data, unsigned int* const d_bin_data, unsigned int N)
{
  /*Work out out thread id*/
  const unsigned int idx = (blockIdx.x * blockDim.x * N) + threadIdx.x;
  const unsigned int idy = (blockIdx.y * blockDim.y) + threadIdx.y;
  const unsigned int tid = idx + idy * gridDim.x * blockDim.x * N;

  /* Clear shared memory */
  d_bin_data_shared[threadIdx.x] = 0;

  /* Wait for all threads to update shared memory */
  __syncthreads();
  for (unsigned int i = 0, tid_offset = 0; i < N; i++ , tid_offset += 256) {
    const unsigned int value_u32 = d_hist_data[tid + tid_offset];
    atomicAdd(&(d_bin_data_shared[(value_u32 & 0x000000FF)]),1);
    atomicAdd(&(d_bin_data_shared[(value_u32 & 0x0000FF00) >> 8]),1);
    atomicAdd(&(d_bin_data_shared[(value_u32 & 0x00FF0000) >> 16]),1);
    atomicAdd(&(d_bin_data_shared[(value_u32 & 0xFF000000) >> 24]),1);
  }

  /* Wait for all threads to update shared memory */
  __syncthreads();

  /* The write the accumulated data back to global memory in blocks, not scattered */
  atomicAdd(&(d_bin_data[threadIdx.x]),d_bin_data_shared[threadIdx.x]);
}
  • 用一个循环变量i进行N次循环，每次循环都处理了在共享内存中的256个字节的直方图数据。每个线程块包含256个线程，每个线程计算一个bin。而循环的次数即每个线程块处理的直方图的个数。
  • 每执行完一次循环，指向内存的指针向后移动256个字节，以指向下一个处理的直方图。

参考：
https://blog.csdn.net/leviopku/article/details/120157690
《CUDA 并行程序设计：GPU编程指南》
https://blog.csdn.net/u011475134/article/details/71577770
https://zhuanlan.zhihu.com/p/349185459
https://zhuanlan.zhihu.com/p/455866677
https://www.cnblogs.com/QZ-CMD/articles/15983695.html
https://zhuanlan.zhihu.com/p/626636620
posted @ 2023-05-08 17:55  人生逆旅，我亦行人  阅读(782)  评论(1)  编辑  收藏  举报