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