在阅读了 CUDA 9 中的协作组之后,我一直在尝试在网格级别进行同步。
我使用的是 Visual Studio 2017、GTX 1060 和 CUDA 9.1。
我将代码更改如下:
__global__ void ExplicitKernel_American(/* ... */) {
int i = threadIdx.x + blockDim.x * blockIdx.x;
auto grid = cooperative_groups::this_grid();
if (i < sizeS) {
//...
for (int j = 1; j < sizeT; ++j) {
// ...
grid.sync(); // __syncthreads();
}
}
}
并且,正如文档中所述,我这样调用我的内核:
void* Explicit_Args[] = { &PDE_Grid, /* ... */, &sizeS, &sizeT };
cudaLaunchCooperativeKernel(
(void*)ExplicitKernel_American,
dim3((sizeS + TPB - 1) / TPB),
dim3(TPB),
Explicit_Args
); // TPB being 256...
不幸的是,当我在内核中添加“网格”部分时,我就遇到了链接错误。
error LNK2001: unresolved external symbol __fatbinwrap_38_cuda_device_runtime_compute_70_cpp1_ii_8b1a5d37
fatal error LNK1120: 1 unresolved externals
我已经设置了 -rdc=true 和 sm_61 但找不到它不起作用的原因......有什么想法吗?
非常感谢!
使用协作内核启动(协作网格 - CG)需要 Pascal 或 Volta GPU,并且需要在 TCC 模式下运行的 Linux 或 Windows 设备。 如果您测试设备属性结构中的
deviceProp.cooperativeLaunch
属性,我想您会发现您的 GPU 在 WDDM 模式下运行时不支持它。
在尝试使用协作网格启动之前,最好在代码中测试此属性。
但是,您询问的问题是编译/链接问题。 为此,我的建议是研究 CG(合作网格)示例代码,例如
6_Advanced/reductionMultiBlockCG
。对于网格同步,绝对需要设置-rdc=true
(即启用可重定位设备代码链接)。根据您设置 -rdc=true
的方式,它可能无法正确应用于您的项目。 这里概述了正确的方法
这里最接近的问题似乎是您没有正确链接到设备运行时库,例如
-lcudadevrt
编辑:最近,WDDM GPU 可能支持。 最佳实践始终是根据提到的属性/属性来决定支持。