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.