CUDA синхронизация и чтение глобальной памяти

Question

У меня есть что-то вроде этого:

__global__ void globFunction(int *arr, int N) {
    int idx = blockIdx.x* blockDim.x+ threadIdx.x;
    // calculating and Writing results to arr ...
    __syncthreads();
    // reading values of another threads(ex i+1)
    int val = arr[idx+1]; // IT IS GIVING OLD VALUE
}


int main() {
    // declare array, alloc memory, copy memory, etc.
    globFunction<<< 4000, 256>>>(arr, N); 
    // do something ...
    return 0;
}

Почему я получаю старое значение, когда читаю arr[idx+1]? Я позвонил __syncthreads, поэтому я ожидаю увидеть обновленное значение. Что я сделал не так? Я читаю кеш или как?

Answer 1

Использование функции __syncthreads() синхронизирует только потоки в текущем блоке.В этом случае это будет 256 потоков на блок, который вы создали при запуске ядра.Таким образом, в данном массиве для каждого значения индекса, которое пересекается с другим блоком потоков, вы в конечном итоге будете читать значение из глобальной памяти, которое не синхронизировано с потоками в текущем блоке.

Единственное, что вы можете сделать, чтобы обойти эту проблему, - создать общее локальное хранилище потоков, используя директиву __shared__ CUDA, которая позволяет потокам в ваших блоках обмениваться информацией между собой, но предотвращает доступ потоков из других блоков к памяти, выделенной для текущего блока,Как только ваши вычисления в блоке завершены (и вы можете использовать __syncthreads() для этой задачи), вы можете скопировать обратно в глобально доступную память значения в общем хранилище уровня блока.

Ваше ядро ​​можетвыглядеть примерно так:

__global__ void globFunction(int *arr, int N) 
{
    __shared__ int local_array[THREADS_PER_BLOCK];  //local block memory cache           
    int idx = blockIdx.x* blockDim.x+ threadIdx.x;

    //...calculate results
    local_array[threadIdx.x] = results;

    //synchronize the local threads writing to the local memory cache
    __syncthreads();

    // read the results of another thread in the current thread
    int val = local_array[(threadIdx.x + 1) % THREADS_PER_BLOCK];

    //write back the value to global memory
    arr[idx] = val;        
}

Если вам нужно синхронизировать потоки между блоками, вам следует искать другой способ решения вашей проблемы, поскольку модель программирования CUDA работает наиболее эффективно, когда проблему можно разбить на блоки, а синхронизация потоков должна происходить только внутри блока.