pointers - 在 CUDA 上原子交换内存指针

Question

我在内存中有两个指针，我想以原子方式交换它，但 CUDA 中的原子操作仅支持 int 类型。有没有办法做下面的交换？

classA* a1 = malloc(...);
classA* a2 = malloc(...);
atomicSwap(a1,a2);

Answer 1

在编写设备端代码时...

虽然 CUDA 提供原子，但它们不能同时覆盖多个（可能是远程的）内存位置。

要执行此交换，您需要使用互斥锁之类的东西“保护”对这两个值的访问，并让想要向它们写入值的任何人在关键部分的持续时间内持有互斥锁（就像在 C++ 的主机中一样 -边std::lock_guard）。这可以使用 CUDA 的实际原子设施来完成，例如比较和交换，并且是这个问题的主题：

在 CUDA 中实现关键部分

@RobertCrovella 提到了对上述内容的警告：如果您可以使用一对 32 位偏移量而不是 64 位指针，那么如果您要将它们存储在 64 位对齐struct中，您可以在整个 struct 上使用 compare-and-exchange 来实现整个struct.

...但它真的是设备端代码吗？

您的代码实际上看起来不像可以在设备上运行的东西：内存分配通常（尽管并非总是）在您启动内核并进行实际工作之前从主机端完成。如果您可以确保这些更改仅发生在主机端（想想 CUDA 事件和回调），并且设备端代码不会受到它们的干扰 - 您可以使用普通的 vanilla C++ 工具进行并发编程（比如lock_guard我上文提到的）。

Answer 2

我设法获得了所需的行为，它不是原子交换，但仍然安全。上下文是一个在 CPU 和 GPU 上工作的单调链表：

template<typename T>
union readablePointer
{
    T* ptr;
    unsigned long long int address;
};

template<typename T>
struct LinkedList
{

    struct Node
    {
        T value;
        readablePointer<Node> previous;
    };

    Node start;
    Node end;

    int size;

    __host__ __device__ void initialize()
    {
        size = 0;

        start.previous.ptr = nullptr;
        end.previous.ptr = &start;
    }

    __host__ __device__ void push_back(T value)
    {
        Node* node = nullptr;
        malloc(&node, sizeof(Node));

        readablePointer<Node> nodePtr;
        nodePtr.ptr = node;

        nodePtr.ptr->value = value;

#ifdef __CUDA_ARCH__
        nodePtr.ptr->previous.address = atomicExch(&end.previous.address, nodePtr.address);
        atomicAdd(&size,1);
#else
        nodePtr.ptr->previous.address = end.previous.address;
        end.previous.address = nodePtr.address;
        size += 1;
#endif

    }

    __host__ __device__ T pop_back()
    {
        assert(end.previous.ptr != &start);

        readablePointer<Node> lastNodePtr;
        lastNodePtr.ptr = nullptr;

#ifdef __CUDA_ARCH__
        lastNodePtr.address = atomicExch(&end.previous.address,end.previous.ptr->previous.address);
        atomicSub(&size,1);
#else
        lastNodePtr.address = end.previous.address;
        end.previous.address = end.previous.ptr->previous.address;
        size -= 1;
#endif
        T toReturn = lastNodePtr.ptr->value;

        free(lastNodePtr.ptr);

        return toReturn;
    }

    __host__ __device__ void clear()
    {
        while(size > 0)
        {
            pop_back();
        }
    }
};