Cuda scan - разные результаты в режимах Debug и Release при использовании общей памяти

Question

Я пытаюсь написать сегментированное сканирование в cuda, где длина сегмента равна длине деформации (32). Вот мое ядро:

__global__ void kernel(int totalSize, unsigned short* result)
{
    __shared__ unsigned short s_data[1024];

    const unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int intraWarpThreadId = threadIdx.x & 31;

    if (tid >= totalSize)
        return;

    s_data[threadIdx.x] = result[tid];
    __syncthreads();

    IntraWarpScan(s_data, threadIdx.x, intraWarpThreadId);
    __syncthreads();

    result[tid] = s_data[threadIdx.x];
}

__device__ void IntraWarpScan(unsigned short* s_data, unsigned int intraBlockThreadId, unsigned int& intraWarpThreadId)
{
    if (intraWarpThreadId >= 1)
        s_data[intraBlockThreadId] += s_data[intraBlockThreadId - 1];

    if (intraWarpThreadId >= 2)
        s_data[intraBlockThreadId] += s_data[intraBlockThreadId - 2];

    if (intraWarpThreadId >= 4)
        s_data[intraBlockThreadId] += s_data[intraBlockThreadId - 4];

    if (intraWarpThreadId >= 8)
        s_data[intraBlockThreadId] += s_data[intraBlockThreadId - 8];

    if (intraWarpThreadId >= 16)
        s_data[intraBlockThreadId] += s_data[intraBlockThreadId - 16];
}

Полагаю, у меня есть некоторые условия гонки в общей памяти, но я не могу понять, почему они случаются. Так как каждый сегмент сканируется внутри деформации, мне не нужна синхронизация внутри процедуры IntraWarpScan, верно? Но без синхронизации после каждой инструкции if в IntraWarpScan я получаю неправильные результаты в сборке выпуска. В Debug я получаю правильные результаты.

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

__global__ void kernel(int totalSize, unsigned short* result)
{
    const unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int intraWarpThreadId = threadIdx.x & 31;

    if (tid >= totalSize)
        return;

    IntraWarpScan(result, tid, intraWarpThreadId);
    __syncthreads();
}

__device__ void IntraWarpScan(unsigned short* s_data, unsigned int intraBlockThreadId, unsigned int& intraWarpThreadId)
{
    if (intraWarpThreadId >= 1)
        s_data[intraBlockThreadId] += s_data[intraBlockThreadId - 1];

    if (intraWarpThreadId >= 2)
        s_data[intraBlockThreadId] += s_data[intraBlockThreadId - 2];

    if (intraWarpThreadId >= 4)
        s_data[intraBlockThreadId] += s_data[intraBlockThreadId - 4];

    if (intraWarpThreadId >= 8)
        s_data[intraBlockThreadId] += s_data[intraBlockThreadId - 8];

    if (intraWarpThreadId >= 16)
        s_data[intraBlockThreadId] += s_data[intraBlockThreadId - 16];
}

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

Answer 1

Спасибо talonmies за указание на проблему. Решение состоит в том, чтобы добавить volatile в декларации совместно используемой памяти. В основном в режиме Release компилятору разрешено хранить значения в регистрах и не хранить промежуточные хранилища в общей памяти. В случае, когда потоки обращаются к расположениям общей памяти, измененным другими потоками, память должна быть объявлена ​​как энергозависимая, чтобы отключить такие оптимизации.

Очень хорошее и подробное объяснение можно найти здесь: Когда использовать volatile с общей памятью CUDA