CUDA Недопустимый доступ к памяти с возможной «недостаточной» общей памятью

Question

У меня есть простое ядро ​​CUDA, которое может выполнять векторное накопление путем базового сокращения. Я масштабирую его, чтобы иметь возможность обрабатывать большие данные, разбивая его на несколько блоков. Тем не менее, мое предположение о назначении соответствующего количества разделяемой памяти, которое будет использоваться ядром, терпит неудачу с незаконным доступом к памяти. Он уходит, когда я увеличиваю этот предел, но я хочу знать, почему. Вот код, который я говорю:

CORE KERNEL:

__global__ static 
    void vec_add(int *buffer, 
       int numElem, // The actual number of elements 
       int numIntermediates) // The next power of two of numElem 
    { 
     extern __shared__ unsigned int interim[]; 

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

     // Copy global intermediate values into shared memory. 
     interim[threadIdx.x] = 
      (index < numElem) ? buffer[index] : 0; 

     __syncthreads(); 

     // numIntermediates2 *must* be a power of two! 
     for (unsigned int s = numIntermediates/2; s > 0; s >>= 1) { 
      if (threadIdx.x < s) { 
       interim[threadIdx.x] += interim[threadIdx.x + s]; 
      } 
      __syncthreads(); 
     } 

     if (threadIdx.x == 0) { 
      buffer[blockIdx.x] = interim[0]; 
     } 
    } 

И это звонящий:

void accumulate (int* buffer, int numElem) 
{ 
    unsigned int numReductionThreads = 
     nextPowerOfTwo(numElem); // A routine to return the next higher power of 2. 

    const unsigned int maxThreadsPerBlock = 1024; // deviceProp.maxThreadsPerBlock 

    unsigned int numThreadsPerBlock, numReductionBlocks, reductionBlockSharedDataSize; 

    while (numReductionThreads > 1) { 

     numThreadsPerBlock = numReductionThreads < maxThreadsPerBlock ?   
      numReductionThreads : maxThreadsPerBlock; 

     numReductionBlocks = (numReductionThreads + numThreadsPerBlock - 1)/numThreadsPerBlock; 

     reductionBlockSharedDataSize = numThreadsPerBlock * sizeof(unsigned int); 

     vec_add <<< numReductionBlocks, numThreadsPerBlock, reductionBlockSharedDataSize >>> 
      (buffer, numElem, numReductionThreads); 

     numReductionThreads = nextPowerOfTwo(numReductionBlocks); 
    } 

} 

Я попробовал этот код с набором образцов 1152 элементов на мой графический процессор со следующей конфигурацией: Тип: Quadro 600 MaxThreadsPerBlock: 1024 MaxSharedMemory: 48KB

ВЫВОД:

Loop 1: numElem = 1152, numReductionThreads = 2048, numReductionBlocks = 2, numThreadsPerBlock = 1024, reductionBlockSharedDataSize = 4096 
Loop 2: numElem = 1152, numReductionThreads = 2, numReductionBlocks = 1, numThreadsPerBlock = 2, reductionBlockSharedDataSize = 8 
CUDA Error 77: an illegal memory access was encountered 

Подозревая, что моя «промежуточная» разделяемая память была причиной незаконного доступа к памяти, я произвольно увеличил общую память в два раза в следующей строке:

reductionBlockSharedDataSize = 2 * numThreadsPerBlock * sizeof(unsigned int); 

И мои ядра начали работаю нормально!

Я не понимаю, почему я должен был предоставить эту дополнительную общую память, чтобы моя проблема исчезла (временно).

В качестве дополнительного эксперимента для проверки этого магического номера я запустил свой код с гораздо большим набором данных с 6912 точками. На этот раз даже 2X или 4X не помогли мне.

Loop 1: numElem = 6912, numReductionThreads = 8192, numReductionBlocks = 8, numThreadsPerBlock = 1024, reductionBlockSharedDataSize = 16384 

Loop 2: numElem = 6912, numReductionThreads = 8, numReductionBlocks = 1, numThreadsPerBlock = 8, reductionBlockSharedDataSize = 128 
CUDA Error 77: an illegal memory access was encountered 

Но проблема снова исчезла, когда я увеличил размер общей памяти на 8X.

Конечно, я не могу произвольно выбирать этот коэффициент масштабирования для больших и больших наборов данных, потому что в скором времени у меня закончится ограничение на общую память 48 КБ. Поэтому я хочу знать законный способ решения проблемы.

Answer 1

Благодаря @havogt для указания доступа вне индекса. Проблема заключалась в том, что я использовал неправильный аргумент как numIntermediates для метода vec_add. Предполагалось, что ядро ​​будет работать с точно таким же количеством точек данных, сколько количество потоков, которые должны были быть 1024 все время. Я установил его с помощью numThreadsPerBlock в качестве аргумента:

vec_add <<< numReductionBlocks, numThreadsPerBlock, reductionBlockSharedDataSize >>> 
     (buffer, numElem, numThreadsPerBlock);