atomicCAS для реализации bool

Question

Я пытаюсь выяснить, есть ли ошибка в ответе (теперь удаленном) о реализации Cuda-like atomicCAS для bool s. Код из ответа (переформатирован):

static __inline__ __device__ bool atomicCAS(bool *address, bool compare, bool val)
{
    unsigned long long addr = (unsigned long long)address;
    unsigned pos = addr & 3;  // byte position within the int
    int *int_addr = (int *)(addr - pos);  // int-aligned address
    int old = *int_addr, assumed, ival;

    do
    {
        assumed = old;
        if(val)
            ival = old | (1 << (8 * pos));
        else
            ival = old & (~((0xFFU) << (8 * pos)));
        old = atomicCAS(int_addr, assumed, ival);
    } while(assumed != old);

    return (bool)(old & ((0xFFU) << (8 * pos)));
}

Согласно в документации , atomicCAS должен установить *address на (*address == compare ? val : *address), но в реализации выше compare аргумент никогда не используется!

Код, который я использую для воспроизведения ошибки:

#include <cstdio>

// atomicCAS definition here

__device__ bool b;


__global__ void kernel()
{
    b = false;
    atomicCAS(&b, true, true); // `(b == true ? true : b)`, where b is false equals to false
    printf("%d\n", b); // b is false => expected output is 0
}


int main()
{
    kernel<<<1, 1>>>();
    cudaDeviceSynchronize();
}

Ожидаемый результат - 0, но фактический результат - 1.

У меня есть предложение о том, как это исправить, но я не уверен на 100%, что это потокобезопасно:

static __inline__ __device__ bool atomicCAS(bool *address, bool compare, bool val)
{
    unsigned long long addr = (unsigned long long)address;
    unsigned pos = addr & 3;  // byte position within the int
    int *int_addr = (int *)(addr - pos);  // int-aligned address
    int old = *int_addr, assumed, ival;

    do
    {
        if(*address != compare) // If we expected that bool to be different, then
            break; // stop trying to update it and just return it's current value

        assumed = old;
        if(val)
            ival = old | (1 << (8 * pos));
        else
            ival = old & (~((0xFFU) << (8 * pos)));
        old = atomicCAS(int_addr, assumed, ival);
    } while(assumed != old);

    return (bool)(old & ((0xFFU) << (8 * pos)));
}

Мои вопросы:

1. Есть ли ошибка в первом примере кода из ответ? Если есть,
2. Исправляет ли последний пример кода это поточно-поточно?

Answer 1

Большое спасибо @RobertCrovella; первый пример кода содержит ошибку, второй - исправляет ее, но не является потокобезопасным (подробности см. в комментариях к вопросу). Поточно-ориентированное исправление:

static __inline__ __device__ bool atomicCAS(bool *address, bool compare, bool val)
{
    unsigned long long addr = (unsigned long long)address;
    unsigned pos = addr & 3;  // byte position within the int
    int *int_addr = (int *)(addr - pos);  // int-aligned address
    int old = *int_addr, assumed, ival;

    bool current_value;

    do
    {
        current_value = (bool)(old & ((0xFFU) << (8 * pos)));

        if(current_value != compare) // If we expected that bool to be different, then
            break; // stop trying to update it and just return it's current value

        assumed = old;
        if(val)
            ival = old | (1 << (8 * pos));
        else
            ival = old & (~((0xFFU) << (8 * pos)));
        old = atomicCAS(int_addr, assumed, ival);
    } while(assumed != old);

    return current_value;
}