CUDA 程序的控制流程

Question

下面是我编写的一个小程序，用于查看在 CUDA 中竞争条件如何发生，但我对输出感到惊讶。

#include<cutil.h>
#include<iostream>
__global__ void testLocal(int *something, int val[]){

 *something = *something/2;


 val[threadIdx.x] = *something;
}

void main(){

    int *a, *c;
    int r =16;

    cudaMalloc((void**)&a, 4*sizeof(int));
    cudaMalloc((void**)&c, sizeof(int));
    cudaMemcpy(c, &r, sizeof(int) , cudaMemcpyHostToDevice);
    testLocal<<<1,4>>>(c,a);
    int *b = (int *)malloc(4 * sizeof(int));
    cudaMemcpy(b,a, 4 * sizeof(int), cudaMemcpyDeviceToHost);

    for( int  j =0 ; j< 4; j++){
        printf("%d\n",b[j]);

    }
    getchar();


}

当我启动 4 个线程时，我希望每个线程将 *something 除以 2 一次。我知道他们划分 *something 的顺序是不固定的。因此，当我尝试打印这些值时，我希望打印的值之一是 8，一个是 4，一个是 2，一个是 1。但是，所有打印的值都是 8。为什么会这样?不应该所有线程都划分*某事一次。

Answer 1

您正在查看的是未定义的行为。因为您要启动一个具有 4 个线程的块，所以所有线程都在同一个 warp 中执行。这意味着

 *something = *something/2;

正在由您启动的所有线程同时执行。 CUDA 编程模型只保证当来自同一个 warp 的多个线程尝试写入同一个内存位置时，其中一个写入会成功。它没有说明哪个线程会成功，以及经线中没有“获胜”的其他线程会发生什么。要获得您期望的行为，需要序列化的内存访问——这只能通过在支持它们的架构上使用原子内存访问原语来实现。

Answer 2

应该是一个强词。你正在做的事情是未指定的，所以它不应该做任何具体的事情。

现在，它可能所做的是在同一个 warp 内的同一个计算单元上运行 4 个线程。 （“SIMT”模型使每个线程作为扭曲的一部分运行）。由于您对something 的操作不是原子操作，因此warp 中的所有线程都以锁步方式读取和写入内存。于是4个线程一起读取*something，然后都将结果除以2，都尝试将8写入内存。

正如您所期望的那样，*something 被原子地读取和写入是通过原子操作实现的，尽管 CUDA 中没有可用的原子除法或乘法运算。所以如果你真的想要这个，你需要自己写（在 atomicCAS 的帮助下）。而且您会开始看到性能急剧下降，因为您现在正在强制努力并行运行的线程以串行方式运行。