Cuda разделяет память вне границ, когда используется только один блок или слишком мало потоков

Question

Я попытался реализовать уменьшение векторной суммы самостоятельно, используя CUDA, и столкнулся с ошибкой, которую мог исправить, но не понял, в чем была проблема.

Я реализовал ядро ​​ниже, которое во многом похоже на то, что использовалось в примерах NVIDIA.

__global__ 
void reduce0(int *input, int *output)
{
    extern __shared__ int s_data[];

    int tid = threadIdx.x;
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    s_data[tid] = input[i];
    __syncthreads();

    for( int s=1; s < blockDim.x; s *= 2) {
        if((tid % 2*s) == 0) {
            s_data[tid] += s_data[tid + s];
        }

        __syncthreads();
    }

    if(tid == 0) {
        output[blockIdx.x] = s_data[0];
    }
}

Кроме того, я рассчитал общее пространство памяти, как показано ниже на стороне хоста

int sharedMemSize = numberOfValues * sizeof(int);

Если используется более 1 блока потоков, код работает нормально. Использование только 1 блока заканчивается в индексе за пределами указанной выше ошибки. Ища мою ошибку, сравнивая мой хост-код с одним из примеров, я нашел следующую строку:

int smemSize = (threads <= 32) ? 2 * threads * sizeof(T) : threads * sizeof(T);

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

1. блок, произвольное количество потоков => сбой кода
2. >2 блока, произвольное количество потоков => код работает нормально
3. 1 блок, произвольное количество потоков, размер разделяемой памяти 2*#threads => код работает нормально

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

ОБНОВЛЕНИЕ: код хоста, вызывающий ядро ​​в соответствии с запросом

int numberOfValues = 1024 ;
int numberOfThreadsPerBlock = 32;
int numberOfBlocks = numberOfValues / numberOfThreadsPerBlock;

int memSize = sizeof(int) * numberOfValues;

int *values = (int *) malloc(memSize);
int *result = (int *) malloc(memSize);

int *values_device, *result_device;
cudaMalloc((void **) &values_device, memSize);
cudaMalloc((void **) &result_device, memSize);

for(int i=0; i < numberOfValues ; i++) {
    values[i] = i+1;
}

cudaMemcpy(values_device, values, memSize, cudaMemcpyHostToDevice);

dim3 dimGrid(numberOfBlocks,1);
dim3 dimBlock(numberOfThreadsPerBlock,1);
int sharedMemSize = numberOfThreadsPerBlock * sizeof(int);

reduce0 <<< dimGrid, dimBlock, sharedMemSize >>>(values_device, result_device);

if (cudaSuccess != cudaGetLastError())
        printf( "Error!\n" );

cudaMemcpy(result, result_device, memSize, cudaMemcpyDeviceToHost);

Answer 1

Может ли ваша проблема быть порядок приоритетов по модулю и умножению.tid % 2*s равно (tid % s)*2 но ты хочешь tid % (s*2)

Причина, почему вам нужно использовать int smemSize = (threads <= 32) ? 2 * threads * sizeof(T) : threads * sizeof(T) для небольшого количества потоков это связано с заграничной индексацией. Один из примеров, когда это происходит, - запуск 29 потоков. когда tid=28 а также s=2 ветка будет взята из-за 28 % (2*2) == 0 и вы будете индексировать в s_data[28+2] но вы выделили общую память только для 29 потоков.