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: 48 КБ

ВЫХОД:

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

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

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

Answer 1

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

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