通过示例点积在 Cuda 中使用锁

Question

我正在阅读"Cuda by Example" by Jason Sanders and Edward Kandrot 这本书，并且对他们使用锁来计算两个数组的点积有疑问（在链接 pdf 的第 254-258 页上找到）。他们定义了他们的 lock.h 头文件：

#ifndef __LOCK_H__
#define __LOCK_H__


struct Lock{
    int *mutex;
    Lock(){
        int state = 0;
        cudaMalloc((void**)&mutex, sizeof(int));
        cudaMemcpy(mutex, &state, sizeof(int), cudaMemcpyHostToDevice);
    }

    ~Lock(){
        cudaFree(mutex);
    }

    __device__ void lock(){
        while (atomicCAS(mutex, 0 ,1) != 0);
    }

    __device__ void unlock(){
        atomicExch(mutex, 0);
    }
};


#endif

然后他们将他们的点积内核称为：

int main()
{
    // bunch of code, initialization etc

    Lock lock;
    dot<<<blocksPerGrid,threadsPerBlock>>>(lock, dev_a, dev_b, dev_c);

    // more code, frees, etc
} 

点内核声明为：

__global__ void dot(Lock lock, float *a, float *b, float *c);

这不会因为 Lock 结构不包含复制构造函数而创建一个无效的空闲吗？我们按值传递锁，它只是按值复制互斥指针。当我们退出内核时，析构函数在这个互斥指针上调用 cudaFree。当我们退出主函数时，我想再次调用析构函数，但现在互斥锁已经被释放了！我问是因为在不同的较大代码中，当使用相同的想法时，我会遇到 cudeErrorInvalidDevicePointer 错误，我认为问题在于没有复制构造函数。

Answer 1

是的，您是对的，将进行额外的 dtor 呼叫 - 您也可以自己检查...

这里有一个简单的例子来验证：

#include <iostream>
#include <cuda_runtime.h>

struct A {
    A()  { std::cout << "ctor for " << this << std::endl << std::flush; }
    ~A() { std::cout << "dtor for " << this << std::endl << std::flush; }
};

__global__ void foo(A device_a) { }

int main(void) {
    A host_a;
    foo<<<1,1>>>(host_a);
    cudaDeviceReset();
    return 0;
}

输出是：

ctor for 0x7ffe584c85bf
dtor for 0x7ffe584c85df
dtor for 0x7ffe584c85bf

因此，您将获得默认的复制 ctor；以及两个副本的构造和销毁发生在主机上（这些是主机专用功能，设备没有<iostream>）。我实际上觉得这有点奇怪，但我想这是 CUDA 的“C-ish”起源的一部分，将内核参数视为 POD。看到 ctor 和 dtor 如何拥有 __device__ 和 __host__ 限定符，这尤其令人惊讶。

但是，在该示例中对 cudaFree() 的额外调用应该不是问题（手指交叉）：CUDA 运行时 API 手册第 3.9 节说调用应该失败并返回 cudaErrorInvalidDevicePointer。

不过，这是一种糟糕的编程习惯，而且 IMO 应该以不同的方式编写结构。