代码之家 › 专栏 › 技术社区 › nglee

CUDA共享内存和扭曲同步

gpu-warp gpu-shared-memory cuda

0

nglee · 技术社区 · 7 年前

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

$ 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。铜

$ 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 ,它给出了不同的结果。

测试1。铜

$ 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");
}
$

如果我编译 测试1。铜 然后运行它,它给出的结果与 test0。铜 测验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
$

扭曲同步不应该与共享内存一起工作吗?

测试1。铜 -arch=sm_61 选项(我正在使用GTX1080进行测试),它给出的结果与 .

$ 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更新的版本,即使我给出了选项

2 回复 | 直到 7 年前

1

Michael Kenzel 7 年前

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

扭曲同步不应该与共享内存一起工作吗?

硬件在锁定步骤中执行翘曲(这并不一定是真正的开始)是完全不相关的,因为它不是硬件,谁读取你的C++代码。无论您使用什么工具链,都将C++代码转换成实际运行在硬件上的机器代码。并根据C++语言的抽象规则,对C++编译器进行优化。

_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汇编程序)已将ifs序列转换为一组指令,这些指令根据if条件设置谓词。信息技术吸引

要修复代码,请使用 explicit warp synchronization

__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和“独立线程调度”使之成为必要时,CUDA 9.0才真正引入了扭曲级同步。在CUDA9.0之前,官方不支持warp同步编程。但编译器过去在实际破坏代码时相当保守,如上面的示例所示。原因可能是这样的“扭曲同步”编程(注意引号)通常是接近峰值性能的唯一方法,没有真正的替代方案,因此,人们一直在这样做。不过,这仍然是未定义的行为,英伟达不断警告我们。在很多情况下,它只是碰巧起作用

2

0

nglee 7 年前

volatile Test code )

然而,正如在 the answer by Michael Kenzel classic parallel reduction(on page 22) 由NVIDIA自己提供。

由于未来的编译器和内存硬件可能会有不同的工作方式,因此依赖它是危险的。使用 __syncwarp() provided by Michael Kenzel 这应该是一个更好的解决方案。借助 this NVIDIA dev blog article

__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];
}