Общая память мьютекса с CUDA - добавление к списку элементов

Question

Моя проблема заключается в следующем: у меня есть изображение, в котором я обнаруживаю некоторые интересные точки, используя GPU. Обнаружение представляет собой тяжеловесный тест с точки зрения обработки, однако только около 1 из 25 баллов проходят тест в среднем. Заключительным этапом алгоритма является составление списка точек. На CPU это будет реализовано как:

forall pixels x,y 
{ 
    if(test_this_pixel(x,y)) 
     vector_of_coordinates.push_back(Vec2(x,y)); 
} 

На графическом процессоре у меня есть каждый блок обработки CUDA 16x16 пикселей. Проблема в том, что мне нужно сделать что-то особенное, чтобы в конечном итоге иметь единый консолидированный список точек в глобальной памяти. В настоящий момент я пытаюсь создать локальный список точек в общей памяти на блок, который в конечном итоге будет записан в глобальную память. Я пытаюсь избежать отправки чего-либо обратно в CPU, потому что после этого есть несколько этапов CUDA.

Я ожидал, что могу использовать атомные операции для реализации функции push_back в общей памяти. Однако я не могу заставить это работать. Есть два вопроса. Первая неприятная проблема заключается в том, что я постоянно сталкиваюсь с следующим сбоем компилятора: «nvcc error:« ptxas »умер с статусом 0xC0000005 (ACCESS_VIOLATION)» при использовании атомных операций. Удар или промах, могу ли я что-то скомпилировать. Кто-нибудь знает, что вызывает это?

Следующая ядро ​​будет воспроизвести ошибку:

__global__ void gpu_kernel(int w, int h, RtmPoint *pPoints, int *pCounts) 
{ 
    __shared__ unsigned int test; 
    atomicInc(&test, 1000); 
} 

Во-вторых, мой код, который включает в себя блокировку мьютекса на разделяемой памяти висит на GPU, и я не понимаю, почему:

__device__ void lock(unsigned int *pmutex) 
{ 
    while(atomicCAS(pmutex, 0, 1) != 0); 
} 

__device__ void unlock(unsigned int *pmutex) 
{ 
    atomicExch(pmutex, 0); 
} 

__global__ void gpu_kernel_non_max_suppress(int w, int h, RtmPoint *pPoints, int *pCounts) 
{ 
    __shared__ RtmPoint localPoints[64]; 
    __shared__ int localCount; 
    __shared__ unsigned int mutex; 

    int x = blockIdx.x * blockDim.x + threadIdx.x; 
    int y = blockIdx.y * blockDim.y + threadIdx.y; 

    int threadid = threadIdx.y * blockDim.x + threadIdx.x; 
    int blockid = blockIdx.y * gridDim.x + blockIdx.x; 

    if(threadid==0) 
    { 
     localCount = 0; 
     mutex = 0; 
    } 

    __syncthreads(); 

    if(x<w && y<h) 
    { 
     if(some_test_on_pixel(x,y)) 
     { 
      RtmPoint point; 
      point.x = x; 
      point.y = y; 

      // this is a local push_back operation 
      lock(&mutex); 
      if(localCount<64) // we should never get >64 points per block 
       localPoints[localCount++] = point; 
      unlock(&mutex); 
     } 
    } 

    __syncthreads(); 

    if(threadid==0) 
     pCounts[blockid] = localCount; 
    if(threadid<localCount) 
     pPoints[blockid * 64 + threadid] = localPoints[threadid]; 
} 

В примере код в this site, автору удается успешно использовать атомные операции в общей памяти, поэтому я смущен тем, почему мой случай не работает. Если я прокомментирую блокировку и разблокировку строк, код будет работать нормально, но явно неправильно добавит в список.

Я был бы признателен за некоторые советы о том, почему эта проблема происходит, а также, возможно, если есть лучшее решение для достижения цели, поскольку я все равно обеспокоен проблемами производительности с использованием атомных операций или блокировок мьютексов.

Answer 1

Предлагаю использовать префикс-sum для реализации этой части для увеличения параллелизма. Для этого вам нужно использовать общий массив. В основном префикс-сумма превратит массив (1,1,0,1) в (0,1,2,2,3), т. Е. Рассчитает на месте эксклюзивную сумму, чтобы вы получили сквозную нить писать индексы.

__shared__ uint8_t vector[NUMTHREADS]; 

.... 

bool emit = (x<w && y<h); 
    emit = emit && some_test_on_pixel(x,y); 
__syncthreads(); 
scan(emit, vector); 
if (emit) { 
    pPoints[blockid * 64 + vector[TID]] = point; 
} 

префиксов сумма пример:

template <typename T> 
__device__ uint32 scan(T mark, T *output) { 
#define GET_OUT (pout?output:values) 
#define GET_INP (pin?output:values) 
    __shared__ T values[numWorkers]; 
    int pout=0, pin=1; 
    int tid = threadIdx.x; 

    values[tid] = mark; 

    syncthreads(); 

    for(int offset=1; offset < numWorkers; offset *= 2) { 
    pout = 1 - pout; pin = 1 - pout; 
    syncthreads(); 
    if (tid >= offset) { 
     GET_OUT[tid] = (GET_INP[tid-offset]) +(GET_INP[tid]); 
    } 
    else { 
     GET_OUT[tid] = GET_INP[tid]; 
    } 
    syncthreads(); 
    } 

    if(!pout) 
    output[tid] =values[tid]; 

    __syncthreads(); 

    return output[numWorkers-1]; 

#undef GET_OUT 
#undef GET_INP 
} 

Answer 2

На основе рекомендаций здесь, я включил код, который я использовал в конце концов. Он использует 16x16 пиксельных блоков. Обратите внимание: теперь я пишу данные в одном глобальном массиве, не разбирая его. Я использовал глобальную функцию atomicAdd для вычисления базового адреса для каждого набора результатов. Поскольку это вызвано только один раз за блок, я не нашел слишком много замедления, в то время как я получил намного больше удобства, выполнив это. Я также избегаю использования общих буферов для ввода и вывода prefix_sum. Значение GlobalCount равно нулю перед вызовом ядра.

#define BLOCK_THREADS 256 

__device__ int prefixsum(int threadid, int data) 
{ 
    __shared__ int temp[BLOCK_THREADS*2]; 

    int pout = 0; 
    int pin = 1; 

    if(threadid==BLOCK_THREADS-1) 
     temp[0] = 0; 
    else 
     temp[threadid+1] = data; 

    __syncthreads(); 

    for(int offset = 1; offset<BLOCK_THREADS; offset<<=1) 
    { 
     pout = 1 - pout; 
     pin = 1 - pin; 

     if(threadid >= offset) 
      temp[pout * BLOCK_THREADS + threadid] = temp[pin * BLOCK_THREADS + threadid] + temp[pin * BLOCK_THREADS + threadid - offset]; 
     else 
      temp[pout * BLOCK_THREADS + threadid] = temp[pin * BLOCK_THREADS + threadid]; 

     __syncthreads(); 
    } 

    return temp[pout * BLOCK_THREADS + threadid]; 
} 

__global__ void gpu_kernel(int w, int h, RtmPoint *pPoints, int *pGlobalCount) 
{ 
    __shared__ int write_base; 

    int x = blockIdx.x * blockDim.x + threadIdx.x; 
    int y = blockIdx.y * blockDim.y + threadIdx.y; 

    int threadid = threadIdx.y * blockDim.x + threadIdx.x; 
    int valid = 0; 

    if(x<w && y<h) 
    { 
     if(test_pixel(x,y)) 
     { 
      valid = 1; 
     } 
    } 

    int index = prefixsum(threadid, valid); 

    if(threadid==BLOCK_THREADS-1) 
    { 
     int total = index + valid; 
     if(total>64) 
      total = 64; // global output buffer is limited to 64 points per block 
     write_base = atomicAdd(pGlobalCount, total); // get a location to write them out 
    } 

    __syncthreads(); // ensure write_base is valid for all threads 

    if(valid) 
    { 
     RtmPoint point; 
     point.x = x; 
     point.y = y; 
     if(index<64) 
      pPoints[write_base + index] = point; 
    } 
}