C++ Virtual methods: Cuda Error ( 718 ): invalid program counter

Accelerated Computing CUDA CUDA NVCC Compiler
1

Please refer the attached nvcc_virtual_method_bug.zip

  1. Attempts to use C++ virtual methods results in Cuda Error ( 718 ): invalid program counter
  2. There is a strange work-around involving adding a kernel __global__ void create_function_local() that is never used (note the suffix _local)

To expose the bug, run: build_and_run.sh

$ ./out/t_with_globals.exe
Getting started … 5 1
Doing something … @t.cu::43::main()::cudaDeviceSynchronize():
Cuda Error ( 718 ): invalid program counter
[…]

$ cuda-gdb ./out/t_with_globals.exe
[…]
CUDA Exception: Warp Invalid PC
The exception was triggered at PC 0x555555bd8aa0

Thread 1 “t_with_globals.” received signal CUDA_EXCEPTION_8, Warp Invalid PC.

nvcc_virtual_method_bug.zip (5.7 KB)