Реализация критического раздела в CUDA
Я пытаюсь реализовать критический раздел в CUDA использование атомарных инструкций, но я столкнулся с некоторой проблемой. Я создал тестовую программу для показа проблемы:
#include <cuda_runtime.h>
#include <cutil_inline.h>
#include <stdio.h>
__global__ void k_testLocking(unsigned int* locks, int n) {
int id = threadIdx.x % n;
while (atomicExch(&(locks[id]), 1u) != 0u) {} //lock
//critical section would go here
atomicExch(&(locks[id]),0u); //unlock
}
int main(int argc, char** argv) {
//initialize the locks array on the GPU to (0...0)
unsigned int* locks;
unsigned int zeros[10]; for (int i = 0; i < 10; i++) {zeros[i] = 0u;}
cutilSafeCall(cudaMalloc((void**)&locks, sizeof(unsigned int)*10));
cutilSafeCall(cudaMemcpy(locks, zeros, sizeof(unsigned int)*10, cudaMemcpyHostToDevice));
//Run the kernel:
k_testLocking<<<dim3(1), dim3(256)>>>(locks, 10);
//Check the error messages:
cudaError_t error = cudaGetLastError();
cutilSafeCall(cudaFree(locks));
if (cudaSuccess != error) {
printf("error 1: CUDA ERROR (%d) {%s}\n", error, cudaGetErrorString(error));
exit(-1);
}
return 0;
}
Этот код, к сожалению, сильные заморозки моя машина в течение нескольких секунд и наконец выходит, распечатывая сообщение:
fcudaSafeCall() Runtime API error in file <XXX.cu>, line XXX : the launch timed out and was terminated.
что означает, что один из тех циклов с условием продолжения не возвращается, но кажется, что это должно работать.
Как напоминание atomicExch(unsigned int* address, unsigned int val) атомарно устанавливает значение ячейки памяти, сохраненной в адресе к val и возвраты old значение. Таким образом, идея позади моего механизма блокировки состоит в том, что это первоначально 0u, таким образом, один поток должен закончить while цикл и все другие потоки должны ожидать на while цикл, так как они будут читать locks[id] как 1u. Затем, когда поток сделан с критическим разделом, он сбрасывает блокировку к 0u таким образом, другой поток может войти.
Что я пропускаю?
Между прочим, я компилирую с:
nvcc -arch sm_11 -Ipath/to/cuda/C/common/inc XXX.cu
2 ответа
Хорошо, я разобрался, и это еще один из куда-парадигма-краски.
Как знает любой хороший куда-программист (заметьте, я не помню, что это делает меня плохим куда-программистом, я думаю), все нити в искривлении должны исполнять один и тот же код. Если бы не этот факт, то код, который я написал, работал бы идеально. Как бы то ни было, скорее всего, будут два потока в одной и той же искривлении, обращающихся к одной и той же блокировке. Если один из них приобретает блокировку, он просто забывает о выполнении цикла, но не может продолжать цикл, пока все остальные потоки в его искривлении не завершат его. К сожалению, другой поток никогда не завершит цикл, потому что ждет разблокировки первого.
Вот кернел, который безошибочно выполнит этот трюк:
__global__ void k_testLocking(unsigned int* locks, int n) {
int id = threadIdx.x % n;
bool leaveLoop = false;
while (!leaveLoop) {
if (atomicExch(&(locks[id]), 1u) == 0u) {
//critical section
leaveLoop = true;
atomicExch(&(locks[id]),0u);
}
}
}
Кстати, вы должны помнить, что пишет глобальная память и! Читает не завершено, где вы пишете их в коде ... так, чтобы это было практикой, вам нужно добавить глобальную мемляцию, т. Е. __threadfence ()