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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
