PyCUDA无法在NVIDIA源代码中找到函数或抛出'可能没有extern“C”Linkage'错误

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

我还尝试在extern "C" { [...] }中包装modStr的内容，并将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加载以从模块中获取已知的受损名称。