[clang] [clang][OpenMP] Generate pointer instead of zero element array (PR #169723)

[Nick Sarnie via cfe-commits]

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 <nick.sarnie at intel.com>
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 <nick.sarnie at intel.com>
---
 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"); }
+}