[clang] [llvm] [clang][OpenMP] Fix assertion in vtable registration for incomplete types (PR #229521)

[via cfe-commits]

llvmorg-github-actions[bot] wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-clang-codegen

Author: David Pagan (ddpagan)

<details>
<summary>Changes</summary>

Assertion failure, "queried property of class with no definition", when mapping a pointer to incomplete type. Function 'emitAndRegisterVTable()' assumed types would always be complete when mapped. When it's not (as in it has no definition), there is no vtable to emit/register, so it should just be skipped rather than assert fail.

---
Full diff: https://github.com/llvm/llvm-project/pull/229521.diff


3 Files Affected:

- (modified) clang/lib/CodeGen/CGOpenMPRuntime.cpp (+1-1) 
- (added) clang/test/OpenMP/target_vtable_codegen_incomplete.cpp (+15) 
- (added) offload/test/offloading/target_map_incomplete_type.cpp (+65) 


``````````diff
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 58bfd48a49d39..40819e61bb6a1 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -1854,7 +1854,7 @@ void CGOpenMPRuntime::emitAndRegisterVTable(CodeGenModule &CGM,
                                             const VarDecl *VD) {
   // Register C++ VTable to OpenMP Offload Entry if it's a new
   // CXXRecordDecl.
-  if (CXXRecord && CXXRecord->isDynamicClass() &&
+  if (CXXRecord && CXXRecord->hasDefinition() && CXXRecord->isDynamicClass() &&
       !CGM.getOpenMPRuntime().VTableDeclMap.contains(CXXRecord)) {
     auto Res = CGM.getOpenMPRuntime().VTableDeclMap.try_emplace(CXXRecord, VD);
     if (Res.second) {
diff --git a/clang/test/OpenMP/target_vtable_codegen_incomplete.cpp b/clang/test/OpenMP/target_vtable_codegen_incomplete.cpp
new file mode 100644
index 0000000000000..c4e9a8fa0cf01
--- /dev/null
+++ b/clang/test/OpenMP/target_vtable_codegen_incomplete.cpp
@@ -0,0 +1,15 @@
+// RUN: %clang_cc1 -verify -fopenmp -Wno-openmp-mapping -x c++ -triple x86_64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -fopenmp-cuda-mode  -emit-llvm-bc %s -o %t-ppc-host.bc -fopenmp-version=52
+// RUN: %clang_cc1 -verify -fopenmp -Wno-openmp-mapping -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -fopenmp-cuda-mode  -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -debug-info-kind=limited -fopenmp-version=52 | FileCheck %s
+// expected-no-diagnostics
+
+// CHECK: target triple = "nvptx64-unknown-unknown"
+// CHECK-NOT: @_ZTV
+
+// A vtable should not be generated when mapping pointer to an incomplete type.
+class Incomplete;
+Incomplete *p;
+
+int main() {
+#pragma omp target enter data map(alloc : p)
+  return 0;
+}
diff --git a/offload/test/offloading/target_map_incomplete_type.cpp b/offload/test/offloading/target_map_incomplete_type.cpp
new file mode 100644
index 0000000000000..06ce20a5c55cd
--- /dev/null
+++ b/offload/test/offloading/target_map_incomplete_type.cpp
@@ -0,0 +1,65 @@
+// Check that mapping a pointer to an incomplete type compiles and offloads
+// correctly. The definition of the pointee must stay below the enter data
+// directive so the type is still incomplete in directive codegen. That is
+// what causes the assertion in vtable registration if the bug is present.
+
+// RUN: %libomptarget-compilexx-run-and-check-generic
+
+// REQUIRES: gpu
+
+#include <cassert>
+#include <cstdio>
+#include <omp.h>
+
+struct Incomplete;
+Incomplete *p;
+
+static void map_incomplete_pointer() {
+#pragma omp target enter data map(alloc : p)
+}
+
+struct Incomplete {
+  int a;
+  int b;
+};
+
+int main() {
+  constexpr int N = 5;
+  auto *host = new Incomplete[N];
+  for (int i = 0; i < N; ++i) {
+    host[i].a = i + 100;
+    host[i].b = i * 2 + 200;
+  }
+  p = host;
+
+  map_incomplete_pointer();
+
+  int isHost = 1;
+#pragma omp target map(tofrom : p[0 : N]) map(from : isHost)
+  {
+    isHost = omp_is_initial_device();
+    for (int i = 0; i < N; ++i) {
+      p[i].a += 1;
+      p[i].b += 2;
+    }
+  }
+
+#pragma omp target exit data map(delete : p)
+
+  for (int i = 0; i < N; ++i) {
+    assert(host[i].a == i + 101);
+    assert(host[i].b == i * 2 + 202);
+  }
+
+  delete[] host;
+
+  // CHECK: Target region executed on the device
+  std::printf("Target region executed on the %s\n", isHost ? "host" : "device");
+
+  if (isHost)
+    return 1;
+
+  // CHECK: PASS
+  std::printf("PASS\n");
+  return 0;
+}

``````````

</details>


https://github.com/llvm/llvm-project/pull/229521