Cuda error CUDA_ERROR_NO_BINARY_FOR_GPU

Question

I have some PTX code that fails to load. I'm running this on a 650M, with OSX. Other CUDA examples run fine on the system, but when loading the module I always get error 209: CUDA_ERROR_NO_BINARY_FOR_GPU

What am I missing?

 .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 is a sm_30 gpu. What happens if you change `.target sm_20` to `.target sm_30` ? Or perhaps I should ask, how did you generate this ptx code? — Robert Crovella, Mar 02 '13 at 00:16
The code is generated via llvm. I've cut down the code to the level you see above. I've tried target models of sm_10, sm_13, sm_30, and sm_35. All the same. Switching from .entry to .func allows the module to load, but then (of course) I can't find the function. — Timothy Baldridge, Mar 02 '13 at 00:29
You might want to create something similar by compiling some code with `nvcc -arch=sm_30 -ptx mymodule.cu` and analyze the differences. What happens if you put `.visible` before `.entry` i.e. `.visible .entry examples_2E_mandelbrot...` And you will want `.target sm_30`, I think. — Robert Crovella, Mar 02 '13 at 00:37

score 6 · Accepted Answer · answered Mar 02 '13 at 01:25

You are getting the error because your PTX contains a syntax error and it is never compiling as a result. The line

mov.b64 func_retval0, %rl0;

references a label func_retval0, but that isn't defined in the PTX file anywhere. You can check this by trying to compile the PTX with the toolchain yourself:

$ 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

Yep! My function should be compiled as returning void. Changing that fixed the issue. Thanks! — Timothy Baldridge, Mar 02 '13 at 05:05

score 1 · Answer 2 · edited Jan 17 '14 at 09:41

1

Great advice on running ptxas. I was getting error 209: Problem turned out to be __shared__ memory was oversubscribed. Used to be this would be a warning on compile. I have Cuda 5.5 and no warnings on compile now -- even with verbose turned on. thanks

edited Jan 17 '14 at 09:41

talonmies

70,661
34
192
269

answered Nov 13 '13 at 03:27

JPM

445
1
5
15

Cuda error CUDA_ERROR_NO_BINARY_FOR_GPU

2 Answers2