Почему использование ключевого слова volatile для разделяемой памяти невозможно, когда атомарные операции выполняются в разделяемой памяти?

У меня есть фрагмент кода CUDA, в котором потоки выполняют атомарные операции с общей памятью. Я думал, так как результат атомарной операции будет виден другим потокам блока в любом случае мгновенно, возможно, было бы хорошо поручить компилятору иметь общую память volatile,
Итак, я изменился

__global__ void CoalescedAtomicOnSharedMem(int* data, uint nElem)
{
    __shared__ int smem_data[BLOCK_SIZE];
    uint tid = (blockIdx.x * blockDim.x) + threadIdx.x;
    for ( uint i = tid; i < nElem; i += blockDim.x*gridDim.x){
        atomicAdd( smem_data+threadIdx.x, 6);
    }
}

в

__global__ void volShared_CoalescedAtomicOnSharedMem(int* data, uint nElem)
{
    volatile __shared__ int smem_data[BLOCK_SIZE];
    uint tid = (blockIdx.x * blockDim.x) + threadIdx.x;
    for ( uint i = tid; i < nElem; i += blockDim.x*gridDim.x){
        atomicAdd( smem_data+threadIdx.x, 6);
    }
}

Ниже приведена ошибка времени компиляции с изменением выше:

error: no instance of overloaded function "atomicAdd" matches the argument list
        argument types are: (volatile int *, int)

Почему не volatile адрес поддерживается в качестве аргумента для атомарных операций? Это потому, что компилятор уже обрабатывает разделяемую память как изменчивую, как только он определяет, что над ней будут атомарные операции?

Источник

2 ответа

Решение

Определение volatile квалификатор приведен в руководстве по программированию. Он указывает компилятору всегда генерировать чтение или запись для этого доступа и никогда не "оптимизировать" его в регистр или какую-либо другую оптимизацию.

Поскольку атомарные операции гарантированно воздействуют на фактические области памяти (совместно используемой или глобальной), комбинация этих двух типов не требуется. Поэтому версии атомарных функций созданы для прототипа volatile классификатор не предоставляется.

Если у вас есть область памяти, которая уже объявлена ​​как volatileпросто приведите его к соответствующемуvolatile введите, когда вы передадите адрес вашей атомарной функции. Поведение будет таким, как ожидалось.( Пример)

Поэтому атомные операции могут работать в местах, указанных как volatile с этой оговоркой.

Тот простой факт, что вы получили доступ к определенному местоположению, используя атомику где-то в вашем коде, не означает, что компилятор будет обрабатывать каждый доступ в другом месте как неявно volatile, Если тебе надо volatile Поведение в другом месте, объявите это явно.

Источник

Предыдущий постер правильно определил проблему: не определена функция atomicAdd, которая принимает переменный параметр.

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

Why isn't a volatile address supported as an argument for atomic operations?

Атомарные операции не являются частью C/C++. В вашем случае они реализуются в библиотеке, которая, вероятно, реализована на ассемблере.

Is it because compiler already treats the shared memory as volatile as soon as it identifies there's going to be atomic operations on it?

Нет, именно так автор библиотеки определил интерфейс функции

Источник
Другие вопросы по тегам