How to compile cuda code with calling one function twice inside one method?

user_3297655 0
Viewed 78

I'am try to compile the piece of code:

struct foo {
    unsigned long long x0;
};

//__device__ __noinline__ foo bar(foo a, foo b){  // << try this
__device__ foo bar(foo a, foo b){
    foo r;
    asm(
    ".reg .u64 my_cool_var;\n\t"
    "add.cc.u64 %0, %1, %2;\n\t"
    : "=l"(r.x0)
    : "l"(a.x0)
      "l"(b.x0)
    );
    return r;
}

__device__ foo func_with2call(foo x, foo y){
    foo res = bar(x, y);
    foo iy =  bar(x, res);
    return iy;
}

__global__ void cuda_test1(foo *a, foo *b, foo *r) {
    *r = func_with2call(*a, *b);
}

Compiler output:

ptxas /tmp/tmpxft_000010f5_00000000-6_main.ptx, line 38; error   : Duplicate definition of variable 'my_cool_var'
ptxas fatal   : Ptx assembly aborted due to errors

I know, that it is a problem with code inlining. For example, if I compile bar function with __noinline__ attribute, then there is no error. Is there a way to keep the inlining (other than duplicating the bar function code with different internal variable names), but still call bar function twice?

ptx cuda
1 Answers

This limitation is discussed in the Inline PTX Assembly Guide. You can get around it by forcing each definition into its own scope, something like:

__device__ foo bar(foo a, foo b){
    foo r;
    asm(
    "{.reg .u64 my_cool_var;\n\t"
    "add.cc.u64 %0, %1, %2;\n\t"
    "}"
    : "=l"(r.x0)
    : "l"(a.x0)
      "l"(b.x0)
    );
    return r;
}

This will safely inline without conflicts.

user_6818650
Related