CUDA разделяет память и синхронизирует деформации - PullRequest
Новая
       27

CUDA разделяет память и синхронизирует деформации

0 голосов
/ 08 января 2019

Следующий код хоста 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");
}
$

Если я скомпилирую test1.cu и запусту его, он даст результат, отличный от test0.cu или test.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.cu и 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.

Ответы [ 2 ]

0 голосов
/ 11 января 2019

Кажется, что я упустил момент, когда объявлял разделяемую память с помощью квалификатора volatile. Это решило проблему. ( Тестовый код )

Однако, как указано в ответе Михаэля Кензеля , такого рода неявного синхронного программирования в основном следует избегать, даже если это введено в классическом параллельном сокращении (на стр. 22). ) предоставлено самой NVIDIA.

Поскольку в будущем аппаратное обеспечение компилятора и памяти может работать по-разному, полагаться на него опасно. Использование __syncwarp() аналогично решению , предоставленному Майклом Кензелем , должно быть лучшим решением. С помощью этой статьи в блоге разработчика 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];
}
0 голосов
/ 08 января 2019

Код вашего устройства имеет неопределенное поведение из-за условий гонки в обоих случаях, при использовании общей памяти или глобальной памяти. У вас есть несколько потоков, которые одновременно читают и изменяют один и тот же объект 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. Он first извлекает all значения, которые ему когда-либо понадобятся, из общей памяти в регистры с использованием условных нагрузок Только после этого он выполняет все добавления и условные хранилища, используя уже загруженные значения. Это совершенно легальная интерпретация вашего кода 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, заключается в том, что синхронизация на уровне деформации была действительно введена в CUDA 9.0 только тогда, когда Volta и «независимое планирование потоков» сделали это необходимостью. До появления CUDA 9.0 синхронное программирование по варпу официально не поддерживалось. Но компиляторы были довольно консервативны, когда дело дошло до взлома кода, как в вашем примере выше. Вероятно, причина в том, что такое «синхронное деформация» программирование (обратите внимание на кавычки) часто было единственным способом приблизиться к пиковой производительности, реальной альтернативы не было, и, таким образом, люди делали это все время. Это все еще было неопределенное поведение, и NVIDIA продолжала предупреждать нас. Во многих случаях это просто работало…

...