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

Use of locks in Cuda by Example dot product

本文关键字：Cuda      更新时间：2023-10-16

我正在阅读Jason Sanders和Edward Kandrot的"Cuda by Example"一书，并对他们使用锁来计算两个数组的点积有疑问（可在链接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 错误，我认为问题是没有复制构造函数。

是的，你是对的，将进行额外的 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。

相关文章：