Cuda Mutex, почему тупик?
Я пытаюсь реализовать мьютекс на основе атома.
Мне это удалось, но у меня есть один вопрос о перекосах / тупиках.
Этот код работает хорошо.
bool blocked = true;
while(blocked) {
if(0 == atomicCAS(&mLock, 0, 1)) {
index = mSize++;
doCriticJob();
atomicExch(&mLock, 0);
blocked = false;
}
}
Но этот не...
while(true) {
if(0 == atomicCAS(&mLock, 0, 1)) {
index = mSize++;
doCriticJob();
atomicExch(&mLock, 0);
break;
}
}
Я думаю, что это позиция выхода из цикла. В первом случае выход происходит там, где есть условие, во втором - в конце if, поэтому поток ожидает другой цикл завершения деформации, но другие потоки также ожидают первый поток... Но я думаю, что Я ошибаюсь, так что если вы можете объяснить мне:).
Спасибо!
1 ответ
Есть и другие вопросы по мьютексам. Возможно, вы захотите взглянуть на некоторые из них. Поиск по "критической секции cuda", например.
Предполагать, что кто-то сработает, а что-то не получится, потому что казалось, что это работает для вашего теста, опасно Управление мьютексами или критическими секциями, особенно когда переговоры находятся между потоками в одной и той же структуре, общеизвестно сложно и хрупко. Общий совет - избегать этого. Как обсуждалось в другом месте, если вы должны использовать мьютексы или критические секции, иметь единый поток в согласовании потокового блока для любого потока, который в этом нуждается, затем управляйте поведением внутри потокового блока, используя механизмы синхронизации внутри потока, такие как
__syncthreads(),На этот вопрос (IMO) невозможно ответить, не глядя на то, как компилятор упорядочивает различные пути выполнения. Поэтому нам нужно взглянуть на код SASS (машинный код). Вы можете использовать двоичные утилиты cuda для этого и, вероятно, захотите ссылаться как на ссылку PTX, так и на ссылку SASS. Это также означает, что вам нужен полный код, а не только предоставленные вами фрагменты.
Вот мой код для анализа:
$ cat t830.cu
#include <stdio.h>
__device__ int mLock = 0;
__device__ void doCriticJob(){
}
__global__ void kernel1(){
int index = 0;
int mSize = 1;
while(true) {
if(0 == atomicCAS(&mLock, 0, 1)) {
index = mSize++;
doCriticJob();
atomicExch(&mLock, 0);
break;
}
}
}
__global__ void kernel2(){
int index = 0;
int mSize = 1;
bool blocked = true;
while(blocked) {
if(0 == atomicCAS(&mLock, 0, 1)) {
index = mSize++;
doCriticJob();
atomicExch(&mLock, 0);
blocked = false;
}
}
}
int main(){
kernel2<<<4,128>>>();
cudaDeviceSynchronize();
}
kernel1 мое представление вашего кода взаимоблокировки, и kernel2 мое представление вашего "рабочего" кода. Когда я компилирую это на Linux под CUDA 7 и запускаю на устройстве cc2.0 (Quadro5000), если я позвоню kernel1 код заблокируется, и если я позвоню kernel2 (как показано) это не так.
я использую cuobjdump -sass сбросить машинный код:
$ cuobjdump -sass ./t830
Fatbin elf code:
================
arch = sm_20
code version = [1,7]
producer = <unknown>
host = linux
compile_size = 64bit
code for sm_20
Fatbin elf code:
================
arch = sm_20
code version = [1,7]
producer = cuda
host = linux
compile_size = 64bit
code for sm_20
Function : _Z7kernel1v
.headerflags @"EF_CUDA_SM20 EF_CUDA_PTX_SM(EF_CUDA_SM20)"
/*0000*/ MOV R1, c[0x1][0x100]; /* 0x2800440400005de4 */
/*0008*/ MOV32I R4, 0x1; /* 0x1800000004011de2 */
/*0010*/ SSY 0x48; /* 0x60000000c0000007 */
/*0018*/ MOV R2, c[0xe][0x0]; /* 0x2800780000009de4 */
/*0020*/ MOV R3, c[0xe][0x4]; /* 0x280078001000dde4 */
/*0028*/ ATOM.E.CAS R0, [R2], RZ, R4; /* 0x54080000002fdd25 */
/*0030*/ ISETP.NE.AND P0, PT, R0, RZ, PT; /* 0x1a8e0000fc01dc23 */
/*0038*/ @P0 BRA 0x18; /* 0x4003ffff600001e7 */
/*0040*/ NOP.S; /* 0x4000000000001df4 */
/*0048*/ ATOM.E.EXCH RZ, [R2], RZ; /* 0x547ff800002fdd05 */
/*0050*/ EXIT; /* 0x8000000000001de7 */
............................
Function : _Z7kernel2v
.headerflags @"EF_CUDA_SM20 EF_CUDA_PTX_SM(EF_CUDA_SM20)"
/*0000*/ MOV R1, c[0x1][0x100]; /* 0x2800440400005de4 */
/*0008*/ MOV32I R0, 0x1; /* 0x1800000004001de2 */
/*0010*/ MOV32I R3, 0x1; /* 0x180000000400dde2 */
/*0018*/ MOV R4, c[0xe][0x0]; /* 0x2800780000011de4 */
/*0020*/ MOV R5, c[0xe][0x4]; /* 0x2800780010015de4 */
/*0028*/ ATOM.E.CAS R2, [R4], RZ, R3; /* 0x54061000004fdd25 */
/*0030*/ ISETP.NE.AND P1, PT, R2, RZ, PT; /* 0x1a8e0000fc23dc23 */
/*0038*/ @!P1 MOV R0, RZ; /* 0x28000000fc0025e4 */
/*0040*/ @!P1 ATOM.E.EXCH RZ, [R4], RZ; /* 0x547ff800004fe505 */
/*0048*/ LOP.AND R2, R0, 0xff; /* 0x6800c003fc009c03 */
/*0050*/ I2I.S32.S16 R2, R2; /* 0x1c00000008a09e84 */
/*0058*/ ISETP.NE.AND P0, PT, R2, RZ, PT; /* 0x1a8e0000fc21dc23 */
/*0060*/ @P0 BRA 0x18; /* 0x4003fffec00001e7 */
/*0068*/ EXIT; /* 0x8000000000001de7 */
............................
Fatbin ptx code:
================
arch = sm_20
code version = [4,2]
producer = cuda
host = linux
compile_size = 64bit
compressed
$
Учитывая единственную деформацию с любым кодом, все потоки должны получить блокировку (через atomicCAS) один раз, чтобы код успешно завершился. При использовании любого кода только один поток в деформации может получить блокировку в любой момент времени, и для того, чтобы другие потоки в деформации (позднее) получили блокировку, этот поток должен иметь возможность снять ее (через atomicExch).
Ключевое различие между этими реализациями заключается в том, как компилятор запланировал atomicExch инструкция относительно условных веток.
Давайте рассмотрим код "тупика" (kernel1). В этом случае ATOM.E.EXCH инструкция не происходит до тех пор, пока не будет завершена одна (и только) условная ветвь @P0 BRA 0x18;) инструкция. Условная ветвь в коде CUDA представляет возможную точку расхождения деформации, и выполнение после расхождения деформации в некоторой степени не определено и зависит от специфики машины. Но, учитывая эту неопределенность, возможно, что поток, который получил блокировку, будет ждать, пока другие потоки завершат свои ветви, прежде чем выполнить atomicExch инструкция, что означает, что другие потоки не смогут получить блокировку, и у нас есть тупик.
Если затем сравнить это с "рабочим" кодом, мы увидим, что как только ATOM.E.CAS инструкция выпущена, между этой точкой и точкой, в которой ATOM.E.EXCH инструкция выпущена, таким образом освобождая только что приобретенную блокировку. Так как каждый поток, который получает блокировку (через ATOM.E.CAS) выпустит его (через ATOM.E.EXCH) до того, как произойдет какое-либо условное ветвление, нет никакой возможности (учитывая эту реализацию кода) для типа тупика, который был засвидетельствован ранее (с kernel1) происходить.
(@P0 это форма предикации, и вы можете прочитать об этом в справочнике по PTX, чтобы понять, как это может привести к условному ветвлению.)
ПРИМЕЧАНИЕ. Я считаю, что оба эти кода опасны и, возможно, имеют недостатки. Несмотря на то, что текущие тесты, кажется, не раскрывают проблему с "рабочим" кодом, я думаю, что вполне возможно, что будущий компилятор CUDA решит иначе планировать вещи и нарушать этот код. Возможно даже, что компиляция для другой архитектуры машины может привести к другому коду. Я считаю механизм, подобный этому, более надежным, который полностью исключает конфликт внутри варпа. Однако даже такой механизм может привести к взаимным блокировкам между потоками. Любой мьютекс должен использоваться при определенных ограничениях программирования и использования.