Cuda 错误 CUDA_ERROR_NO_BINARY_FOR_GPU答案

【问题标题】：Cuda error CUDA_ERROR_NO_BINARY_FOR_GPUCuda 错误 CUDA_ERROR_NO_BINARY_FOR_GPU
【发布时间】：2013-03-02 00:02:39
【问题描述】：

我有一些无法加载的 PTX 代码。我在 650M 上使用 OSX 运行它。其他 CUDA 示例在系统上运行良好，但在加载模块时，我总是收到错误 209：CUDA_ERROR_NO_BINARY_FOR_GPU

我错过了什么？

 .version 3.1
.target sm_20, texmode_independent
.address_size 64


    // .globl   examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx
.entry examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx(
    .param .u64 .ptr .global .align 8 examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_0,
    .param .f64 examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_1,
    .param .f64 examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_2,
    .param .f64 examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_3
)
{
    .reg .pred %p<396>;
    .reg .s16 %rc<396>;
    .reg .s16 %rs<396>;
    .reg .s32 %r<396>;
    .reg .s64 %rl<396>;
    .reg .f32 %f<396>;
    .reg .f64 %fl<396>;

    ld.param.u64    %rl0, [examples_2E_mandelbrot_2F_calc_2D_mandelbrot_2D_ptx_param_0];
    mov.b64 func_retval0, %rl0;
    ret;
}

【问题讨论】：

GT650M 是 sm_30 gpu。如果将 .target sm_20 更改为 .target sm_30 会发生什么？或者我应该问，你是如何生成这个 ptx 代码的？
代码是通过llvm生成的。我已将代码缩减到您在上面看到的级别。我试过 sm_10、sm_13、sm_30 和 sm_35 的目标模型。都一样。从 .entry 切换到 .func 可以加载模块，但是（当然）我找不到该函数。
即cuModuleGetFunction返回CUDA_ERROR_NOT_FOUND
您可能想通过使用nvcc -arch=sm_30 -ptx mymodule.cu 编译一些代码来创建类似的东西并分析差异。如果你把.visible 放在.entry 之前会发生什么@ 即.visible .entry examples_2E_mandelbrot... 我想你会想要.target sm_30。

标签： cuda ptx

【解决方案1】：

您收到错误是因为您的 PTX 包含语法错误，因此它永远不会编译。线

mov.b64 func_retval0, %rl0;

引用了一个标签func_retval0，但这并没有在 PTX 文件的任何地方定义。您可以尝试自己使用工具链编译 PTX 来检查这一点：

$ ptxas -arch=sm_20 own.ptx 
ptxas own.ptx, line 24; error   : Arguments mismatch for instruction 'mov'
ptxas own.ptx, line 24; error   : Unknown symbol 'func_retval0'
ptxas own.ptx, line 24; error   : Label expected for forward reference of 'func_retval0'
ptxas fatal   : Ptx assembly aborted due to errors

【讨论】：

是的！我的函数应该编译为返回 void。改变它解决了这个问题。谢谢！

【解决方案2】：

关于运行 ptxas 的好建议。我收到错误 209：问题原来是 __shared__ 内存被超额订阅。以前这将是编译时的警告。我有 Cuda 5.5，现在编译时没有警告——即使打开了详细信息。谢谢

【讨论】：