0

以下主机代码test.c和设备代码test0.cu旨在提供相同的结果。

test.c

$ cat test.c
#include <stdio.h>
#include <string.h>

int main()
{
        int data[32];
        int dummy[32];

        for (int i = 0; i < 32; i++)
                data[i] = i;

        memcpy(dummy, data, sizeof(data));
        for (int i = 1; i < 32; i++)
                data[i] += dummy[i - 1];
        memcpy(dummy, data, sizeof(data));
        for (int i = 2; i < 32; i++)
                data[i] += dummy[i - 2];
        memcpy(dummy, data, sizeof(data));
        for (int i = 4; i < 32; i++)
                data[i] += dummy[i - 4];
        memcpy(dummy, data, sizeof(data));
        for (int i = 8; i < 32; i++)
                data[i] += dummy[i - 8];
        memcpy(dummy, data, sizeof(data));
        for (int i = 16; i < 32; i++)
                data[i] += dummy[i - 16];

        printf("kernel  : ");
        for (int i = 0; i < 32; i++)
                printf("%4i ", data[i]);
        printf("\n");
}
$

test0.cu

$ cat test0.cu
#include <stdio.h>

__global__ void kernel0(int *data)
{
        size_t t_id = threadIdx.x;

        if (1 <= t_id)
                data[t_id] += data[t_id - 1];
        if (2 <= t_id)
                data[t_id] += data[t_id - 2];
        if (4 <= t_id)
                data[t_id] += data[t_id - 4];
        if (8 <= t_id)
                data[t_id] += data[t_id - 8];
        if (16 <= t_id)
                data[t_id] += data[t_id - 16];
}

int main()
{
        int data[32];
        int result[32];

        int *data_d;
        cudaMalloc(&data_d, sizeof(data));

        for (int i = 0; i < 32; i++)
                data[i] = i;

        dim3 gridDim(1);
        dim3 blockDim(32);

        cudaMemcpy(data_d, data, sizeof(data), cudaMemcpyHostToDevice);
        kernel0<<<gridDim, blockDim>>>(data_d);
        cudaMemcpy(result, data_d, sizeof(data), cudaMemcpyDeviceToHost);

        printf("kernel0 : ");
        for (int i = 0; i < 32; i++)
                printf("%4i ", result[i]);
        printf("\n");
}
$

如果我编译并运行它们,它们会给出与我预期相同的结果。

$ gcc -o test test.c
$ ./test
kernel  :    0    1    3    6   10   15   21   28   36   45   55   66   78   91  105  120  136  153  171  190  210  231  253  276  300  325  351  378  406  435  465  496
$ nvcc -o test_dev0 test0.cu
$ ./test_dev0
kernel0 :    0    1    3    6   10   15   21   28   36   45   55   66   78   91  105  120  136  153  171  190  210  231  253  276  300  325  351  378  406  435  465  496
$

但是,如果我在设备代码中使用共享内存而不是全局内存,如 中test1.cu所示,它会给出不同的结果。

test1.cu

$ cat test1.cu
#include <stdio.h>

__global__ void kernel1(int *data)
{
        __shared__ int data_s[32];

        size_t t_id = threadIdx.x;

        data_s[t_id] = data[t_id];

        if (1 <= t_id)
                data_s[t_id] += data_s[t_id - 1];
        if (2 <= t_id)
                data_s[t_id] += data_s[t_id - 2];
        if (4 <= t_id)
                data_s[t_id] += data_s[t_id - 4];
        if (8 <= t_id)
                data_s[t_id] += data_s[t_id - 8];
        if (16 <= t_id)
                data_s[t_id] += data_s[t_id - 16];

        data[t_id] = data_s[t_id];
}

int main()
{
        int data[32];
        int result[32];

        int *data_d;
        cudaMalloc(&data_d, sizeof(data));

        for (int i = 0; i < 32; i++)
                data[i] = i;

        dim3 gridDim(1);
        dim3 blockDim(32);

        cudaMemcpy(data_d, data, sizeof(data), cudaMemcpyHostToDevice);
        kernel1<<<gridDim, blockDim>>>(data_d);
        cudaMemcpy(result, data_d, sizeof(data), cudaMemcpyDeviceToHost);

        printf("kernel1 : ");
        for (int i = 0; i < 32; i++)
                printf("%4i ", result[i]);
        printf("\n");
}
$

如果我编译并运行它,它会给出与ortest1.cu不同的结果。test0.cutest.c

$ nvcc -o test_dev1 test1.cu
$ ./test_dev1
kernel1 :    0    1    2    3    4    5    6    7    8    9   10   11   12   13   14   15   16   17   18   19   20   21   22   23   24   25   26   27   28   29   30   31
$

经纱同步不应该与共享内存一起使用吗?


对此问题的一些调查:

使用 CUDA8.0 时,如果我test1.cu使用-arch=sm_61选项进行编译(我正在使用 GTX 1080 进行测试),它会给出与test0.cuand相同的结果test.c

$ nvcc -o test_dev1_arch -arch=sm_61 test1.cu
$ ./test_dev1_arch
kernel1 :    0    1    3    6   10   15   21   28   36   45   55   66   78   91  105  120  136  153  171  190  210  231  253  276  300  325  351  378  406  435  465  496
$

但这不适用于较新版本的 CUDA。如果我使用任何比 8.0 更新的版本,即使我给出-arch=sm_61选项,测试结果也会有所不同。

4

2 回答 2

1

由于两种情况下的竞争条件,使用共享内存或使用全局内存,您的设备代码具有未定义的行为。您有多个线程同时读取和修改同一个int对象。

经纱同步不应该与共享内存一起使用吗?

我在您的代码中看不到任何扭曲同步。

硬件在锁步中执行扭曲的事实(一开始不一定是真的)是完全不相关的,因为读取您的 C++ 代码的不是硬件。它是您用来将 C++ 代码转换为将在您的硬件上实际运行的机器代码的任何工具链。并且允许 C++ 编译器基于 C++ 语言的抽象规则进行优化。

让我们看看为您的示例实际生成的机器代码(在我的机器上使用 CUDA 10):

_Z7kernel1Pi:
        /*0008*/                   MOV R1, c[0x0][0x20] ;
        /*0010*/                   S2R R9, SR_TID.X ;
        /*0018*/                   SHL R8, R9.reuse, 0x2 ;
        /*0028*/                   SHR.U32 R0, R9, 0x1e ;
        /*0030*/                   IADD R2.CC, R8, c[0x0][0x140] ;
        /*0038*/                   IADD.X R3, R0, c[0x0][0x144] ;
        /*0048*/                   LDG.E R0, [R2] ;
        /*0050*/                   ISETP.NE.AND P0, PT, R9.reuse, RZ, PT ;
        /*0058*/                   ISETP.GE.U32.AND P1, PT, R9, 0x2, PT ;
        /*0068*/               @P0 LDS.U.32 R5, [R8+-0x4] ;
        /*0070*/         {         ISETP.GE.U32.AND P2, PT, R9.reuse, 0x4, PT ;
        /*0078*/               @P1 LDS.U.32 R6, [R8+-0x8]         }
        /*0088*/                   ISETP.GE.U32.AND P3, PT, R9, 0x8, PT ;
        /*0090*/               @P2 LDS.U.32 R7, [R8+-0x10] ;
        /*0098*/         {         ISETP.GE.U32.AND P4, PT, R9, 0x10, PT   SLOT 0;
        /*00a8*/               @P3 LDS.U.32 R9, [R8+-0x20]   SLOT 1        }
        /*00b0*/               @P4 LDS.U.32 R10, [R8+-0x40] ;
        /*00b8*/         {         MOV R4, R0 ;
        /*00c8*/                   STS [R8], R0         }
        /*00d0*/               @P0 IADD R5, R4, R5 ;
        /*00d8*/         {     @P0 MOV R4, R5 ;
        /*00e8*/               @P0 STS [R8], R5         }
        /*00f0*/               @P1 IADD R6, R4, R6 ;
        /*00f8*/         {     @P1 MOV R4, R6 ;
        /*0108*/               @P1 STS [R8], R6         }
        /*0110*/               @P2 IADD R7, R4, R7 ;
        /*0118*/         {     @P2 MOV R4, R7 ;
        /*0128*/               @P2 STS [R8], R7         }
        /*0130*/               @P3 IADD R9, R4, R9 ;
        /*0138*/         {     @P3 MOV R4, R9 ;
        /*0148*/               @P3 STS [R8], R9         }
        /*0150*/               @P4 IADD R10, R4, R10 ;
        /*0158*/               @P4 STS [R8], R10 ;
        /*0168*/               @P4 MOV R4, R10 ;
        /*0170*/                   STG.E [R2], R4 ;
        /*0178*/                   EXIT ;
.L_1:
        /*0188*/                   BRA `(.L_1) ;
.L_14:

如您所见,编译器(在这种特殊情况下,“罪魁祸首”实际上是 PTX 汇编程序)已将您的 if 序列转换为一组指令,这些指令根据 if 条件设置谓词。它首先使用条件加载将它需要的所有值从共享内存中提取到寄存器中。只有在那之后,它才会使用已经加载的值执行所有的加法和条件存储。这是对您的 C++ 代码的完全合法解释。由于您没有指定任何同步或内存排序约束,编译器可以在没有潜在并发冲突的假设下运行,并且所有这些加载和存储都可以以它认为合适的任何方式重新排序。

要修复您的代码,请使用显式扭曲同步

__global__ void kernel1(int *data)
{
        __shared__ int data_s[32];

        size_t t_id = threadIdx.x;

        data_s[t_id] = data[t_id];

        __syncwarp();
        if (1 <= t_id)
                data_s[t_id] += data_s[t_id - 1];
        __syncwarp();
        if (2 <= t_id)
                data_s[t_id] += data_s[t_id - 2];
        __syncwarp();
        if (4 <= t_id)
                data_s[t_id] += data_s[t_id - 4];
        __syncwarp();
        if (8 <= t_id)
                data_s[t_id] += data_s[t_id - 8];
        __syncwarp();
        if (16 <= t_id)
                data_s[t_id] += data_s[t_id - 16];

        data[t_id] = data_s[t_id];
}

这个问题只从 CUDA 9.0 开始出现的原因是,当 Volta 和“独立线程调度”使其成为必要时,warp 级同步才真正在 CUDA 9.0 中引入。在 CUDA 9.0 之前,官方不支持 warp-synchronous 编程。但是编译器在实际破坏代码(如上面的示例中)时非常保守。原因可能是这种“扭曲同步”编程(注意引号)通常是接近峰值性能的唯一方法,没有真正的替代方案,因此人们一直在这样做。不过,这仍然是未定义的行为,NVIDIA 一直在警告我们。它只是碰巧在很多情况下都有效……</p>

于 2019-01-08T12:44:00.783 回答
0

看来我错过的一点是用volatile限定符声明共享内存。这解决了这个问题。(测试代码

但是,正如Michael Kenzel 的回答中所述,通常应避免这种隐式扭曲同步编程,即使这是在NVIDIA 本身提供的经典并行缩减(第 22 页)中引入的。

由于未来的编译器和内存硬件可能会以不同的方式工作,因此依赖它会很危险。使用__syncwarp()类似于Michael Kenzel 提供的解决方案应该是更好的解决方案。在这篇 NVIDIA 开发者博客文章的帮助下,安全的解决方案是:

__global__ void kernel(int *data)
{
    __shared__ int data_s[32];

    size_t t_id = threadIdx.x;

    data_s[t_id] = data[t_id];

    int v = data_s[t_id];

    unsigned mask = 0xffffffff;     __syncwarp(mask);

    mask = __ballot_sync(0xffffffff, 1 <= t_id);
    if (1 <= t_id) {
        v += data_s[t_id - 1];  __syncwarp(mask);
        data_s[t_id] = v;       __syncwarp(mask);
    }
    mask = __ballot_sync(0xffffffff, 2 <= t_id);
    if (2 <= t_id) {
        v += data_s[t_id - 2];  __syncwarp(mask);
        data_s[t_id] = v;       __syncwarp(mask);
    }
    mask = __ballot_sync(0xffffffff, 4 <= t_id);
    if (4 <= t_id) {
        v += data_s[t_id - 4];  __syncwarp(mask);
        data_s[t_id] = v;       __syncwarp(mask);
    }
    mask = __ballot_sync(0xffffffff, 8 <= t_id);
    if (8 <= t_id) {
        v += data_s[t_id - 8];  __syncwarp(mask);
        data_s[t_id] = v;       __syncwarp(mask);
    }
    mask = __ballot_sync(0xffffffff, 16 <= t_id);
    if (16 <= t_id) {
        v += data_s[t_id - 16]; __syncwarp(mask);
        data_s[t_id] = v;
    }

    data[t_id] = data_s[t_id];
}
于 2019-01-11T02:17:04.503 回答