Стек устройств CUDA и синхронизация; Инструкция SSY
Редактировать: этот вопрос является новой версией оригинала, поэтому первые несколько ответов могут быть неактуальными.
Мне любопытно, какое влияние вызов функции устройства с принудительным отсутствием встраивания оказывает на синхронизацию внутри функции устройства. У меня есть простое тестовое ядро, которое иллюстрирует рассматриваемое поведение.
Ядро берет буфер и передает его функции устройства вместе с общим буфером и индикаторной переменной, которая идентифицирует отдельный поток как поток "босс". Функция устройства имеет расходящийся код: поток босса сначала тратит время на выполнение простых операций с общим буфером, затем записывает в глобальный буфер. После вызова синхронизации все потоки записывают в глобальный буфер. После вызова ядра хост печатает содержимое глобального буфера. Вот код:
Код CUDA:
test_main.cu
#include<cutil_inline.h>
#include "test_kernel.cu"
int main()
{
int scratchBufferLength = 100;
int *scratchBuffer;
int *d_scratchBuffer;
int b = 1;
int t = 64;
// copy scratch buffer to device
scratchBuffer = (int *)calloc(scratchBufferLength,sizeof(int));
cutilSafeCall( cudaMalloc(&d_scratchBuffer,
sizeof(int) * scratchBufferLength) );
cutilSafeCall( cudaMemcpy(d_scratchBuffer, scratchBuffer,
sizeof(int)*scratchBufferLength, cudaMemcpyHostToDevice) );
// kernel call
testKernel<<<b, t>>>(d_scratchBuffer);
cudaThreadSynchronize();
// copy data back to host
cutilSafeCall( cudaMemcpy(scratchBuffer, d_scratchBuffer,
sizeof(int) * scratchBufferLength, cudaMemcpyDeviceToHost) );
// print results
printf("Scratch buffer contents: \t");
for(int i=0; i < scratchBufferLength; ++i)
{
if(i % 25 == 0)
printf("\n");
printf("%d ", scratchBuffer[i]);
}
printf("\n");
//cleanup
cudaFree(d_scratchBuffer);
free(scratchBuffer);
return 0;
}
test_kernel.cu
#ifndef __TEST_KERNEL_CU
#define __TEST_KERNEL_CU
#define IS_BOSS() (threadIdx.x == blockDim.x - 1)
__device__
__noinline__
void testFunc(int *sA, int *scratchBuffer, bool isBoss) {
if(isBoss) { // produces unexpected output-- "broken" code
//if(IS_BOSS()) { // produces expected output-- "working" code
for (int c = 0; c < 10000; c++) {
sA[0] = 1;
}
}
if(isBoss) {
scratchBuffer[0] = 1;
}
__syncthreads();
scratchBuffer[threadIdx.x ] = threadIdx.x;
return;
}
__global__
void testKernel(int *scratchBuffer)
{
__shared__ int sA[4];
bool isBoss = IS_BOSS();
testFunc(sA, scratchBuffer, isBoss);
return;
}
#endif
Я скомпилировал этот код из CUDA SDK, чтобы воспользоваться функциями "cutilsafecall()" в test_main.cu, но, конечно, их можно было бы удалить, если вы хотите скомпилировать вне SDK. Я скомпилировал с помощью CUDA Driver/Toolkit версии 4.0, вычислил возможности 2.0, и код был запущен на GeForce GTX 480, который имеет архитектуру Fermi.
Ожидаемый результат
0 1 2 3 ... blockDim.x-1
Тем не менее, я получаю вывод
1 1 2 3 ... blockDim.x-1
Похоже, это указывает на то, что поток босса выполнил условное "scratchBuffer[0] = 1;" оператор ПОСЛЕ того, как все потоки выполнят "scratchBuffer[threadIdx.x] = threadIdx.x;" оператора, даже если они разделены барьером __syncthreads().
Это происходит, даже если поток босса получает команду записать значение часового в позицию буфера потока в том же самом искажении; часовой является последним значением, присутствующим в буфере, а не соответствующим threadIdx.x .
Одна из модификаций, которая приводит к тому, что код выдает ожидаемый результат, - это изменение условного оператора
if(isBoss) {
в
if (IS_BOSS ()) {
; то есть, чтобы изменить переменную, контролирующую расхождение, с сохранения в регистре параметров на вычисление в макрофункции. (Обратите внимание на комментарии к соответствующим строкам в исходном коде.) Именно на этом конкретном изменении я сосредоточился, чтобы попытаться отследить проблему. При рассмотрении дизассемблированных.cubins ядра с условным условием 'isBoss' (т. Е. Неработающим кодом) и условным условием 'IS_BOSS()' (т. Е. Рабочим кодом), наиболее заметным отличием в инструкциях кажется отсутствие инструкция SSY в разобранном битом коде.
Вот дизассемблированные ядра, сгенерированные путем дизассемблирования файлов.cubin с помощью "cuobjdump -sass test_kernel.cubin" . все до первого "EXIT" - это ядро, а все, что после этого - функция устройства. Единственные отличия в функции устройства.
РАЗБОР КОДА ОБЪЕКТА:
"сломанный" код
code for sm_20
Function : _Z10testKernelPi
/*0000*/ /*0x00005de428004404*/ MOV R1, c [0x1] [0x100];
/*0008*/ /*0x20009de428004000*/ MOV R2, c [0x0] [0x8];
/*0010*/ /*0x84001c042c000000*/ S2R R0, SR_Tid_X;
/*0018*/ /*0xfc015de428000000*/ MOV R5, RZ;
/*0020*/ /*0x00011de428004000*/ MOV R4, c [0x0] [0x0];
/*0028*/ /*0xfc209c034800ffff*/ IADD R2, R2, 0xfffff;
/*0030*/ /*0x9001dde428004000*/ MOV R7, c [0x0] [0x24];
/*0038*/ /*0x80019de428004000*/ MOV R6, c [0x0] [0x20];
/*0040*/ /*0x08001c03110e0000*/ ISET.EQ.U32.AND R0, R0, R2, pt;
/*0048*/ /*0x01221f841c000000*/ I2I.S32.S32 R8, -R0;
/*0050*/ /*0x2001000750000000*/ CAL 0x60;
/*0058*/ /*0x00001de780000000*/ EXIT;
/*0060*/ /*0x20201e841c000000*/ I2I.S32.S8 R0, R8;
/*0068*/ /*0xfc01dc231a8e0000*/ ISETP.NE.AND P0, pt, R0, RZ, pt;
/*0070*/ /*0xc00021e740000000*/ @!P0 BRA 0xa8;
/*0078*/ /*0xfc001de428000000*/ MOV R0, RZ;
/*0080*/ /*0x04001c034800c000*/ IADD R0, R0, 0x1;
/*0088*/ /*0x04009de218000000*/ MOV32I R2, 0x1;
/*0090*/ /*0x4003dc231a8ec09c*/ ISETP.NE.AND P1, pt, R0, 0x2710, pt;
/*0098*/ /*0x00409c8594000000*/ ST.E [R4], R2;
/*00a0*/ /*0x600005e74003ffff*/ @P1 BRA 0x80;
/*00a8*/ /*0x040001e218000000*/ @P0 MOV32I R0, 0x1;
/*00b0*/ /*0x0060008594000000*/ @P0 ST.E [R6], R0;
/*00b8*/ /*0xffffdc0450ee0000*/ BAR.RED.POPC RZ, RZ;
/*00c0*/ /*0x84001c042c000000*/ S2R R0, SR_Tid_X;
/*00c8*/ /*0x10011c03200dc000*/ IMAD.U32.U32 R4.CC, R0, 0x4, R6;
/*00d0*/ /*0x10009c435000c000*/ IMUL.U32.U32.HI R2, R0, 0x4;
/*00d8*/ /*0x08715c4348000000*/ IADD.X R5, R7, R2;
/*00e0*/ /*0x00401c8594000000*/ ST.E [R4], R0;
/*00e8*/ /*0x00001de790000000*/ RET;
.................................
"рабочий" код
code for sm_20
Function : _Z10testKernelPi
/*0000*/ /*0x00005de428004404*/ MOV R1, c [0x1] [0x100];
/*0008*/ /*0x20009de428004000*/ MOV R2, c [0x0] [0x8];
/*0010*/ /*0x84001c042c000000*/ S2R R0, SR_Tid_X;
/*0018*/ /*0xfc015de428000000*/ MOV R5, RZ;
/*0020*/ /*0x00011de428004000*/ MOV R4, c [0x0] [0x0];
/*0028*/ /*0xfc209c034800ffff*/ IADD R2, R2, 0xfffff;
/*0030*/ /*0x9001dde428004000*/ MOV R7, c [0x0] [0x24];
/*0038*/ /*0x80019de428004000*/ MOV R6, c [0x0] [0x20];
/*0040*/ /*0x08001c03110e0000*/ ISET.EQ.U32.AND R0, R0, R2, pt;
/*0048*/ /*0x01221f841c000000*/ I2I.S32.S32 R8, -R0;
/*0050*/ /*0x2001000750000000*/ CAL 0x60;
/*0058*/ /*0x00001de780000000*/ EXIT;
/*0060*/ /*0x20009de428004000*/ MOV R2, c [0x0] [0x8];
/*0068*/ /*0x8400dc042c000000*/ S2R R3, SR_Tid_X;
/*0070*/ /*0x20201e841c000000*/ I2I.S32.S8 R0, R8;
/*0078*/ /*0x4000000760000001*/ SSY 0xd0;
/*0080*/ /*0xfc209c034800ffff*/ IADD R2, R2, 0xfffff;
/*0088*/ /*0x0831dc031a8e0000*/ ISETP.NE.U32.AND P0, pt, R3, R2, pt;
/*0090*/ /*0xc00001e740000000*/ @P0 BRA 0xc8;
/*0098*/ /*0xfc009de428000000*/ MOV R2, RZ;
/*00a0*/ /*0x04209c034800c000*/ IADD R2, R2, 0x1;
/*00a8*/ /*0x04021de218000000*/ MOV32I R8, 0x1;
/*00b0*/ /*0x4021dc231a8ec09c*/ ISETP.NE.AND P0, pt, R2, 0x2710, pt;
/*00b8*/ /*0x00421c8594000000*/ ST.E [R4], R8;
/*00c0*/ /*0x600001e74003ffff*/ @P0 BRA 0xa0;
/*00c8*/ /*0xfc01dc33190e0000*/ ISETP.EQ.AND.S P0, pt, R0, RZ, pt;
/*00d0*/ /*0x040021e218000000*/ @!P0 MOV32I R0, 0x1;
/*00d8*/ /*0x0060208594000000*/ @!P0 ST.E [R6], R0;
/*00e0*/ /*0xffffdc0450ee0000*/ BAR.RED.POPC RZ, RZ;
/*00e8*/ /*0x10311c03200dc000*/ IMAD.U32.U32 R4.CC, R3, 0x4, R6;
/*00f0*/ /*0x10309c435000c000*/ IMUL.U32.U32.HI R2, R3, 0x4;
/*00f8*/ /*0x84001c042c000000*/ S2R R0, SR_Tid_X;
/*0100*/ /*0x08715c4348000000*/ IADD.X R5, R7, R2;
/*0108*/ /*0x00401c8594000000*/ ST.E [R4], R0;
/*0110*/ /*0x00001de790000000*/ RET;
.................................
Инструкция "SSY" присутствует в рабочем коде, но не в сломанном коде. Руководство cuobjdump описывает инструкцию с помощью "Установить точку синхронизации; используется перед потенциально расходящимися инструкциями". Это заставляет меня думать, что по какой-то причине компилятор не распознает возможность расхождения в неработающем коде.
Я также обнаружил, что если я закомментирую директиву __noinline__, то код выдает ожидаемый результат, и действительно сборка, создаваемая в противном случае "сломанной" и "рабочей" версиями, точно идентична. Таким образом, это заставляет меня думать, что когда переменная передается через стек вызовов, эта переменная не может использоваться для управления расхождением и последующим вызовом синхронизации; компилятор, похоже, не распознает возможность расхождения в этом случае и поэтому не вставляет инструкцию "SSY". Кто-нибудь знает, действительно ли это законное ограничение CUDA, и если да, если это где-то задокументировано?
Заранее спасибо.
1 ответ
Похоже, это просто ошибка компилятора, исправленная в CUDA 4.1/4.2. Не воспроизводит для спрашивающего на CUDA 4.2.