Fix runtime crash inside __kmpc_init_allocator

[jyu2]

Summary

It seems load of traits.addr should be passed in runtime call. Currently
the load of load traits.addr gets passed cause runtime to fail.

To fix this, skip the call to EmitLoadOfScalar for extra load.

[ABataev]

LG

[ronlieb]

observing failure in our amdgpu buildbot , 3 tests
see: https://lab.llvm.org/staging/#/builders/247/builds/947

7: tgt_target_table *tgt_rtl_load_binary(int32_t, __tgt_device_image *): Assertion `Table != nullptr && "Invalid table"' failed.

if you have a quick fix, please apply, otherwise please revert and fix.
if you need someone to work with on our end, let me know.

also , i see you marked one test unsupported, but it seems like it should be reported.
please consult with us to see if there is a better solution than instantly xfailing the test.

also i see two additional failures

libomptarget :: amdgcn-amd-amdhsa :: jit/empty_kernel_lvl1.c
libomptarget :: amdgcn-amd-amdhsa :: jit/empty_kernel_lvl2.c

[ronlieb]

i think the extra two failures are not related to you patch. so ignore comment about these
libomptarget :: amdgcn-amd-amdhsa :: jit/empty_kernel_lvl1.c
libomptarget :: amdgcn-amd-amdhsa :: jit/empty_kernel_lvl2.c

[ronlieb]

apologies, i just realized its a new test that you xfailed.

clang/lib/CodeGen/CGOpenMPRuntime.cpp

llvm::Value *NumTraits = llvm::ConstantInt::get(
                      ->getSize()
                      .getLimitedValue());
   LValue AllocatorTraitsLVal = CGF.EmitLValue(AllocatorTraits);
   Address Addr = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
       AllocatorTraitsLVal.getAddress(CGF), CGF.VoidPtrPtrTy, CGF.VoidPtrTy);
   AllocatorTraitsLVal = CGF.MakeAddrLValue(Addr, CGF.getContext().VoidPtrTy,
                                            AllocatorTraitsLVal.getBaseInfo(),
                                            AllocatorTraitsLVal.getTBAAInfo());
-  llvm::Value *Traits =
-      CGF.EmitLoadOfScalar(AllocatorTraitsLVal, AllocatorTraits->getExprLoc());
+  llvm::Value *Traits = Addr.getPointer();
 
   llvm::Value *AllocatorVal =
       CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
                               CGM.getModule(), OMPRTL___kmpc_init_allocator),
                           {ThreadId, MemSpaceHandle, NumTraits, Traits});
   // Store to allocator.
   CGF.EmitVarDecl(*cast<VarDecl>(
       cast<DeclRefExpr>(Allocator->IgnoreParenImpCasts())->getDecl()));

clang/test/OpenMP/target_parallel_for_simd_uses_allocators_codegen.cpp

#pragma omp target parallel for simd uses_allocators(omp_null_allocator, omp_thread_mem_alloc, my_allocator(traits))
   for (int i = 0; i < 10; ++i)
     ;
 }
 
 // CHECK: define internal void @[[TGT_REGION]](ptr {{.+}})
 // CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
 // CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
 // CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS_ADDR]])
 // CHECK: [[CONV:%.+]] = ptrtoint ptr [[ALLOCATOR]] to i64
 // CHECK: store i64 [[CONV]], ptr [[MY_ALLOCATOR_ADDR]],
 
 // Destroy allocator upon exit from the region.
 // CHECK: [[ALLOCATOR:%.+]] = load i64, ptr [[MY_ALLOCATOR_ADDR]],
 // CHECK: [[CONV:%.+]] = inttoptr i64 [[ALLOCATOR]] to ptr
 // CHECK: call void @__kmpc_destroy_allocator(i32 %{{.+}}, ptr [[CONV]])
 // CHECK: ret void
 
 #endif

clang/test/OpenMP/target_parallel_for_uses_allocators_codegen.cpp

#pragma omp target parallel for uses_allocators(omp_null_allocator, omp_thread_mem_alloc, my_allocator(traits))
   for (int i = 0; i < 10; ++i)
     ;
 }
 
 // CHECK: define internal void @[[TGT_REGION]](ptr {{.+}})
 // CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
 // CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
 // CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS_ADDR]])
 // CHECK: [[CONV:%.+]] = ptrtoint ptr [[ALLOCATOR]] to i64
 // CHECK: store i64 [[CONV]], ptr [[MY_ALLOCATOR_ADDR]],
 
 // Destroy allocator upon exit from the region.
 // CHECK: [[ALLOCATOR:%.+]] = load i64, ptr [[MY_ALLOCATOR_ADDR]],
 // CHECK: [[CONV:%.+]] = inttoptr i64 [[ALLOCATOR]] to ptr
 // CHECK: call void @__kmpc_destroy_allocator(i32 %{{.+}}, ptr [[CONV]])
 // CHECK: ret void
 
 #endif

clang/test/OpenMP/target_parallel_uses_allocators_codegen.cpp

 #pragma omp target parallel uses_allocators(omp_null_allocator, omp_thread_mem_alloc, my_allocator(traits))
   ;
 }
 
 // CHECK: define internal void @[[TGT_REGION]](ptr {{.+}})
 // CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
 // CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
 // CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS_ADDR]])
 // CHECK: [[CONV:%.+]] = ptrtoint ptr [[ALLOCATOR]] to i64
 // CHECK: store i64 [[CONV]], ptr [[MY_ALLOCATOR_ADDR]],
 
 // Destroy allocator upon exit from the region.
 // CHECK: [[ALLOCATOR:%.+]] = load i64, ptr [[MY_ALLOCATOR_ADDR]],
 // CHECK: [[CONV:%.+]] = inttoptr i64 [[ALLOCATOR]] to ptr
 // CHECK: call void @__kmpc_destroy_allocator(i32 %{{.+}}, ptr [[CONV]])
 // CHECK: ret void
 
 #endif

clang/test/OpenMP/target_simd_uses_allocators_codegen.cpp

#pragma omp target simd uses_allocators(omp_null_allocator, omp_thread_mem_alloc, my_allocator(traits))
   for (int i = 0; i < 10; ++i)
     ;
 }
 
 // CHECK: define internal void @[[TGT_REGION]](ptr {{.+}})
 // CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
 // CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
 // CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS_ADDR]])
 // CHECK: [[CONV:%.+]] = ptrtoint ptr [[ALLOCATOR]] to i64
 // CHECK: store i64 [[CONV]], ptr [[MY_ALLOCATOR_ADDR]],
 
 // Destroy allocator upon exit from the region.
 // CHECK: [[ALLOCATOR:%.+]] = load i64, ptr [[MY_ALLOCATOR_ADDR]],
 // CHECK: [[CONV:%.+]] = inttoptr i64 [[ALLOCATOR]] to ptr
 // CHECK: call void @__kmpc_destroy_allocator(i32 %{{.+}}, ptr [[CONV]])
 // CHECK: ret void
 
 #endif

clang/test/OpenMP/target_teams_distribute_parallel_for_simd_uses_allocators_codegen.cpp

#pragma omp target teams distribute parallel for simd uses_allocators(omp_null_allocator, omp_thread_mem_alloc, my_allocator(traits))
   for (int i = 0; i < 10; ++i)
     ;
 }
 
 // CHECK: define internal void @[[TGT_REGION]](ptr {{.+}})
 // CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
 // CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
 // CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS_ADDR]])
 // CHECK: [[CONV:%.+]] = ptrtoint ptr [[ALLOCATOR]] to i64
 // CHECK: store i64 [[CONV]], ptr [[MY_ALLOCATOR_ADDR]],
 
 // Destroy allocator upon exit from the region.
 // CHECK: [[ALLOCATOR:%.+]] = load i64, ptr [[MY_ALLOCATOR_ADDR]],
 // CHECK: [[CONV:%.+]] = inttoptr i64 [[ALLOCATOR]] to ptr
 // CHECK: call void @__kmpc_destroy_allocator(i32 %{{.+}}, ptr [[CONV]])
 // CHECK: ret void
 
 #endif

clang/test/OpenMP/target_teams_distribute_parallel_for_uses_allocators_codegen.cpp

#pragma omp target teams distribute parallel for uses_allocators(omp_null_allocator, omp_thread_mem_alloc, my_allocator(traits))
   for (int i = 0; i < 10; ++i)
     ;
 }
 
 // CHECK: define internal void @[[TGT_REGION]](ptr {{.+}})
 // CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
 // CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
 // CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS_ADDR]])
 // CHECK: [[CONV:%.+]] = ptrtoint ptr [[ALLOCATOR]] to i64
 // CHECK: store i64 [[CONV]], ptr [[MY_ALLOCATOR_ADDR]],
 
 // Destroy allocator upon exit from the region.
 // CHECK: [[ALLOCATOR:%.+]] = load i64, ptr [[MY_ALLOCATOR_ADDR]],
 // CHECK: [[CONV:%.+]] = inttoptr i64 [[ALLOCATOR]] to ptr
 // CHECK: call void @__kmpc_destroy_allocator(i32 %{{.+}}, ptr [[CONV]])
 // CHECK: ret void
 
 #endif

clang/test/OpenMP/target_teams_distribute_simd_uses_allocators_codegen.cpp

#pragma omp target teams distribute simd uses_allocators(omp_null_allocator, omp_thread_mem_alloc, my_allocator(traits))
   for (int i = 0; i < 10; ++i)
     ;
 }
 
 // CHECK: define internal void @[[TGT_REGION]](ptr {{.+}})
 // CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
 // CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
 // CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS_ADDR]])
 // CHECK: [[CONV:%.+]] = ptrtoint ptr [[ALLOCATOR]] to i64
 // CHECK: store i64 [[CONV]], ptr [[MY_ALLOCATOR_ADDR]],
 
 // Destroy allocator upon exit from the region.
 // CHECK: [[ALLOCATOR:%.+]] = load i64, ptr [[MY_ALLOCATOR_ADDR]],
 // CHECK: [[CONV:%.+]] = inttoptr i64 [[ALLOCATOR]] to ptr
 // CHECK: call void @__kmpc_destroy_allocator(i32 %{{.+}}, ptr [[CONV]])
 // CHECK: ret void
 
 #endif

clang/test/OpenMP/target_teams_distribute_uses_allocators_codegen.cpp

#pragma omp target teams distribute uses_allocators(omp_null_allocator, omp_thread_mem_alloc, my_allocator(traits))
   for (int i = 0; i < 10; ++i)
     ;
 }
 
 // CHECK: define internal void @[[TGT_REGION]](ptr {{.+}})
 // CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
 // CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
 // CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS_ADDR]])
 // CHECK: [[CONV:%.+]] = ptrtoint ptr [[ALLOCATOR]] to i64
 // CHECK: store i64 [[CONV]], ptr [[MY_ALLOCATOR_ADDR]],
 
 // Destroy allocator upon exit from the region.
 // CHECK: [[ALLOCATOR:%.+]] = load i64, ptr [[MY_ALLOCATOR_ADDR]],
 // CHECK: [[CONV:%.+]] = inttoptr i64 [[ALLOCATOR]] to ptr
 // CHECK: call void @__kmpc_destroy_allocator(i32 %{{.+}}, ptr [[CONV]])
 // CHECK: ret void
 
 #endif

clang/test/OpenMP/target_teams_uses_allocators_codegen.cpp

 #pragma omp target teams uses_allocators(omp_null_allocator, omp_thread_mem_alloc, my_allocator(traits))
   ;
 }
 
 // CHECK: define internal void @[[TGT_REGION]](ptr {{.+}})
 // CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
 // CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
 // CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS_ADDR]])
 // CHECK: [[CONV:%.+]] = ptrtoint ptr [[ALLOCATOR]] to i64
 // CHECK: store i64 [[CONV]], ptr [[MY_ALLOCATOR_ADDR]],
 
 // Destroy allocator upon exit from the region.
 // CHECK: [[ALLOCATOR:%.+]] = load i64, ptr [[MY_ALLOCATOR_ADDR]],
 // CHECK: [[CONV:%.+]] = inttoptr i64 [[ALLOCATOR]] to ptr
 // CHECK: call void @__kmpc_destroy_allocator(i32 %{{.+}}, ptr [[CONV]])
 // CHECK: ret void
 
 #endif

clang/test/OpenMP/target_uses_allocators.c

 // CHECK-NEXT: %.x..void.addr = call ptr @__kmpc_alloc(i32 %[[#R0]], i64 4, ptr inttoptr (i64 8 to ptr))
 // CHECK-NEXT: %[[#R1:]] = load i32, ptr %x.addr, align 4
 // CHECK-NEXT: store i32 %[[#R1]], ptr %.x..void.addr, align 4
 // CHECK-NEXT: call void @__kmpc_free(i32 %[[#R0]], ptr %.x..void.addr, ptr inttoptr (i64 8 to ptr))
 
 // CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
 // CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
 // CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS_ADDR]])
 // CHECK: [[CONV:%.+]] = ptrtoint ptr [[ALLOCATOR]] to i64
 // CHECK: store i64 [[CONV]], ptr [[MY_ALLOCATOR_ADDR]],
 
 // Destroy allocator upon exit from the region.
 // CHECK: [[ALLOCATOR:%.+]] = load i64, ptr [[MY_ALLOCATOR_ADDR]],
 // CHECK: [[CONV:%.+]] = inttoptr i64 [[ALLOCATOR]] to ptr
 // CHECK: call void @__kmpc_destroy_allocator(i32 %{{.+}}, ptr [[CONV]])

clang/test/OpenMP/target_uses_allocators_codegen.cpp

 #pragma omp target uses_allocators(omp_null_allocator, omp_thread_mem_alloc, my_allocator(traits))
   ;
 }
 
 // CHECK: define internal void @[[TGT_REGION]](ptr {{.+}})
 // CHECK: [[TRAITS_ADDR_REF:%.+]] = alloca ptr,
 // CHECK: [[MY_ALLOCATOR_ADDR:%.+]] = alloca i64,
 // CHECK: [[TRAITS_ADDR:%.+]] = load ptr, ptr [[TRAITS_ADDR_REF]],
-// CHECK: [[TRAITS:%.+]] = load ptr, ptr [[TRAITS_ADDR]],
-// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS]])
+// CHECK: [[ALLOCATOR:%.+]] = call ptr @__kmpc_init_allocator(i32 %{{.+}}, ptr null, i32 10, ptr [[TRAITS_ADDR]])
 // CHECK: [[CONV:%.+]] = ptrtoint ptr [[ALLOCATOR]] to i64
 // CHECK: store i64 [[CONV]], ptr [[MY_ALLOCATOR_ADDR]],
 
 // Destroy allocator upon exit from the region.
 // CHECK: [[ALLOCATOR:%.+]] = load i64, ptr [[MY_ALLOCATOR_ADDR]],
 // CHECK: [[CONV:%.+]] = inttoptr i64 [[ALLOCATOR]] to ptr
 // CHECK: call void @__kmpc_destroy_allocator(i32 %{{.+}}, ptr [[CONV]])
 // CHECK: ret void
 
 #endif

openmp/libomptarget/test/mapping/target_uses_allocator.c

+// RUN: %libomptarget-compile-run-and-check-generic
+
+#include <omp.h>
+#include <stdio.h>
+
+#define N 1024
+
+int test_omp_aligned_alloc_on_device() {
+  int errors = 0;
+
+  omp_memspace_handle_t memspace = omp_default_mem_space;
+  omp_alloctrait_t traits[2] = {{omp_atk_alignment, 64}, {omp_atk_access, 64}};
+  omp_allocator_handle_t alloc =
+      omp_init_allocator(omp_default_mem_space, 1, traits);
+
+#pragma omp target map(tofrom : errors) uses_allocators(alloc(traits))
+  {
+    int *x;
+    int not_correct_array_values = 0;
+
+    x = (int *)omp_aligned_alloc(64, N * sizeof(int), alloc);
+    if (x == NULL) {
+      errors++;
+    } else {
+#pragma omp parallel for simd simdlen(16) aligned(x : 64)
+      for (int i = 0; i < N; i++) {
+        x[i] = i;
+      }
+
+#pragma omp parallel for simd simdlen(16) aligned(x : 64)
+      for (int i = 0; i < N; i++) {
+        if (x[i] != i) {
+#pragma omp atomic write
+          not_correct_array_values = 1;
+        }
+      }
+      if (not_correct_array_values) {
+        errors++;
+      }
+      omp_free(x, alloc);
+    }
+  }
+
+  omp_destroy_allocator(alloc);
+
+  return errors;
+}
+
+int main() {
+  int errors = 0;
+  if (test_omp_aligned_alloc_on_device())
+    printf("FAILE\n");
+  else
+    // CHECK: PASSED
+    printf("PASSED\n");
+}