PyCUDA 要么无法在 NVIDIA 源代码中找到函数，要么抛出“可能没有外部“C”链接”错误

Question

我正在尝试使用（并从中学习）Mark Harris's optimized reduction kernel，将他的源代码复制到一个简单的 pycuda 应用程序中（下面列出了我尝试的完整源代码）。

不幸的是，我遇到了以下两个错误之一。

1. cuda内核没有编译，抛出如下错误信息。

   kernel.cu(3): error: this declaration may not have extern "C" linkage
   

2. 如果我在编译内核的行中包含参数no_extern_c=True，则会引发以下错误：

   pycuda._driver.LogicError: cuModuleGetFunction failed: not found
   

我还尝试将 modStr 的内容包装在 extern "C" { [...] } 中，并将 no_extern_c 变量设置为 True 或 False，但没有任何成功。

问题似乎涉及template <unsigned int blockSize> 行，好像我将函数的主体注释掉它仍然会引发错误。但我对这个问题的理解还不够深入，无法对如何解决它有更多的想法。

任何意见/建议/帮助将不胜感激 -- 提前致谢！

from pylab import *

import pycuda.gpuarray as gpuarray
import pycuda.autoinit
import pycuda.driver as drv
from pycuda.compiler import SourceModule

modStr = """
template <unsigned int blockSize>
__global__ void reduce6(int *g_idata, int *g_odata, unsigned int n) {
    extern __shared__ int sdata[];
    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x*(blockSize*2) + tid;
    unsigned int gridSize = blockSize*2*gridDim.x;
    sdata[tid] = 0;
    while (i < n) { 
        sdata[tid] += g_idata[i] + g_idata[i+blockSize]; i += gridSize; 
    }
    __syncthreads();
    if (blockSize >= 512) { if (tid < 256) { sdata[tid] += sdata[tid + 256]; } __syncthreads(); }
    if (blockSize >= 256) { if (tid < 128) { sdata[tid] += sdata[tid + 128]; } __syncthreads(); }
    if (blockSize >= 128) { if (tid < 64) { sdata[tid] += sdata[tid + 64]; } __syncthreads(); }
    if (tid < 32) {
        if (blockSize >= 64) sdata[tid] += sdata[tid + 32];
        if (blockSize >= 32) sdata[tid] += sdata[tid + 16];
        if (blockSize >= 16) sdata[tid] += sdata[tid + 8];
        if (blockSize >= 8) sdata[tid] += sdata[tid + 4];
        if (blockSize >= 4) sdata[tid] += sdata[tid + 2];
        if (blockSize >= 2) sdata[tid] += sdata[tid + 1];
    }
    if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}
"""

mod = SourceModule(modStr,no_extern_c=True) 
# With no_extern_c = True, the error is :
# pycuda._driver.LogicError: cuModuleGetFunction failed: not found
# With no_extern_c = False, the error is :
# kernel.cu(3): error: this declaration may not have extern "C" linkage


cuda_reduce_fn = mod.get_function("reduce6")
iData = arange(32).astype(np.float32)
oData = zeros_like(iData)

cuda_reduce_fn(
    drv.In(iData),
    drv.Out(oData),
    np.int32(32),
    block=(32,1,1), grid=(1,1))

print(iData)
print(oData)

Answer 1

在 C++ 中使用带有 C 链接的模板函数是非法的，这就是您在第一种情况下得到错误的原因。

在第二种情况下，您会收到一个未找到的错误，因为您实际上并没有在我能看到的任何地方实例化模板，因此编译器不会发出任何输出。

当你添加一个实例时，你会得到同样的错误，因为设备的编译代码对象有一个mangled name。您将需要在 get_function 调用中使用损坏的名称。矛盾的是，当从源代码进行 JIT 编译时，您不知道损坏的名称，因为您需要查看编译器输出并且不知道先验（任何编译器消息、PTX、cubin 或目标文件都会为您提供损坏的名称)。

如果您想在 PyCUDA 中使用模板化内核，我建议您将它们编译为使用工具链进行 cubin，然后从 PyCUDA 中的 cubin 加载以从模块中获取已知的重整名称。