[Openmp-dev] CUDA error is: device kernel image is invalid

[Itaru Kitayama via Openmp-dev]

Here's the backtrace of the app.

(gdb) r
Starting program: /p/project/cjzam11/kitayama1/a.out
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/usr/lib64/libthread_db.so.1".

Program received signal SIGSEGV, Segmentation fault.
0x0000000000000000 in ?? ()
(gdb) where
#0  0x0000000000000000 in ?? ()
#1  0x00001555555380f9 in RTLsTy::LoadRTLs() ()
   from /p/project/cjzam11/kitayama1/projects/spack/opt/spack/linux-centos8-zen2/gcc-9.3.0/llvm-master-wbxxqwlovyuijyspx4beh4yn7lsxkcfu/lib/libomptarget.so
#2  0x000015555547c269 in std::__1::__call_once(unsigned long
volatile&, void*, void (*)(void*)) ()
   from /p/project/cjzam11/kitayama1/projects/spack/opt/spack/linux-centos8-zen2/gcc-9.3.0/llvm-master-wbxxqwlovyuijyspx4beh4yn7lsxkcfu/lib/libc++.so.1
#3  0x0000155555538f02 in RTLsTy::RegisterLib(__tgt_bin_desc*) ()
   from /p/project/cjzam11/kitayama1/projects/spack/opt/spack/linux-centos8-zen2/gcc-9.3.0/llvm-master-wbxxqwlovyuijyspx4beh4yn7lsxkcfu/lib/libomptarget.so
#4  0x0000000000401120 in omp_offloading.descriptor_reg ()
#5  0x000000000040208d in __libc_csu_init ()
#6  0x000015555410a62e in __libc_start_main () from /usr/lib64/libc.so.6
#7  0x000000000040117e in _start ()

On Sun, Dec 13, 2020 at 5:52 PM Itaru Kitayama <itaru.kitayama at gmail.com> wrote:
>
> Hi,
> On JURECA-DC at JSC, I am seeing a run-time error like below:
>
> Libomptarget --> Init target library!
> Libomptarget --> Loading RTLs...
> Libomptarget --> Loading library 'libomptarget.rtl.ppc64.so'...
> Libomptarget --> Unable to load library 'libomptarget.rtl.ppc64.so':
> libomptarget.rtl.ppc64.so: cannot open shared object file: No such
> file or directory!
> Libomptarget --> Loading library 'libomptarget.rtl.x86_64.so'...
> Libomptarget --> Successfully loaded library 'libomptarget.rtl.x86_64.so'!
> Libomptarget --> Registering RTL libomptarget.rtl.x86_64.so supporting
> 4 devices!
> Libomptarget --> Loading library 'libomptarget.rtl.cuda.so'...
> Target CUDA RTL --> Start initializing CUDA
> Libomptarget --> Successfully loaded library 'libomptarget.rtl.cuda.so'!
> Libomptarget --> Registering RTL libomptarget.rtl.cuda.so supporting 4 devices!
> Libomptarget --> Loading library 'libomptarget.rtl.aarch64.so'...
> Libomptarget --> Unable to load library 'libomptarget.rtl.aarch64.so':
> libomptarget.rtl.aarch64.so: cannot open shared object file: No such
> file or directory!
> Libomptarget --> Loading library 'libomptarget.rtl.ve.so'...
> Libomptarget --> Unable to load library 'libomptarget.rtl.ve.so':
> libomptarget.rtl.ve.so: cannot open shared object file: No such file
> or directory!
> Libomptarget --> Loading library 'libomptarget.rtl.amdgpu.so'...
> Libomptarget --> Unable to load library 'libomptarget.rtl.amdgpu.so':
> libomptarget.rtl.amdgpu.so: cannot open shared object file: No such
> file or directory!
> Libomptarget --> RTLs loaded!
> Libomptarget --> Image 0x00000000004031e0 is NOT compatible with RTL
> libomptarget.rtl.x86_64.so!
> Libomptarget --> Image 0x00000000004031e0 is compatible with RTL
> libomptarget.rtl.cuda.so!
> Libomptarget --> RTL 0x00000000021f5910 has index 0!
> Libomptarget --> Registering image 0x00000000004031e0 with RTL
> libomptarget.rtl.cuda.so!
> Libomptarget --> Done registering entries!
> Libomptarget --> Call to omp_get_num_devices returning 4
> Libomptarget --> Default TARGET OFFLOAD policy is now mandatory
> (devices were found)
> Libomptarget --> Entering target region with entry point
> 0x0000000000403058 and device Id -1
> Libomptarget --> Checking whether device 0 is ready.
> Libomptarget --> Is the device 0 (local ID 0) initialized? 0
> Target CUDA RTL --> Init requires flags to 1
> Target CUDA RTL --> Getting device 0
> Target CUDA RTL --> The primary context is inactive, set its flags to
> CU_CTX_SCHED_BLOCKING_SYNC
> Target CUDA RTL --> Max CUDA blocks per grid 2147483647 exceeds the
> hard team limit 65536, capping at the hard limit
> Target CUDA RTL --> Using 1024 CUDA threads per block
> Target CUDA RTL --> Using warp size 32
> Target CUDA RTL --> Device supports up to 65536 CUDA blocks and 1024
> threads with a warp size of 32
> Target CUDA RTL --> Default number of teams set according to library's
> default 128
> Target CUDA RTL --> Default number of threads set according to
> library's default 128
> Libomptarget --> Device 0 is ready to use.
> Target CUDA RTL --> Load data from image 0x00000000004031e0
> Target CUDA RTL --> Error returned from cuModuleLoadDataEx
> Target CUDA RTL --> CUDA error is: device kernel image is invalid
> Libomptarget --> Unable to generate entries table for device id 0.
> Libomptarget --> Failed to init globals on device 0
> Libomptarget --> Failed to get device 0 ready
> Libomptarget error: run with env LIBOMPTARGET_INFO>1 to dump
> host-target pointer maps
> Libomptarget error: Build with debug information to provide more
> informationLibomptarget fatal error 1: failure of target construct
> while offloading is mandatory
> /var/spool/parastation/jobs/8813449: line 22: 22331 Aborted
>      (core dumped) ./a.out
>
> Here's the sample code:
>
> #include <stdlib.h>
> #include <stdio.h>
> #include <unistd.h>
>
> #include <omp.h>
>
> static double inline calc_pi(size_t n)
> {
>         double h = 1.0 / n;
>         double sum = 0.0, x;
>         size_t i;
>
>         #pragma omp parallel for simd private(x) reduction(+:sum)
>         for (i = 0; i < n; i++) {
>                 x = h * ((double) i + 0.5);
>                 sum += 4.0 / (1.0 + x * x);
>         }
>
>         return sum * h;
> }
> //#pragma omp declare target
> int t1(void) {
>         printf("eeee\n");
>         return 0;
> }
> //#pragma omp end declare target
> struct A {
>
> };
>
> const size_t N = 2000000000;
>
> int main()
> {
>         double start, pi, end;
>
>         #pragma omp target
>         {}
>
>         printf(">> host: start\n");
>         start = omp_get_wtime();
>         pi = calc_pi(N);
>         end = omp_get_wtime();
>         printf(">> host: pi = %f (%f ms)\n", pi, (end - start) * 1000);
>
>         printf(">> target: start\n");
>         start = omp_get_wtime();
>
>         double h = 1.0 / N;
>         double sum = 0.0, x;
>         size_t i;
>         #pragma omp target teams distribute parallel for simd
> private(x) map(sum) reduction(+:sum)
>         for (i = 0; i < N; i++) {
>                 x = h * ((double) i + 0.5);
>                 sum += 4.0 / (1.0 + x * x);
>                 //t1();
>                 //struct A a;
>         }
>
>         pi = sum * h;
>         end = omp_get_wtime();
>         printf(">> target: pi = %f (%f ms)\n", pi, (end - start) * 1000);
>
>         return EXIT_SUCCESS;
> }