[BUG?] NVRTC does not understand "statement and declaration in expression", but NVCC can handle

Accelerated Computing CUDA CUDA Programming and Performance
1

Please check the code bellow.

This is a simplest example that tries to use a temporary variable inside of expression.
nvcc accepts this statment, and works perfectly. However, NVRTC raises an compilation error for this statement.

#ifndef __CUDACC_RTC__
#include <cuda_runtime.h>
#include <stdio.h>
#endif

typedef struct {
    char    isnull;
    char    value;
} pg_bool_t;

#define IS_TRUE(x)      (!(x).isnull && (x).value)
#define IS_FALSE(x)     (!(x).isnull && (x).value)

#define AND(x,y)                                \
    ({pg_bool_t v = (x); !IS_TRUE(v) ? (v) : (y);})
#define OR(x,y)                                 \
    ({pg_bool_t v = (x);  IS_TRUE(v) ? (v) : (y);})

__device__ static pg_bool_t fn(int x)
{
    pg_bool_t   ret;

    if (x < 0)
        ret.isnull = 1;
    else
    {
        ret.isnull = 0;
        ret.value  = x;
    }
    printf("X {isnull=%d value=%d}\n", ret.isnull, ret.value);
    return ret;
}

__global__ void kernel_func(int a, int b, int c, int d)
{
    pg_bool_t   r;

    r = OR(fn(a), OR(fn(b), OR(fn(c), fn(d))));

    printf("R {isnull=%d value=%d}\n", r.isnull, r.value);
}

#ifndef __CUDACC_RTC__
int main(int argc, char *argv[])
{
    kernel_func<<<1,1>>>(0, -1, 1, 2);

    cudaStreamSynchronize(NULL);

    return 0;
}
#endif

In case of nvcc (standalone compiler).

$ nvcc ~/hoge.cu
$ ./a.out
x = 0
x = -1
x = 1
isnull=0 value=1

In case of nvrtc (run-time compiler with test program).

$ ./nvrtcc ~/hoge.cu
/home/kaigai/hoge.cu(38): error: expected an expression

/home/kaigai/hoge.cu(19): warning: function "fn" was declared but never referenced

1 error detected in the compilation of "/home/kaigai/hoge.cu".

The test program of nvrtc can be checked out from:
toybox/nvrtc at master · heterodb/toybox · GitHub

This incompatibility is inconvenient for our usage (auto GPU code generation from SQL statement), and hope this problem is fixed soon.

2

BTW, “Statements and Declarations in Expressions” is an extension of gcc, even though nvcc also supports.

3

The stated support for CUDA device code is available in the programming guide.

If you’d like to see a feature added to CUDA, the best suggestion is to file a bug. The instructions are linked in a sticky post at the top of this sub-forum.

(BTW the printout you have shown here doesn’t match the code you have posted.)

4

(BTW the printout you have shown here doesn’t match the code you have posted.)

Sorry, binary was older than the code I posted.

[kaigai@saba nvrtc]$ ./a.out
X {isnull=0 value=0}
X {isnull=1 value=-1}
X {isnull=0 value=1}
R {isnull=0 value=1}