我正在阅读杰森·桑德斯（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结构不包含复制构造函数，这是否会创建无效的free？我们按值传递锁，该锁仅按值复制互斥体指针。当我们退出内核时，析构函数在此互斥指针上调用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 Runtime API手册的3.9节说，该调用应该失败并返回cudaErrorInvalidDevicePointer。
尽管如此，那仍然是糟糕的编程实践，并且应该以不同的方式编写IMO的结构。
本文收集自互联网，转载请注明来源。
如有侵权，请联系[email protected] 删除。
我来说两句