В какой памяти идет создание статического массива при объявлении его в __global__ или __device__?

В приведенном ниже примере идет подсчет гистограммы (count_by_key), т.е. подсчет в массиве количества каждого элемента размером в 1 байт. В kernel (histogram_gpu) создается временный статический массив и в основном цикле программы используется обычное НЕ атомарное инкрементирование элементов этого массива и все работает без data race condition.

1. Переменные объявленные в __global__ или __device__ создаются для каждого потока свои, а вот статические массивы, в частности (unsigned short int temp[256];) тоже создаются для каждого потока?
2. И создание статического массива в __global__ или __device__ создает его в какой памяти: глобальной, локальной, разделяемой или регистрах?
3. И есть какие-то варианты ускорить подсчет гистограммы?

  1. // speed 1 - 3 GB/sec on GTX460 SE
  2. #define BLOCKS_GPU 24
  3. #define THREADS_PER_BLOCK 128  
  4. #define ALL_THREADS (BLOCKS_GPU*THREADS_PER_BLOCK)
  5. #define UNROLL_ROW 4
  6.  
  7. __global__ void histogram_gpu(unsigned int *counts, unsigned char *src, unsigned int size) {
  8.                 unsigned int index_thread = threadIdx.x + blockIdx.x * blockDim.x;
  9.                 unsigned short int temp[256];
  10.  
  11.                 #pragma unroll 256
  12.                 for (unsigned int m = 0; m < 256; ++m)
  13.                         temp[m] = 0;
  14.  
  15.                 for (unsigned int k=0; k < size; k += ALL_THREADS*UNROLL_ROW) {
  16.                         const unsigned int offset = index_thread*UNROLL_ROW + k;
  17.                        
  18.                         #pragma unroll UNROLL_ROW
  19.                         for (unsigned int i = 0; i < UNROLL_ROW; ++i)
  20.                                 ++*(temp + src[offset + i]);  // даже без atomicAdd считает точно без race-condition
  21.                 }
  22.     __syncthreads();
  23.                                
  24.                         #pragma unroll 256
  25.                         for (unsigned int m = 0; m < 256; ++m)
  26.                                 atomicAdd(counts + m, temp[m]);
  27. }

Forums: 

temp - пытается создаться в

temp - пытается создаться в регистрах, а если их не хватает - создается в глобальной памяти, индивидуальный массив для каждого потока. В результате все будет безобразно медленно.

В примерах CUDA есть пример Histogram, посмотрите на него. На моей GTX480 он работает со скоростью 17Gb/sec

Да, пробовал histogram_256 из

Да, пробовал histogram_256 из примеров. У меня массив небольшой 50 МБ, на нем она работает со скоростью 2.4 GB/sec на GTX460 SE.
В то же время вот здесь http://www.fastvideo.ru/info/cuda/cuda-histogram.htm
говорят "Для тестового массива данных размером 143 МБайт скорость расчёта гистограммы 256-bin получается 74 ГБайт/с, а для гистограммы 64-bin мы получили производительность 110 ГБайт/с."

Кстати, http://www.cvg.ethz.ch/teaching/2011spring/gpgpu/cuda_memory.pdf
Нумерации страниц нет, искать по CUDA Type Qualifiers

Variable declaration Memory Scope Lifetime
int LocalVar; register thread thread
int LocalArray[10]; local thread thread
[__device__] __shared__ int SharedVar; shared block block
__device__ int GlobalVar; global grid application
[__device__] __constant__ int ConstantVar; constant grid application

Пишут, что статический массив даже не будет пытаться разместиться в регистрах и сразу создастся в локальной памяти (кэшированной глобальной)?

Лично я бы пробовал бы

Лично я бы пробовал бы обойтись без atomic, одним из двух вариантов
а) использование регистров. Да, в Fermi их всего 63 на thread, но можно попробовать использовать, скажем, 32 на поток и один и тот же массив данных - грузить в shared mem и одно и то же значение анализировать 8 раз (сдвиг и битовый AND, величина сдвига зависит от номера потока).
б) Наоборот, использование shared для персональной гистограммы каждого потока, притом таким образом, чтобы гарантированно не было бы конфликта банков.
Есть 48k shared, нужно 32 (по 1 килобайту на поток), на SM будет один warp - ну и хрен с ним.
Латентность прятать разворотом цикла эдак в 16.

Пока писал - решил, что путь б) явно более многообещающий.

А у меня первый пример уже

А у меня первый пример уже без Atomic :)
И кстати в экзамплах hostogram256 есть дефайн для двух реализаций с атомиками и без, без работает заметно быстрее :)

Вот пробовал реализовать тот что у вас вариант Б, уж не знаю есть ли там конфликты банок, но работает у меня слабовато 0.6 GB/sec:

  1. #define SLICES 48
  2. #define BLOCKS_GPU 6
  3. #define THREADS_PER_BLOCK SLICES        // 1 thread - 1 slice (not need atomic)
  4. #define ALL_THREADS (BLOCKS_GPU*THREADS_PER_BLOCK)
  5. #define UNROLL_ROW 32
  6.  
  7. __global__ void histogram_gpu(unsigned int *counts, unsigned char *src, unsigned int size) {
  8.  
  9.     __shared__ unsigned int temp[SLICES*256];           // 4 bytes*256*48 = 48 KB
  10.                 const unsigned int offset_slice = (threadIdx.x % SLICES)*256;
  11.                 const unsigned int index_thread = threadIdx.x + blockIdx.x * blockDim.x;
  12.  
  13.                 #pragma unroll 256
  14.                 for (unsigned int m = 0; m < 256; ++m)
  15.                         temp[offset_slice + m] = 0;
  16.                 __syncthreads();
  17.  
  18.                 unsigned int *const ptr_temp_for_slice = temp + offset_slice;
  19.  
  20.                 for (unsigned int k=0; k < size; k += ALL_THREADS*UNROLL_ROW) {
  21.                         const unsigned int offset = index_thread*UNROLL_ROW + k;
  22.                        
  23.                         #pragma unroll UNROLL_ROW
  24.                         for (unsigned int i = 0; i < UNROLL_ROW; ++i)
  25.                                 //atomicAdd(ptr_temp_for_slice + src[offset + i], 1);
  26.                                 ++*(ptr_temp_for_slice + src[offset + i]);              // 1 thread - 1 slice (not need atomic)
  27.                 }
  28.     __syncthreads();
  29.  
  30.                 if (threadIdx.x < SLICES) {
  31.                         #pragma unroll 256
  32.                         for (unsigned int m = 0; m < 256; ++m)
  33.                                 atomicAdd(counts + m, temp[offset_slice + m]);
  34.                 }
  35. }

Нет. Самый первый пример в

Нет.
Самый первый пример в локальной памяти и работает от 1 до 3 GB/sec (в зависимости от данных).
А второй пример в разделяемой памяти и работает 0.6 GB/sec.
Оба без atomic.