Фильтрация в 3D пространстве с помощью CUDA, горизонтальный доступ быстрее, чем вертикальный доступ?

Question

Я пытаюсь применить фильтр 1x3, 3x1 в 3D-структуре (том).
Например, если имеется 20 (смещ_по_столбцы) х 10 (строк) х 10 (глубина) структуры,

for(int depth = 0; depth < 10; depth++) 
    Apply image filter(depth); 

применить фильтр к 2d изображений (20x10) 10 раз. каждый фрагмент изображения отличается.

во-первых, я выделить 3D структуру, как

// COLS = 450, ROWS = 375, MAX_DISPARITY = 60 
cudaPitchedPtr volume; 
cudaExtent volumeExtent = make_cudaExtent(COLS, ROWS, MAX_DISPARITY); 
HANDLE_ERROR(cudaMalloc3D(&volume, volumeExtent)); 

и памяти равным нулю для стабильного производства. пока все хорошо, пока копировать изображение в том объеме.

При применении фильтра 3x1, как показано ниже, его расчет составлял 6 мсек.

Apply_3by1 << <ROWS, COLS, COLS>> > (volume, COLS, ROWS); 

__global__ void Apply_3by1 (cudaPitchedPtr src, unsigned int COLS, unsigned int ROWS) 
{ 
    const unsigned int x = threadIdx.x; 
    const unsigned int y = blockIdx.x; 

    extern __shared__ unsigned char SharedMemory[]; 
    for (int dispCnt = 0; dispCnt < MAX_DISPARITY; dispCnt++) 
    { 
     if (x < dispCnt) continue;//exception for my algorithm. 

     unsigned char dst_val = *GET_UCHAR_PTR_3D(src, x, y, dispCnt); 
     SharedMemory[x] = dst_val; 
     __syncthreads(); 


     unsigned char left; 
     int leftIdx = x - 3; 
     if (leftIdx < 0)//index underflow 
      left = 0; 
     else 
      left = SharedMemory[leftIdx]; 


     unsigned char right;//index overflow 
     int rightIdx = x + 3; 
     if (COLS < rightIdx) 
      right = 0; 
     else 
      right = SharedMemory[rightIdx]; 


     *GET_UCHAR_PTR_3D(src, x, y, dispCnt) += left + right; 
    } 
} 

, но когда я применяю вертикальное направление фильтра 1x3, время его вычисления составляло 46 мсек.

Apply_1by3 << <COLS, ROWS, ROWS >> > (volume, COLS, ROWS); 

__global__ void Apply_1by3 (cudaPitchedPtr src, unsigned int COLS, unsigned int ROWS) 
{ 
    const unsigned int x = threadIdx.x; 
    const unsigned int y = blockIdx.x; 

    extern __shared__ unsigned char SharedMemory[]; 
    for (int dispCnt = 0; dispCnt < MAX_DISPARITY; dispCnt++) 
    { 
     unsigned char my_val = *GET_UCHAR_PTR_3D(src, y, x, dispCnt); 
     SharedMemory[x] = my_val; 
     __syncthreads(); 

     if (y < dispCnt) continue; 

     int topIdx = x - 3; 
     unsigned char top_value; 
     if (topIdx < 0) 
      top_value = 0; 
     else 
      top_value = SharedMemory[topIdx]; 

     int bottomIdx = x + 3; 
     unsigned char bottom_value; 
     if (ROWS <= bottomIdx) 
      bottom_value = 0; 
     else 
      bottom_value = SharedMemory[bottomIdx]; 

     *GET_UCHAR_PTR_3D(src, y, x, dispCnt) += bottom_value + top_value; 
    } 
} 

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

Мои извинения, я забыл добавить

#define GET_UCHAR_PTR_3D(pptr, x, y, d) \ 
(unsigned char*)((char*)(pptr).ptr + (sizeof(unsigned char)* x) + ((pptr).pitch * y) + ((pptr).pitch * (pptr).ysize * d)) 

Answer 1

Рассмотрим глобальную память доступа и сливающихся поведения между этими двумя случаями. Не имеет значения, рассматриваются ли операция загрузки:

unsigned char my_val = *GET_UCHAR_PTR_3D(src, y, x, dispCnt); 

или операции сохранения:

*GET_UCHAR_PTR_3D(src, y, x, dispCnt) += bottom_value + top_value; 

Давайте распаковать ваш макрос и подставят в фактических значениях x и y в каждом конкретном случае:

define GET_UCHAR_PTR_3D(pptr, x, y, d) \ 
(unsigned char*)((char*)(pptr).ptr + (sizeof(unsigned char)* x) + ((pptr).pitch * y) + ((pptr).pitch * (pptr).ysize * d)) 

мы имеем:

(a pointer) + (1*x) + (pitch*y) + offset 

Теперь, если х = threadIdx.x и у = blockIdx.x мы имеем:

(a pointer) + (1*threadIdx.x) + (pitch*blockIdx.x) + offset 

который становится:

(a pointer) + (some offset) + threadIdx.x 

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

Теперь, что произойдет, если x = blockIdx.x и y = threadIdx.x?мы имеем:

(a pointer) + (1*blockIdx.x) + (pitch*threadIdx.x) + offset 

, которая становится:

(a pointer) + (some offset) + (pitch*threadIdx.x) 

Это означает, что соседние нити в перекоса являются не чтение смежных местах в памяти, но вместо этого читают места, которые отделены друг от друга значением pitch , Это будет не объединиться, и переведет на многие более глобальные запросы, чтобы удовлетворить активность warp. Это «плохой случай».

GPU, такие как «горизонтальный» доступ к памяти в деформации. Им не нравится «вертикальный» доступ к памяти в пределах основы. Это приведет к очень существенной разнице в производительности между этими двумя случаями. Нередко можно увидеть разницу в 10x perf между этими двумя случаями, и теоретически это может быть так же высоко, как разница 32x perf.

Если вы хотите получить больше информации об оптимизации для совместного доступа к глобальной памяти, попробуйте this презентацию, особенно слайды 30-48.