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

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:
https://github.com/heterodb/toybox/tree/master/nvrtc

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

BTW, “Statements and Declarations in Expressions” is an extension of gcc, even though nvcc also supports.
https://gcc.gnu.org/onlinedocs/gcc/Statement-Exprs.html

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.)

(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}