CUDA 调试无效内核映像错误答案

【问题标题】：CUDA debug invalid kernel image errorCUDA 调试无效内核映像错误
【发布时间】：2014-03-25 20:35:16
【问题描述】：

我编写了以下 CUDA 内核并尝试将其加载到模块中：

#include <stdio.h>

extern "C"   // ensure function name to be exactly "vadd"
{
    __global__ void vadd(const float *a, const float *b, float *c)
    {
        int i = threadIdx.x + blockIdx.x * blockDim.x;
        printf("Thread id %d\n", i);
        c[i] = a[i] + b[i];
    }
}

我使用以下命令将其编译为 ptx 代码：

nvcc -ptx -arch=sm_20 vadd.cu

当尝试使用cuModuleLoad 将此文件加载到模块中时，我收到 CUDA 200 错误（无效的内核映像）。如何找出内核映像有什么问题？我试过ptxas，但据此生成的ptx代码是没问题的。

编辑：这是我用来加载模块的代码：

#include "cuda.h"
#include <cassert>
#include <dlfcn.h>
#include <stdio.h>

void check(CUresult err) {
  if (err != CUDA_SUCCESS) {
    printf("Error %i\n", err);
  }
  assert(err == CUDA_SUCCESS);
}

int main(int argc, char **argv) {
    void *cuda = dlopen("libcuda.so", RTLD_NOW | RTLD_DEEPBIND | RTLD_GLOBAL);
    assert(cuda != NULL);

    printf("cuInit\n");
    CUresult (*Init)() = (CUresult (*)()) dlsym(cuda, "cuInit");
    check(Init());

    printf("cuDeviceGet\n");
    CUresult (*DeviceGet)(CUdevice *, int) = (CUresult (*)(CUdevice *, int)) dlsym(cuda, "cuDeviceGet");
    CUdevice device;
    check(DeviceGet(&device, 0));

    printf("cuCtxCreate\n");
    CUresult (*CtxCreate)(CUcontext * , unsigned int, CUdevice) = (CUresult (*)(CUcontext * , unsigned int, CUdevice)) dlsym(cuda, "cuCtxCreate");
    CUcontext context;
    check(CtxCreate(&context, 0, device));

    printf("cuModuleLoad\n");
    CUresult (*ModuleLoad)(CUmodule *, const char*) = (CUresult (*)(CUmodule *, const char*)) dlsym(cuda, "cuModuleLoad");
    CUmodule mod;
    check(ModuleLoad(&mod, "vadd.ptx"));

    return 0;
}

【问题讨论】：

顺便说一句：根据文档，cuModuleLoad 永远不会导致CUDA_ERROR_INVALID_IMAGE。您是否还有其他调用通常有效的情况（例如，一个空内核，没有 printf 语句，针对 sm_10 左右编译）？
@Marco13 你确定：CUDA_ERROR_INVALID_IMAGE during cuModuleLoad?
对不起，我只是提到文档中没有提到它（docs.nvidia.com/cuda/cuda-driver-api/…），但它可能无论如何都会在实践中发生
您能否显示正在尝试加载和 JIT PTX 的主机代码。根据我的经验，这种失败通常意味着您正在尝试加载一个不是有效 PTX 或 CUBIN 图像的文件

标签： cuda

【解决方案1】：

这与Why cuCtxCreate creates old context? 相关：您正在直接使用cuCtxCreate，这会为您提供与您使用printf 不兼容的旧API 上下文（v3.1）。您可以使用cuCtxGetApiVersion 查看 API 版本。如果您切换到 cuCtxCreate_v2（通常通过 cuda.h 中的一些 #define 使用），您将获得更新的 API 上下文。

为了发现这种差异，我使用LD_DEBUG=symbols 运行了您的示例，并将其与直接使用 CUDA API 进行了比较（因为它可以正确运行您的示例 PTX）。比较符号分辨率，最大的区别在于对 cuCtxCreate 的调用：

cuCtxCreate(...)
    symbol=cuCtxCreate_v2;  lookup in file=./test [0]
    symbol=cuCtxCreate_v2;  lookup in file=/usr/lib/x86_64-linux-gnu/libcuda.so.1 [0]

...在您的原始代码中，使用dlsym(..., "cuCtxCreate") 直接映射到cuCtxCreate。

【讨论】：