[llvm-bugs] [Bug 51552] New: -O0 built openmp offload for nvidia, occasionally fails trying to copy data

[via llvm-bugs]

https://bugs.llvm.org/show_bug.cgi?id=51552

            Bug ID: 51552
           Summary: -O0 built openmp offload for nvidia, occasionally
                    fails trying to copy data
           Product: OpenMP
           Version: unspecified
          Hardware: Other
                OS: Linux
            Status: NEW
          Severity: normal
          Priority: P
         Component: Runtime Library
          Assignee: unassignedbugs at nondot.org
          Reporter: rlieberm at amd.com
                CC: llvm-bugs at lists.llvm.org

cmake command to build trunk for nvptx
cmake -DCMAKE_INSTALL_PREFIX=/xtmp/rlieberm/trunkNvptx_1.0 \
      -DCMAKE_BUILD_TYPE=Release \
      -DLLVM_ENABLE_PROJECTS="clang" \
      -DLLVM_TARGETS_TO_BUILD="X86;NVPTX" \
      -DLLVM_ENABLE_ASSERTIONS=ON                        \
      -DLLVM_CCACHE_BUILD=$OPTION_CCACHE                 \
      -DLLVM_ENABLE_RUNTIMES="openmp" \
      -DLIBOMPTARGET_CMAKE_BUILD_TYPE=debug \
      ../llvm


compile command
 /xtmp/rlieberm/trunkNvptx_1.0//bin/clang  -O0  -target x86_64-pc-linux-gnu
-fopenmp -fopenmp-targets=nvptx64-nvidia-cuda
-Xopenmp-target=nvptx64-nvidia-cuda -march=sm_50  nested_loop.c -o nested_loop
-L/usr/local/cuda/targets/x86_64-linux/lib -lcudart


execution, this is intermittent in failure, so it may reproduce , i hope so
$ LIBOMPTARGET_INFO=16  LD_LIBRARY_PATH=/xtmp/rlieberm/trunkNvptx_1.0/lib
./nested_loop
CUDA device 0 info: Device supports up to 65536 CUDA blocks and 1024 threads
with a warp size of 32
CUDA device 0 info: Device heap size is 8388608 Bytes, device stack size is
1024 Bytes per thread
CUDA device 0 info: Launching kernel __omp_offloading_811_181cfa_main_l9 with 1
blocks and 38 threads in Generic mode
CUDA error: an illegal memory access was encountered
Libomptarget error: Copying data from device failed.
Libomptarget error: Call to targetDataEnd failed, abort target.
Libomptarget error: Failed to process data after launching the kernel.
Libomptarget error: Run with LIBOMPTARGET_INFO=4 to dump host-target pointer
mappings.
Libomptarget error: Source location information not present. Compile with -g or
-gline-tables-only.
Libomptarget fatal error 1: failure of target construct while offloading is
mandatory
Aborted (core dumped)



Test case: 
#include <stdio.h>
#include "assert.h"
#include <unistd.h>

int t1 = 6;
int main() {
  int wayout=0;

  #pragma omp target map(tofrom: wayout)
  #pragma omp teams num_teams(1  ) thread_limit(t1)
  {   wayout =1;
  }
  printf("wayout=%d\n", wayout);
  if(wayout==1){
        printf("Succeeded\n");
        return 0;
  } else {
        printf("Failed\n");
        return 1;
  }
}

-- 
You are receiving this mail because:
You are on the CC list for the bug.
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-bugs/attachments/20210819/6625c36d/attachment.html>