TL; DR: код поврежден при проверке. Модель потоков CUDA не гарантирует продвижение вперед какого-либо конкретного потока , за исключением случаев, когда соблюдаются следующие условия:
- Прогресс вперед будет доставлен как минимум в 1 (выдаваемый, не выбывший) ) потока, при условии, что есть хотя бы один из них.
- будет соблюдена семантика барьера выполнения
Какой поток или потоки будут выбраны для элемента 1, не определены моделью программирования CUDA , До тех пор, пока программист не получит явное управление с использованием барьеров выполнения, модель потоков CUDA может свободно планировать один поток так долго, как ему нравится, до тех пор, пока этот поток не прекратит работу или не достигнет явного барьера выполнения.
Поскольку предоставленный код не имеет никаких барьеров выполнения, планировщик работ CUDA (в отношении семантики CUDA) может свободно планировать, например, поток 0 и никаких других потоков. Если мы применим эту концепцию к предоставленному коду, очевидно, что поток 0, если он выполняется сам по себе, будет показывать бесконечное значение l oop.
Дольше:
Это как раз и есть наблюдаемое поведение, хотя, если бы это был я, я бы не связал их. Причина зависания (согласно тому, как я бы попытался это описать) не в «правильности», этот код зависит от гарантий, которые не предоставляются моделью программирования CUDA », хотя я считаю, что это верное утверждение. Чтобы понять причину зависания, я предлагаю проверить поведение машины низкого уровня с учетом SASS (код сборки машины). На самом деле у меня нет возможности преследовать эту топи c до изнеможения, поэтому я представлю ограниченное представление об этом.
Зачем проводить это различие? Поскольку относительно небольшие изменения в предоставленном коде, которые на самом деле не решают проблему корректности, могут привести к тому, что компилятор не будет зависать. Отсутствие осторожного обращения может привести к выводу, что, поскольку он не зависает, все должно быть в порядке. Дело в том, что зависание кода зависит от того, является ли он правильным. Я доказал это себе. Однако я не могу предоставить этот код. Правильно сделать правильный дизайн кода. Ниже приведена моя попытка сделать это.
Прежде чем мы углубимся в SASS, я хотел бы указать на еще один недостаток в коде. Компилятор CUDA может «оптимизировать» любые глобальные данные в регистры, сохраняя при этом однопоточность семанти c корректность. Компилятор в основном имеет один поток, и это может сбить с толку программистов, которые зависят от взаимодействия между потоками (как этот код делает). Для корректности в этом коде данные, измененные потоком x, должны быть видны (в конце концов) потоку x-1. Этот вид видимости между потоками не гарантируется моделью программирования CUDA, и компилятор обычно не обеспечивает ее. Для правильности необходимо сообщить компилятору, чтобы эти данные были видны, и заказать загрузки и сохранения, чтобы это произошло. Есть несколько способов сделать это sh. Для простоты я предлагаю пометить данные volatile, хотя это может быть возможно сделать с помощью барьеров выполнения (например, __syncthreads(), __syncwarp()), которые также имеют встроенные барьеры памяти. в . Независимо от метода, выбранного для обеспечения видимости данных между потоками, без него код нарушается независимо от других соображений.
Поэтому перед погружением в SASS я предложу следующую модификацию прилагаемого кода вместе с SASS, следующим за ним:
$ cat t1691.cu
__device__ volatile float data[10] = {0,0,0,0,0,0,0,0,0,1};
__global__ void gradually_set_global_data() {
while (1) {
if (data[threadIdx.x + 1]) {
atomicAdd((float *)&data[threadIdx.x], data[threadIdx.x + 1]);
break;
}
}
}
int main() {
gradually_set_global_data<<<1, 9>>>();
cudaDeviceReset();
return 0;
}
$ nvcc -o t1691 t1691.cu
$ cuobjdump -sass ./t1691
Fatbin elf code:
================
arch = sm_30
code version = [1,7]
producer = <unknown>
host = linux
compile_size = 64bit
code for sm_30
Fatbin elf code:
================
arch = sm_30
code version = [1,7]
producer = <unknown>
host = linux
compile_size = 64bit
code for sm_30
Function : _Z25gradually_set_global_datav
.headerflags @"EF_CUDA_SM30 EF_CUDA_PTX_SM(EF_CUDA_SM30)"
/* 0x22f2c04272004307 */
/*0008*/ MOV R1, c[0x0][0x44]; /* 0x2800400110005de4 */
/*0010*/ S2R R0, SR_TID.X; /* 0x2c00000084001c04 */
/*0018*/ MOV32I R3, 0x0; /* 0x180000000000dde2 */
/*0020*/ SSY 0x68; /* 0x6000000100001c07 */
/*0028*/ IMAD R2.CC, R0, 0x4, R3; /* 0x2007c00010009ca3 */
/*0030*/ MOV32I R3, 0x0; /* 0x180000000000dde2 */
/*0038*/ IMAD.U32.U32.HI.X R3, R0, 0x4, R3; /* 0x2086c0001000dc43 */
/* 0x22f043f2f2e2c3f7 */
/*0048*/ LD.E.CV R0, [R2+0x4]; /* 0x8400000010201f85 */
/*0050*/ FSETP.NEU.AND P0, PT, R0, RZ, PT; /* 0x268e0000fc01dc00 */
/*0058*/ @!P0 BRA 0x40; /* 0x4003ffff800021e7 */
/*0060*/ NOP.S; /* 0x4000000000001df4 */
/*0068*/ LD.E.CV R4, [R2+0x4]; /* 0x8400000010211f85 */
/*0070*/ RED.E.ADD.F32.FTZ.RN [R2], R4; /* 0x2c00000000211e05 */
/*0078*/ EXIT; /* 0x8000000000001de7 */
/*0080*/ BRA 0x80; /* 0x4003ffffe0001de7 */
/*0088*/ NOP; /* 0x4000000000001de4 */
/*0090*/ NOP; /* 0x4000000000001de4 */
/*0098*/ NOP; /* 0x4000000000001de4 */
/*00a0*/ NOP; /* 0x4000000000001de4 */
/*00a8*/ NOP; /* 0x4000000000001de4 */
/*00b0*/ NOP; /* 0x4000000000001de4 */
/*00b8*/ NOP; /* 0x4000000000001de4 */
.........................................
Fatbin ptx code:
================
arch = sm_30
code version = [6,4]
producer = <unknown>
host = linux
compile_size = 64bit
compressed
$
Согласно моим тестам на cc3.5 и Устройства cc7.0, приведенный выше код все еще зависает, поэтому мы не изменили его наблюдаемое поведение с этими изменениями. (обратите внимание, что приведенный выше код SASS предназначен для cc3.0, скомпилированного с CUDA 10.1.243).
Код будет демонстрировать отклоняющееся от деформации поведение, и IMO это важно для понимания зависания, поэтому мы сосредоточимся на условная область кода SASS:
/*0038*/ IMAD.U32.U32.HI.X R3, R0, 0x4, R3; /* 0x2086c0001000dc43 */
/* 0x22f043f2f2e2c3f7 */
/*0048*/ LD.E.CV R0, [R2+0x4]; /* 0x8400000010201f85 */
/*0050*/ FSETP.NEU.AND P0, PT, R0, RZ, PT; /* 0x268e0000fc01dc00 */
/*0058*/ @!P0 BRA 0x40; /* 0x4003ffff800021e7 */
/*0060*/ NOP.S; /* 0x4000000000001df4 */
/*0068*/ LD.E.CV R4, [R2+0x4]; /* 0x8400000010211f85 */
/*0070*/ RED.E.ADD.F32.FTZ.RN [R2], R4; /* 0x2c00000000211e05 */
/*0078*/ EXIT; /* 0x8000000000001de7 */
По строке 0038 все работы по настройке завершены. В строке 0048 поток загружает свое значение __device__ data из глобальной памяти, и в строке 0050 выполняется условный тест, а в строке 0058 - условная ветвь. Если поток получил ненулевое значение, он продолжит работу. на линию 0060 (и, в конце концов, выполните операцию atomi c и выйдите). Если нет, то go вернется к строке 0040, чтобы повторить загрузку и тестирование.
Теперь мы наблюдаем зависание. Пропускаемые потоки и потоки, которые не проходят условный тест, не планируются одновременно планировщиком деформации. Это должно запланировать одну группу (например, прохождение) или другую (например, неудача). Это же решение должно приниматься планировщиком основы несколько раз. Если мы наблюдаем зависание, единственный возможный вывод состоит в том, что потоки, которые не проходят условный тест, многократно планируются (выбираются для выпуска), а потоки, которые проходят условный тест, не планируются.
Это законно, и любые выводы о том, что проходящие потоки должны планироваться «в конечном итоге», являются неверными выводами, согласно модели программирования CUDA и этой конструкции кода. Единственный способ гарантировать, что проходящие потоки будут запланированы, - передать планировщику деформации ситуацию, при которой для него не будет другого выбора, в соответствии с принципом 1 в верхней части этого ответа.
(В сторону Обратите внимание, что мы могли также заметить, что планировщик деформации выбрал прохождение потоков вместо сбойных потоков для планирования / выпуска. В этом случае, поскольку эти проходящие потоки в конечном итоге завершаются / удаляются в этой реализации, я ожидаю, что это привело бы к код, который не зависает. В конечном итоге все проходящие потоки будут удалены, и планировщик деформации будет вынужден, с помощью пункта 1 вверху этого ответа, начать планировать сбойные потоки. Не зависание здесь будет равно достоверное и возможное наблюдение , в той степени, в которой здесь изложены характеристики планирования деформации. Но основывать любые выводы о правильности на этом результате все равно будет неправильно.)
Если продолжить эту идею, то можно спросить: есть законный способ осознать этот шаблон? Я предполагаю, что теперь мы знаем, что нам, вероятно, понадобятся барьеры исполнения, если мы собираемся сделать эту работу. Давайте выберем __syncwarp(). Для этого барьера, законное использование барьера, как правило, потребует, чтобы у нас был полностью неповрежденный перекос (или перекос). Таким образом, нам нужно будет переделать код, чтобы активировать полную деформацию, но только нужные потоки (всего 9) выполняют «работу».
Ниже приведен один из возможных способов достижения этого. Я уверен, что есть и другие способы. Этот код, в соответствии с моим тестированием, не зависает на устройствах cc3.5 или cc7.0:
__device__ volatile float data[10] = {0,0,0,0,0,0,0,0,0,1};
__global__ void gradually_set_global_data(int sz) {
int tflag = (threadIdx.x < sz) ? 1:0;
unsigned wflag = 1;
while (wflag) {
if (tflag)
if (data[threadIdx.x + 1]) {
atomicAdd((float *)&data[threadIdx.x], data[threadIdx.x + 1]);
tflag = 0;
}
__syncwarp();
wflag = __ballot_sync(0xFFFFFFFFU, tflag);
}
}
int main() {
gradually_set_global_data<<<1, 32>>>(9);
cudaDeviceReset();
return 0;
}
Обратите внимание, что, если мы хотим подняться еще ближе к предоставленному коду, вышеприведенное может быть изменено с a while(1) l oop, а внутри l oop выведите break, если wflag равно нулю (после операции голосования). Я не думаю, что в этой реализации есть какая-либо значимая разница.
Я все еще не претендую на правильность этого кода или любого другого кода, который я публикую. Любой, кто использует любой код, который я публикую, делает это на свой страх и риск. Я просто утверждаю, что попытался устранить недостатки, обнаруженные в первоначальной публикации, и приведу некоторые объяснения этого. Я не утверждаю, что мой код не имеет дефектов или что он подходит для какой-либо конкретной цели. Используйте его (или нет) на свой страх и риск.