https://github.com/sarnex created https://github.com/llvm/llvm-project/pull/169723
None >From 8ca3936f72205bc5e6626b7b7d05b1e81b6ce444 Mon Sep 17 00:00:00 2001 From: Nick Sarnie <[email protected]> Date: Wed, 26 Nov 2025 12:21:43 -0800 Subject: [PATCH] [clang][OpenMP] Generate pointer instead of zero element array Signed-off-by: Nick Sarnie <[email protected]> --- clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp | 7 +++++-- clang/test/OpenMP/spirv_empty_array.cpp | 17 +++++++++++++++++ 2 files changed, 22 insertions(+), 2 deletions(-) create mode 100644 clang/test/OpenMP/spirv_empty_array.cpp diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp index 572d59edb99b2..212eb3578a6b6 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp @@ -1258,9 +1258,12 @@ void CGOpenMPRuntimeGPU::emitParallelCall( // TODO: Is that needed? CodeGenFunction::OMPPrivateScope PrivateArgScope(CGF); + llvm::Type *CapturedVarsAddrType = + CapturedVars.size() ? cast<llvm::Type>(llvm::ArrayType::get( + CGM.VoidPtrTy, CapturedVars.size())) + : CGM.VoidPtrTy; Address CapturedVarsAddrs = CGF.CreateDefaultAlignTempAlloca( - llvm::ArrayType::get(CGM.VoidPtrTy, CapturedVars.size()), - "captured_vars_addrs"); + CapturedVarsAddrType, "captured_vars_addrs"); // There's something to share. if (!CapturedVars.empty()) { // Prepare for parallel region. Indicate the outlined function. diff --git a/clang/test/OpenMP/spirv_empty_array.cpp b/clang/test/OpenMP/spirv_empty_array.cpp new file mode 100644 index 0000000000000..67b687c68d5d8 --- /dev/null +++ b/clang/test/OpenMP/spirv_empty_array.cpp @@ -0,0 +1,17 @@ +// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -fopenmp-targets=spirv64-intel -emit-llvm-bc %s -o %t-host.bc +// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple spirv64-intel -fopenmp-targets=spirv64-intel -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-host.bc -o - | FileCheck -implicit-check-not="[0 x {{.*}}]" %s + +// expected-no-diagnostics + +// CHECK: [[CAPTURED_VARS:%captured_vars.*]] = alloca ptr addrspace(4), align 8 +// CHECK: [[CAPTURED_VARS_AS_CAST:%.*]] = addrspacecast ptr [[CAPTURED_VARS]] to ptr addrspace(4) +// CHECK: call spir_func addrspace(9) void @__kmpc_parallel_51(ptr addrspace(4) addrspacecast (ptr addrspace(1) @{{.*}} to ptr addrspace(4)), +// CHECK-SAME: i32 %{{.*}}, i32 1, i32 2, i32 -1, ptr addrspace(9) @__omp_offloading_{{.*}}_omp_outlined_omp_outlined, ptr addrspace(9) null, ptr addrspace(4) [[CAPTURED_VARS_AS_CAST]], i64 0) + +int puts( const char* str ); + +int main() { +#pragma omp target teams num_teams(4) +#pragma omp parallel num_threads(2) + { puts("foo"); } +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
