[flang-commits] [clang] [flang] [llvm] [mlir] [OpenMP] Introduce the ompx_name clause for kernel naming (PR #200301)
Johannes Doerfert via flang-commits
flang-commits at lists.llvm.org
Tue Jun 23 14:55:17 PDT 2026
https://github.com/jdoerfert updated https://github.com/llvm/llvm-project/pull/200301
>From 642e8b6aaf7b7dfcf78358fe7ed2402df6aee079 Mon Sep 17 00:00:00 2001
From: Johannes Doerfert <johannes at jdoerfert.de>
Date: Mon, 1 Jun 2026 13:04:56 -0700
Subject: [PATCH 1/2] [OpenMP] Use ext linkage for kernels handles and globals
handles keep linkage
Host handles are now emmitted with external linkage to clash if two
kernels with the same name are registered. This could have happen right
now and silently corrupt the program, but it can happen more easily once
we allow users to name their kernels.
In the same patch we make global variable handles retain the linkage of
the global variable, forcing clashes for external ones and continue to
support weak use cases.
---
clang/test/OpenMP/amdgcn_weak_alias.c | 4 ++--
clang/test/OpenMP/declare_target_codegen.cpp | 2 +-
clang/test/OpenMP/target_codegen.cpp | 22 +++++++++----------
clang/test/OpenMP/target_depend_codegen.cpp | 4 ++--
clang/test/OpenMP/target_indirect_codegen.cpp | 8 +++----
.../OpenMP/target_parallel_depend_codegen.cpp | 4 ++--
.../target_parallel_for_depend_codegen.cpp | 4 ++--
...arget_parallel_for_simd_depend_codegen.cpp | 4 ++--
clang/test/OpenMP/target_simd_codegen.cpp | 16 +++++++-------
.../OpenMP/target_simd_depend_codegen.cpp | 4 ++--
.../OpenMP/target_teams_depend_codegen.cpp | 4 ++--
...target_teams_distribute_depend_codegen.cpp | 4 ++--
...distribute_parallel_for_depend_codegen.cpp | 4 ++--
...ibute_parallel_for_simd_depend_codegen.cpp | 4 ++--
...t_teams_distribute_simd_depend_codegen.cpp | 4 ++--
.../llvm/Frontend/Offloading/Utility.h | 8 +++----
llvm/lib/Frontend/Offloading/Utility.cpp | 13 ++++++-----
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 7 +++---
.../omptarget-declare-target-llvm-host.mlir | 16 +++++++-------
.../omptarget-declare-target-to-host.mlir | 2 +-
20 files changed, 70 insertions(+), 68 deletions(-)
diff --git a/clang/test/OpenMP/amdgcn_weak_alias.c b/clang/test/OpenMP/amdgcn_weak_alias.c
index 4cc54b9f15b43..6292bb5640a79 100644
--- a/clang/test/OpenMP/amdgcn_weak_alias.c
+++ b/clang/test/OpenMP/amdgcn_weak_alias.c
@@ -10,9 +10,9 @@
// HOST: @__Two_var = global i32 2, align 4
// HOST: @__Three_var = global i32 3, align 4
// HOST: @.offloading.entry_name = internal unnamed_addr constant [10 x i8] c"__Two_var\00", section ".llvm.rodata.offloading", align 1
-// HOST: @.offloading.entry.__Two_var = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @__Two_var, ptr @.offloading.entry_name, i64 4, i64 0, ptr null }, section "llvm_offload_entries", align 8
+// HOST: @.offloading.entry.__Two_var = constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @__Two_var, ptr @.offloading.entry_name, i64 4, i64 0, ptr null }, section "llvm_offload_entries", align 8
// HOST: @.offloading.entry_name.1 = internal unnamed_addr constant [12 x i8] c"__Three_var\00", section ".llvm.rodata.offloading", align 1
-// HOST: @.offloading.entry.__Three_var = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @__Three_var, ptr @.offloading.entry_name.1, i64 4, i64 0, ptr null }, section "llvm_offload_entries", align 8
+// HOST: @.offloading.entry.__Three_var = constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @__Three_var, ptr @.offloading.entry_name.1, i64 4, i64 0, ptr null }, section "llvm_offload_entries", align 8
// HOST: @One = weak alias i32 (), ptr @__One
// HOST: @One_ = alias i32 (), ptr @__One
// HOST: @One_var = weak alias i32, ptr @__One_var
diff --git a/clang/test/OpenMP/declare_target_codegen.cpp b/clang/test/OpenMP/declare_target_codegen.cpp
index 6add3eed1f226..89899a1cb4f30 100644
--- a/clang/test/OpenMP/declare_target_codegen.cpp
+++ b/clang/test/OpenMP/declare_target_codegen.cpp
@@ -30,7 +30,7 @@
// CHECK-DAG: @dx = {{protected | }}global i32 0,
// CHECK-DAG: @dy = {{protected | }}global i32 0,
// CHECK-DAG: @bbb = {{protected | }}global i32 0,
-// CHECK-DAG: weak constant %struct.__tgt_offload_entry {
+// CHECK-DAG: constant %struct.__tgt_offload_entry {
// CHECK-DAG: @ccc = external global i32,
// CHECK-DAG: @ddd = {{protected | }}global i32 0,
// CHECK-DAG: @hhh_decl_tgt_ref_ptr = weak global ptr null
diff --git a/clang/test/OpenMP/target_codegen.cpp b/clang/test/OpenMP/target_codegen.cpp
index 431b09e81714b..34a02d85858ac 100644
--- a/clang/test/OpenMP/target_codegen.cpp
+++ b/clang/test/OpenMP/target_codegen.cpp
@@ -102,17 +102,17 @@
// CHECK-DAG: @{{.*}} = weak constant i8 0
// CHECK-DAG: @{{.*}} = weak constant i8 0
-// TCHECK: @{{.+}} = weak constant [[ENTTY]]
-// TCHECK: @{{.+}} = weak constant [[ENTTY]]
-// TCHECK: @{{.+}} = weak constant [[ENTTY]]
-// TCHECK: @{{.+}} = weak constant [[ENTTY]]
-// TCHECK: @{{.+}} = weak constant [[ENTTY]]
-// TCHECK: @{{.+}} = weak constant [[ENTTY]]
-// TCHECK: @{{.+}} = weak constant [[ENTTY]]
-// TCHECK: @{{.+}} = weak constant [[ENTTY]]
-// TCHECK: @{{.+}} = weak constant [[ENTTY]]
-// TCHECK: @{{.+}} = weak constant [[ENTTY]]
-// TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
+// TCHECK-NOT: @{{.+}} = constant [[ENTTY]]
template<typename tx, typename ty>
struct TT{
diff --git a/clang/test/OpenMP/target_depend_codegen.cpp b/clang/test/OpenMP/target_depend_codegen.cpp
index 86b70dd73680d..cf221386eaee9 100644
--- a/clang/test/OpenMP/target_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_depend_codegen.cpp
@@ -47,9 +47,9 @@
// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [4 x i64] [i64 544, i64 800, i64 3, i64 288]
// CHECK-DAG: @{{.*}} = weak constant i8 0
-// TCHECK: @{{.+}} = weak constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
// TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]]
-// TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]]
+// TCHECK-NOT: @{{.+}} = constant [[ENTTY]]
template<typename tx, typename ty>
struct TT{
diff --git a/clang/test/OpenMP/target_indirect_codegen.cpp b/clang/test/OpenMP/target_indirect_codegen.cpp
index fd8b6c76d0881..ba161ff8cf94d 100644
--- a/clang/test/OpenMP/target_indirect_codegen.cpp
+++ b/clang/test/OpenMP/target_indirect_codegen.cpp
@@ -23,13 +23,13 @@
// HOST: @indirect_foo = global ptr @_Z3foov, align 8
// HOST: @indirect_array = global [3 x ptr] [ptr @_Z3foov, ptr @_ZL3barv, ptr @_Z3bazv], align 8
// HOST: @[[FOO_ENTRY_NAME:.+]] = internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[FOO_NAME:__omp_offloading_[0-9a-z]+_[0-9a-z]+_foo_l[0-9]+]]\00"
-// HOST: @.offloading.entry.[[FOO_NAME]] = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 8, ptr @_Z3foov, ptr @[[FOO_ENTRY_NAME]], i64 8, i64 0, ptr null }
+// HOST: @.offloading.entry.[[FOO_NAME]] = weak_odr constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 8, ptr @_Z3foov, ptr @[[FOO_ENTRY_NAME]], i64 8, i64 0, ptr null }
// HOST: @[[BAZ_ENTRY_NAME:.+]] = internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[BAZ_NAME:__omp_offloading_[0-9a-z]+_[0-9a-z]+_baz_l[0-9]+]]\00"
-// HOST: @.offloading.entry.[[BAZ_NAME]] = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 8, ptr @_Z3bazv, ptr @[[BAZ_ENTRY_NAME]], i64 8, i64 0, ptr null }
+// HOST: @.offloading.entry.[[BAZ_NAME]] = weak_odr constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 8, ptr @_Z3bazv, ptr @[[BAZ_ENTRY_NAME]], i64 8, i64 0, ptr null }
// HOST: @[[VAR_ENTRY_NAME:.+]] = internal unnamed_addr constant [4 x i8] c"var\00"
-// HOST: @.offloading.entry.var = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @[[VAR]], ptr @[[VAR_ENTRY_NAME]], i64 1, i64 0, ptr null }
+// HOST: @.offloading.entry.var = constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @[[VAR]], ptr @[[VAR_ENTRY_NAME]], i64 1, i64 0, ptr null }
// HOST: @[[BAR_ENTRY_NAME:.+]] = internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[BAR_NAME:__omp_offloading_[0-9a-z]+_[0-9a-z]+_bar_l[0-9]+]]\00"
-// HOST: @.offloading.entry.[[BAR_NAME]] = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 8, ptr @_ZL3barv, ptr @[[BAR_ENTRY_NAME]], i64 8, i64 0, ptr null }
+// HOST: @.offloading.entry.[[BAR_NAME]] = weak_odr constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 8, ptr @_ZL3barv, ptr @[[BAR_ENTRY_NAME]], i64 8, i64 0, ptr null }
//.
// DEVICE: @[[FOO_NAME:__omp_offloading_[0-9a-z]+_[0-9a-z]+_foo_l[0-9]+]] = protected addrspace(1) constant {{ptr|ptr addrspace\(9\)}} @_Z3foov
// DEVICE: @[[BAZ_NAME:__omp_offloading_[0-9a-z]+_[0-9a-z]+_baz_l[0-9]+]] = protected addrspace(1) constant {{ptr|ptr addrspace\(9\)}} @_Z3bazv
diff --git a/clang/test/OpenMP/target_parallel_depend_codegen.cpp b/clang/test/OpenMP/target_parallel_depend_codegen.cpp
index ae41454a5c376..6966f1bbc5a7b 100644
--- a/clang/test/OpenMP/target_parallel_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_parallel_depend_codegen.cpp
@@ -47,9 +47,9 @@
// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [3 x i64] [i64 544, i64 800, i64 288]
// CHECK-DAG: @{{.*}} = weak constant i8 0
-// TCHECK: @{{.+}} = weak constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
// TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]]
-// TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]]
+// TCHECK-NOT: @{{.+}} = constant [[ENTTY]]
template<typename tx, typename ty>
struct TT{
diff --git a/clang/test/OpenMP/target_parallel_for_depend_codegen.cpp b/clang/test/OpenMP/target_parallel_for_depend_codegen.cpp
index 7cccb2549c2c8..8d79b37ea46c9 100644
--- a/clang/test/OpenMP/target_parallel_for_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_parallel_for_depend_codegen.cpp
@@ -47,9 +47,9 @@
// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [3 x i64] [i64 544, i64 800, i64 288]
// CHECK-DAG: @{{.*}} = weak constant i8 0
-// TCHECK: @{{.+}} = weak constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
// TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]]
-// TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]]
+// TCHECK-NOT: @{{.+}} = constant [[ENTTY]]
template<typename tx, typename ty>
struct TT{
diff --git a/clang/test/OpenMP/target_parallel_for_simd_depend_codegen.cpp b/clang/test/OpenMP/target_parallel_for_simd_depend_codegen.cpp
index f1391cd26e2d4..cacde85ca6e82 100644
--- a/clang/test/OpenMP/target_parallel_for_simd_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_parallel_for_simd_depend_codegen.cpp
@@ -47,9 +47,9 @@
// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [3 x i64] [i64 544, i64 800, i64 288]
// CHECK-DAG: @{{.*}} = weak constant i8 0
-// TCHECK: @{{.+}} = weak constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
// TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]]
-// TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]]
+// TCHECK-NOT: @{{.+}} = constant [[ENTTY]]
template<typename tx, typename ty>
struct TT{
diff --git a/clang/test/OpenMP/target_simd_codegen.cpp b/clang/test/OpenMP/target_simd_codegen.cpp
index 0c2dde23f6c46..141fa6ffe385b 100644
--- a/clang/test/OpenMP/target_simd_codegen.cpp
+++ b/clang/test/OpenMP/target_simd_codegen.cpp
@@ -101,14 +101,14 @@
// CHECK-DAG: @{{.*}} = weak constant i8 0
// CHECK-DAG: @{{.*}} = weak constant i8 0
-// TCHECK: @{{.+}} = weak constant [[ENTTY]]
-// TCHECK: @{{.+}} = weak constant [[ENTTY]]
-// TCHECK: @{{.+}} = weak constant [[ENTTY]]
-// TCHECK: @{{.+}} = weak constant [[ENTTY]]
-// TCHECK: @{{.+}} = weak constant [[ENTTY]]
-// TCHECK: @{{.+}} = weak constant [[ENTTY]]
-// TCHECK: @{{.+}} = weak constant [[ENTTY]]
-// TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
+// TCHECK-NOT: @{{.+}} = constant [[ENTTY]]
template<typename tx, typename ty>
struct TT{
diff --git a/clang/test/OpenMP/target_simd_depend_codegen.cpp b/clang/test/OpenMP/target_simd_depend_codegen.cpp
index e399998869cf5..53a4f6ce9897b 100644
--- a/clang/test/OpenMP/target_simd_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_simd_depend_codegen.cpp
@@ -47,9 +47,9 @@
// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [3 x i64] [i64 544, i64 800, i64 288]
// CHECK-DAG: @{{.*}} = weak constant i8 0
-// TCHECK: @{{.+}} = weak constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
// TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]]
-// TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]]
+// TCHECK-NOT: @{{.+}} = constant [[ENTTY]]
template<typename tx, typename ty>
struct TT{
diff --git a/clang/test/OpenMP/target_teams_depend_codegen.cpp b/clang/test/OpenMP/target_teams_depend_codegen.cpp
index 1b7e25ee7e936..3bc16dc41c610 100644
--- a/clang/test/OpenMP/target_teams_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_depend_codegen.cpp
@@ -47,9 +47,9 @@
// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [3 x i64] [i64 544, i64 800, i64 288]
// CHECK-DAG: @{{.*}} = weak constant i8 0
-// TCHECK: @{{.+}} = weak constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
// TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]]
-// TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]]
+// TCHECK-NOT: @{{.+}} = constant [[ENTTY]]
template<typename tx, typename ty>
struct TT{
diff --git a/clang/test/OpenMP/target_teams_distribute_depend_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_depend_codegen.cpp
index 5bf4615fe7b70..c146a36ec9b90 100644
--- a/clang/test/OpenMP/target_teams_distribute_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_depend_codegen.cpp
@@ -47,9 +47,9 @@
// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [3 x i64] [i64 544, i64 800, i64 288]
// CHECK-DAG: @{{.*}} = weak constant i8 0
-// TCHECK: @{{.+}} = weak constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
// TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]]
-// TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]]
+// TCHECK-NOT: @{{.+}} = constant [[ENTTY]]
template<typename tx, typename ty>
struct TT{
diff --git a/clang/test/OpenMP/target_teams_distribute_parallel_for_depend_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_parallel_for_depend_codegen.cpp
index 9fd3ca822a38b..f4d6c005d7d54 100644
--- a/clang/test/OpenMP/target_teams_distribute_parallel_for_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_depend_codegen.cpp
@@ -47,9 +47,9 @@
// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [3 x i64] [i64 544, i64 800, i64 288]
// CHECK-DAG: @{{.*}} = weak constant i8 0
-// TCHECK: @{{.+}} = weak constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
// TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]]
-// TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]]
+// TCHECK-NOT: @{{.+}} = constant [[ENTTY]]
template<typename tx, typename ty>
struct TT{
diff --git a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_depend_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_depend_codegen.cpp
index 9393d9d0474bd..fc8114ed70f7f 100644
--- a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_depend_codegen.cpp
@@ -47,9 +47,9 @@
// CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [3 x i64] [i64 544, i64 800, i64 288]
// CHECK-DAG: @{{.*}} = weak constant i8 0
-// TCHECK: @{{.+}} = weak constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
// TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]]
-// TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]]
+// TCHECK-NOT: @{{.+}} = constant [[ENTTY]]
template<typename tx, typename ty>
struct TT{
diff --git a/clang/test/OpenMP/target_teams_distribute_simd_depend_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_simd_depend_codegen.cpp
index fd5cea7ebd9a0..47cef10da1b4e 100644
--- a/clang/test/OpenMP/target_teams_distribute_simd_depend_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_simd_depend_codegen.cpp
@@ -61,9 +61,9 @@
// OMP50-DAG: @{{.*}} = weak constant i8 0
-// TCHECK: @{{.+}} = weak constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
// TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]]
-// TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]]
+// TCHECK-NOT: @{{.+}} = constant [[ENTTY]]
template<typename tx, typename ty>
struct TT{
diff --git a/llvm/include/llvm/Frontend/Offloading/Utility.h b/llvm/include/llvm/Frontend/Offloading/Utility.h
index 4c0bc87786dfb..b681e8bb59bc9 100644
--- a/llvm/include/llvm/Frontend/Offloading/Utility.h
+++ b/llvm/include/llvm/Frontend/Offloading/Utility.h
@@ -91,10 +91,10 @@ LLVM_ABI StructType *getEntryTy(Module &M);
LLVM_ABI StringRef getOffloadEntrySection(Module &M);
/// \return The emitted global variable containing the offloading entry.
-LLVM_ABI GlobalVariable *
-emitOffloadingEntry(Module &M, object::OffloadKind Kind, Constant *Addr,
- StringRef Name, uint64_t Size, uint32_t Flags,
- uint64_t Data, Constant *AuxAddr = nullptr);
+LLVM_ABI GlobalVariable *emitOffloadingEntry(
+ Module &M, object::OffloadKind Kind, Constant *Addr, StringRef Name,
+ uint64_t Size, uint32_t Flags, uint64_t Data, Constant *AuxAddr = nullptr,
+ GlobalValue::LinkageTypes Linkage = GlobalValue::WeakAnyLinkage);
/// Create a constant struct initializer used to register this global at
/// runtime.
diff --git a/llvm/lib/Frontend/Offloading/Utility.cpp b/llvm/lib/Frontend/Offloading/Utility.cpp
index c07d276244ee1..bec13d67bb9ae 100644
--- a/llvm/lib/Frontend/Offloading/Utility.cpp
+++ b/llvm/lib/Frontend/Offloading/Utility.cpp
@@ -104,7 +104,8 @@ getOffloadEntryBoundarySymbols(const Triple &T, StringRef SectionName) {
GlobalVariable *offloading::emitOffloadingEntry(
Module &M, object::OffloadKind Kind, Constant *Addr, StringRef Name,
- uint64_t Size, uint32_t Flags, uint64_t Data, Constant *AuxAddr) {
+ uint64_t Size, uint32_t Flags, uint64_t Data, Constant *AuxAddr,
+ GlobalValue::LinkageTypes Linkage) {
const llvm::Triple &Triple = M.getTargetTriple();
StringRef SectionName = getOffloadEntrySection(M);
@@ -113,11 +114,11 @@ GlobalVariable *offloading::emitOffloadingEntry(
StringRef Prefix =
Triple.isNVPTX() ? "$offloading$entry$" : ".offloading.entry.";
- auto *Entry = new GlobalVariable(
- M, getEntryTy(M),
- /*isConstant=*/true, GlobalValue::WeakAnyLinkage, EntryInitializer,
- Prefix + Name, nullptr, GlobalValue::NotThreadLocal,
- M.getDataLayout().getDefaultGlobalsAddressSpace());
+ auto *Entry =
+ new GlobalVariable(M, getEntryTy(M),
+ /*isConstant=*/true, Linkage, EntryInitializer,
+ Prefix + Name, nullptr, GlobalValue::NotThreadLocal,
+ M.getDataLayout().getDefaultGlobalsAddressSpace());
// The entry has to be created in the section the linker expects it to be.
if (Triple.isOSBinFormatCOFF())
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index 92eb7de0d882f..d1de4c7bd03f8 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -11794,12 +11794,13 @@ std::unique_ptr<CodeExtractor> DeviceSharedMemOutlineInfo::createCodeExtractor(
void OpenMPIRBuilder::createOffloadEntry(Constant *ID, Constant *Addr,
uint64_t Size, int32_t Flags,
- GlobalValue::LinkageTypes,
+ GlobalValue::LinkageTypes Linkage,
StringRef Name) {
if (!Config.isGPU()) {
llvm::offloading::emitOffloadingEntry(
M, object::OffloadKind::OFK_OpenMP, ID,
- Name.empty() ? Addr->getName() : Name, Size, Flags, /*Data=*/0);
+ Name.empty() ? Addr->getName() : Name, Size, Flags, /*Data=*/0,
+ /*AuxAddr*/ nullptr, Linkage);
return;
}
// TODO: Add support for global variables on the device after declare target
@@ -11910,7 +11911,7 @@ void OpenMPIRBuilder::createOffloadEntriesAndInfoMetadata(
}
createOffloadEntry(CE->getID(), CE->getAddress(),
/*Size=*/0, CE->getFlags(),
- GlobalValue::WeakAnyLinkage);
+ GlobalValue::ExternalLinkage);
} else if (const auto *CE = dyn_cast<
OffloadEntriesInfoManager::OffloadEntryInfoDeviceGlobalVar>(
E.first)) {
diff --git a/mlir/test/Target/LLVMIR/omptarget-declare-target-llvm-host.mlir b/mlir/test/Target/LLVMIR/omptarget-declare-target-llvm-host.mlir
index 92c85738dbc72..425d8d52f0a82 100644
--- a/mlir/test/Target/LLVMIR/omptarget-declare-target-llvm-host.mlir
+++ b/mlir/test/Target/LLVMIR/omptarget-declare-target-llvm-host.mlir
@@ -53,7 +53,7 @@ module attributes {llvm.target_triple = "x86_64-unknown-linux-gnu", omp.is_targe
// CHECK-DAG: @_QMtest_0Edata_extended_to_1 = global float 2.000000e+00
// CHECK-DAG: @.offloading.entry_name{{.*}} = internal unnamed_addr constant [29 x i8] c"_QMtest_0Edata_extended_to_1\00"
- // CHECK-DAG: @.offloading.entry._QMtest_0Edata_extended_to_1 = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @_QMtest_0Edata_extended_to_1, ptr @.offloading.entry_name{{.*}}, i64 4, i64 0, ptr null }, section "llvm_offload_entries"
+ // CHECK-DAG: @.offloading.entry._QMtest_0Edata_extended_to_1 = constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @_QMtest_0Edata_extended_to_1, ptr @.offloading.entry_name{{.*}}, i64 4, i64 0, ptr null }, section "llvm_offload_entries"
// CHECK-DAG: !{{.*}} = !{i32 {{.*}}, !"_QMtest_0Edata_extended_to_1", i32 {{.*}}, i32 {{.*}}}
llvm.mlir.global external @_QMtest_0Edata_extended_to_1() {addr_space = 0 : i32, omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>} : f32 {
%0 = llvm.mlir.constant(2.000000e+00 : f32) : f32
@@ -62,7 +62,7 @@ module attributes {llvm.target_triple = "x86_64-unknown-linux-gnu", omp.is_targe
// CHECK-DAG: @_QMtest_0Edata_extended_enter_1 = global float 2.000000e+00
// CHECK-DAG: @.offloading.entry_name{{.*}} = internal unnamed_addr constant [32 x i8] c"_QMtest_0Edata_extended_enter_1\00"
- // CHECK-DAG: @.offloading.entry._QMtest_0Edata_extended_enter_1 = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @_QMtest_0Edata_extended_enter_1, ptr @.offloading.entry_name{{.*}}, i64 4, i64 0, ptr null }, section "llvm_offload_entries"
+ // CHECK-DAG: @.offloading.entry._QMtest_0Edata_extended_enter_1 = constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @_QMtest_0Edata_extended_enter_1, ptr @.offloading.entry_name{{.*}}, i64 4, i64 0, ptr null }, section "llvm_offload_entries"
// CHECK-DAG: !{{.*}} = !{i32 {{.*}}, !"_QMtest_0Edata_extended_enter_1", i32 {{.*}}, i32 {{.*}}}
llvm.mlir.global external @_QMtest_0Edata_extended_enter_1() {addr_space = 0 : i32, omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>} : f32 {
%0 = llvm.mlir.constant(2.000000e+00 : f32) : f32
@@ -71,7 +71,7 @@ module attributes {llvm.target_triple = "x86_64-unknown-linux-gnu", omp.is_targe
// CHECK-DAG: @_QMtest_0Edata_extended_to_2 = global float 3.000000e+00
// CHECK-DAG: @.offloading.entry_name{{.*}} = internal unnamed_addr constant [29 x i8] c"_QMtest_0Edata_extended_to_2\00"
- // CHECK-DAG: @.offloading.entry._QMtest_0Edata_extended_to_2 = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @_QMtest_0Edata_extended_to_2, ptr @.offloading.entry_name{{.*}}, i64 4, i64 0, ptr null }, section "llvm_offload_entries"
+ // CHECK-DAG: @.offloading.entry._QMtest_0Edata_extended_to_2 = constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @_QMtest_0Edata_extended_to_2, ptr @.offloading.entry_name{{.*}}, i64 4, i64 0, ptr null }, section "llvm_offload_entries"
// CHECK-DAG: !{{.*}} = !{i32 {{.*}}, !"_QMtest_0Edata_extended_to_2", i32 {{.*}}, i32 {{.*}}}
llvm.mlir.global external @_QMtest_0Edata_extended_to_2() {addr_space = 0 : i32, omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>} : f32 {
%0 = llvm.mlir.constant(3.000000e+00 : f32) : f32
@@ -80,7 +80,7 @@ module attributes {llvm.target_triple = "x86_64-unknown-linux-gnu", omp.is_targe
// CHECK-DAG: @_QMtest_0Edata_extended_enter_2 = global float 3.000000e+00
// CHECK-DAG: @.offloading.entry_name{{.*}} = internal unnamed_addr constant [32 x i8] c"_QMtest_0Edata_extended_enter_2\00"
- // CHECK-DAG: @.offloading.entry._QMtest_0Edata_extended_enter_2 = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @_QMtest_0Edata_extended_enter_2, ptr @.offloading.entry_name{{.*}}, i64 4, i64 0, ptr null }, section "llvm_offload_entries"
+ // CHECK-DAG: @.offloading.entry._QMtest_0Edata_extended_enter_2 = constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @_QMtest_0Edata_extended_enter_2, ptr @.offloading.entry_name{{.*}}, i64 4, i64 0, ptr null }, section "llvm_offload_entries"
// CHECK-DAG: !{{.*}} = !{i32 {{.*}}, !"_QMtest_0Edata_extended_enter_2", i32 {{.*}}, i32 {{.*}}}
llvm.mlir.global external @_QMtest_0Edata_extended_enter_2() {addr_space = 0 : i32, omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>} : f32 {
%0 = llvm.mlir.constant(3.000000e+00 : f32) : f32
@@ -99,7 +99,7 @@ module attributes {llvm.target_triple = "x86_64-unknown-linux-gnu", omp.is_targe
// CHECK-DAG: @_QMtest_0Edata_int_clauseless_to = global i32 1
// CHECK-DAG: @.offloading.entry_name{{.*}} = internal unnamed_addr constant [33 x i8] c"_QMtest_0Edata_int_clauseless_to\00"
- // CHECK-DAG: @.offloading.entry._QMtest_0Edata_int_clauseless_to = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @_QMtest_0Edata_int_clauseless_to, ptr @.offloading.entry_name{{.*}}, i64 4, i64 0, ptr null }, section "llvm_offload_entries"
+ // CHECK-DAG: @.offloading.entry._QMtest_0Edata_int_clauseless_to = constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @_QMtest_0Edata_int_clauseless_to, ptr @.offloading.entry_name{{.*}}, i64 4, i64 0, ptr null }, section "llvm_offload_entries"
// CHECK-DAG: !{{.*}} = !{i32 {{.*}}, !"_QMtest_0Edata_int_clauseless_to", i32 {{.*}}, i32 {{.*}}}
llvm.mlir.global external @_QMtest_0Edata_int_clauseless_to() {addr_space = 0 : i32, omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>} : i32 {
%0 = llvm.mlir.constant(1 : i32) : i32
@@ -108,7 +108,7 @@ module attributes {llvm.target_triple = "x86_64-unknown-linux-gnu", omp.is_targe
// CHECK-DAG: @_QMtest_0Edata_int_clauseless_enter = global i32 1
// CHECK-DAG: @.offloading.entry_name{{.*}} = internal unnamed_addr constant [36 x i8] c"_QMtest_0Edata_int_clauseless_enter\00"
- // CHECK-DAG: @.offloading.entry._QMtest_0Edata_int_clauseless_enter = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @_QMtest_0Edata_int_clauseless_enter, ptr @.offloading.entry_name{{.*}}, i64 4, i64 0, ptr null }, section "llvm_offload_entries"
+ // CHECK-DAG: @.offloading.entry._QMtest_0Edata_int_clauseless_enter = constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @_QMtest_0Edata_int_clauseless_enter, ptr @.offloading.entry_name{{.*}}, i64 4, i64 0, ptr null }, section "llvm_offload_entries"
// CHECK-DAG: !{{.*}} = !{i32 {{.*}}, !"_QMtest_0Edata_int_clauseless_enter", i32 {{.*}}, i32 {{.*}}}
llvm.mlir.global external @_QMtest_0Edata_int_clauseless_enter() {addr_space = 0 : i32, omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>} : i32 {
%0 = llvm.mlir.constant(1 : i32) : i32
@@ -117,7 +117,7 @@ module attributes {llvm.target_triple = "x86_64-unknown-linux-gnu", omp.is_targe
// CHECK-DAG: @_QMtest_0Edata_int_to = global i32 5
// CHECK-DAG: @.offloading.entry_name{{.*}} = internal unnamed_addr constant [22 x i8] c"_QMtest_0Edata_int_to\00"
- // CHECK-DAG: @.offloading.entry._QMtest_0Edata_int_to = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @_QMtest_0Edata_int_to, ptr @.offloading.entry_name{{.*}}, i64 4, i64 0, ptr null }, section "llvm_offload_entries"
+ // CHECK-DAG: @.offloading.entry._QMtest_0Edata_int_to = constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @_QMtest_0Edata_int_to, ptr @.offloading.entry_name{{.*}}, i64 4, i64 0, ptr null }, section "llvm_offload_entries"
// CHECK-DAG: !{{.*}} = !{i32 {{.*}}, !"_QMtest_0Edata_int_to", i32 {{.*}}, i32 {{.*}}}
llvm.mlir.global external @_QMtest_0Edata_int_to() {addr_space = 0 : i32, omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>} : i32 {
%0 = llvm.mlir.constant(5 : i32) : i32
@@ -126,7 +126,7 @@ module attributes {llvm.target_triple = "x86_64-unknown-linux-gnu", omp.is_targe
// CHECK-DAG: @_QMtest_0Edata_int_enter = global i32 5
// CHECK-DAG: @.offloading.entry_name{{.*}} = internal unnamed_addr constant [25 x i8] c"_QMtest_0Edata_int_enter\00"
- // CHECK-DAG: @.offloading.entry._QMtest_0Edata_int_enter = weak constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @_QMtest_0Edata_int_enter, ptr @.offloading.entry_name{{.*}}, i64 4, i64 0, ptr null }, section "llvm_offload_entries"
+ // CHECK-DAG: @.offloading.entry._QMtest_0Edata_int_enter = constant %struct.__tgt_offload_entry { i64 0, i16 1, i16 1, i32 0, ptr @_QMtest_0Edata_int_enter, ptr @.offloading.entry_name{{.*}}, i64 4, i64 0, ptr null }, section "llvm_offload_entries"
// CHECK-DAG: !{{.*}} = !{i32 {{.*}}, !"_QMtest_0Edata_int_enter", i32 {{.*}}, i32 {{.*}}}
llvm.mlir.global external @_QMtest_0Edata_int_enter() {addr_space = 0 : i32, omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>} : i32 {
%0 = llvm.mlir.constant(5 : i32) : i32
diff --git a/mlir/test/Target/LLVMIR/omptarget-declare-target-to-host.mlir b/mlir/test/Target/LLVMIR/omptarget-declare-target-to-host.mlir
index 11be736c820a9..5da3b0f80b565 100644
--- a/mlir/test/Target/LLVMIR/omptarget-declare-target-to-host.mlir
+++ b/mlir/test/Target/LLVMIR/omptarget-declare-target-to-host.mlir
@@ -4,7 +4,7 @@ module attributes {llvm.target_triple = "x86_64-unknown-linux-gnu", omp.is_gpu =
// CHECK-DAG: @_QMtest_0Ezii = global [11 x float] zeroinitializer
// CHECK-DAG: @.offload_sizes = private unnamed_addr constant [2 x i64] [i64 48, i64 0]
// CHECK-DAG: @.offload_maptypes = private unnamed_addr constant [2 x i64] [i64 3, i64 288]
- // CHECK-DAG: @.offloading.entry._QMtest_0Ezii = weak constant %struct.__tgt_offload_entry {{.*}} ptr @_QMtest_0Ezii, {{.*}}, i64 44,{{.*}}
+ // CHECK-DAG: @.offloading.entry._QMtest_0Ezii = constant %struct.__tgt_offload_entry {{.*}} ptr @_QMtest_0Ezii, {{.*}}, i64 44,{{.*}}
llvm.mlir.global external @_QMtest_0Ezii() {addr_space = 0 : i32, omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>} : !llvm.array<11 x f32> {
%0 = llvm.mlir.zero : !llvm.array<11 x f32>
llvm.return %0 : !llvm.array<11 x f32>
>From 7e0d6584fe4cfd482e0e6cf297b8e5a4cedb6e36 Mon Sep 17 00:00:00 2001
From: Johannes Doerfert <johannes at jdoerfert.de>
Date: Fri, 29 May 2026 15:17:30 -0700
Subject: [PATCH 2/2] [OpenMP] Introduce the ompx_name clause for kernel naming
This adds support for the ompx_name clause that allows users to specify
custom kernel names for OpenMP target offloading regions. The clause
accepts a string literal and overrides the default compiler-generated
kernel names.
Example usage:
#pragma omp target ompx_name("my_kernel")
{ ... }
Kernel names need to be unique or they are diagnosed at compile or link
time as errors.
Co-Authored-By: Claude (claude-sonnet-4.5) <noreply at anthropic.com>
---
clang/include/clang/AST/OpenMPClause.h | 33 ++++++++
clang/include/clang/AST/RecursiveASTVisitor.h | 9 ++-
.../clang/Basic/DiagnosticSemaKinds.td | 6 ++
clang/include/clang/Sema/SemaOpenMP.h | 15 ++++
clang/lib/AST/OpenMPClause.cpp | 6 ++
clang/lib/AST/StmtProfile.cpp | 5 ++
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 3 +
clang/lib/Parse/ParseOpenMP.cpp | 1 +
clang/lib/Sema/SemaOpenMP.cpp | 36 +++++++++
clang/lib/Sema/TreeTransform.h | 20 +++++
clang/lib/Serialization/ASTReader.cpp | 8 ++
clang/lib/Serialization/ASTWriter.cpp | 5 ++
clang/test/OpenMP/ompx_name_codegen.cpp | 53 +++++++++++++
.../test/OpenMP/ompx_name_messages_errors.cpp | 62 +++++++++++++++
clang/tools/libclang/CIndex.cpp | 1 +
flang/include/flang/Lower/OpenMP/Clauses.h | 1 +
flang/lib/Semantics/check-omp-structure.cpp | 1 +
llvm/include/llvm/Frontend/OpenMP/ClauseT.h | 7 +-
llvm/include/llvm/Frontend/OpenMP/OMP.td | 17 ++++
.../llvm/Frontend/OpenMP/OMPIRBuilder.h | 13 ++--
llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 12 ++-
offload/test/offloading/ompx_name.c | 78 +++++++++++++++++++
.../offloading/ompx_name_duplicate_link.c | 40 ++++++++++
23 files changed, 422 insertions(+), 10 deletions(-)
create mode 100644 clang/test/OpenMP/ompx_name_codegen.cpp
create mode 100644 clang/test/OpenMP/ompx_name_messages_errors.cpp
create mode 100644 offload/test/offloading/ompx_name.c
create mode 100644 offload/test/offloading/ompx_name_duplicate_link.c
diff --git a/clang/include/clang/AST/OpenMPClause.h b/clang/include/clang/AST/OpenMPClause.h
index 8ceafc4669297..9efef195113fb 100644
--- a/clang/include/clang/AST/OpenMPClause.h
+++ b/clang/include/clang/AST/OpenMPClause.h
@@ -10475,6 +10475,39 @@ class OMPXBareClause : public OMPNoChildClause<llvm::omp::OMPC_ompx_bare> {
OMPXBareClause() = default;
};
+/// This represents the 'ompx_name' clause in the '#pragma omp target'
+/// directive.
+///
+/// \code
+/// #pragma omp target ompx_name("foo")
+/// \endcode
+/// In this example directive '#pragma omp target' has simple 'ompx_name'
+/// clause with the name "foo".
+class OMPXNameClause final
+ : public OMPOneStmtClause<llvm::omp::OMPC_ompx_name, OMPClause> {
+ friend class OMPClauseReader;
+
+ /// Set name.
+ void setName(Expr *A) { setStmt(A); }
+
+public:
+ /// Build 'ompx_name' clause with the given name.
+ ///
+ /// \param A Name.
+ /// \param StartLoc Starting location of the clause.
+ /// \param LParenLoc Location of '('.
+ /// \param EndLoc Ending location of the clause.
+ OMPXNameClause(Expr *A, SourceLocation StartLoc, SourceLocation LParenLoc,
+ SourceLocation EndLoc)
+ : OMPOneStmtClause(A, StartLoc, LParenLoc, EndLoc) {}
+
+ /// Build an empty clause.
+ OMPXNameClause() : OMPOneStmtClause() {}
+
+ /// Returns name.
+ Expr *getName() const { return getStmtAs<Expr>(); }
+};
+
} // namespace clang
#endif // LLVM_CLANG_AST_OPENMPCLAUSE_H
diff --git a/clang/include/clang/AST/RecursiveASTVisitor.h b/clang/include/clang/AST/RecursiveASTVisitor.h
index 529d657fc01f5..549b23de75b31 100644
--- a/clang/include/clang/AST/RecursiveASTVisitor.h
+++ b/clang/include/clang/AST/RecursiveASTVisitor.h
@@ -3483,7 +3483,14 @@ bool RecursiveASTVisitor<Derived>::VisitOMPAllocatorClause(
}
template <typename Derived>
-bool RecursiveASTVisitor<Derived>::VisitOMPAllocateClause(OMPAllocateClause *C) {
+bool RecursiveASTVisitor<Derived>::VisitOMPXNameClause(OMPXNameClause *C) {
+ TRY_TO(TraverseStmt(C->getName()));
+ return true;
+}
+
+template <typename Derived>
+bool RecursiveASTVisitor<Derived>::VisitOMPAllocateClause(
+ OMPAllocateClause *C) {
TRY_TO(TraverseStmt(C->getAllocator()));
TRY_TO(VisitOMPClauseList(C));
return true;
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index f84cd8dca6d4c..5fdd384c8aaac 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12044,6 +12044,12 @@ let CategoryName = "OpenMP Issue" in {
// OpenMP support.
def err_omp_expected_var_arg : Error<
"%0 is not a global variable, static local variable or static data member">;
+def err_ompx_name_argument_not_string : Error<
+ "argument to 'ompx_name' clause must be a string literal">;
+def err_ompx_name_duplicate : Error<
+ "OpenMP target kernel name '%0' is used more than once in this translation unit">;
+def note_ompx_name_previous : Note<
+ "previous use of this kernel name is here">;
def err_omp_expected_var_arg_suggest : Error<
"%0 is not a global variable, static local variable or static data member; "
"did you mean %1">;
diff --git a/clang/include/clang/Sema/SemaOpenMP.h b/clang/include/clang/Sema/SemaOpenMP.h
index 3621ce96b8724..1009001c557b9 100644
--- a/clang/include/clang/Sema/SemaOpenMP.h
+++ b/clang/include/clang/Sema/SemaOpenMP.h
@@ -29,6 +29,8 @@
#include "clang/Sema/Ownership.h"
#include "clang/Sema/SemaBase.h"
#include "llvm/ADT/DenseMap.h"
+#include "llvm/ADT/DenseSet.h"
+#include "llvm/ADT/StringMap.h"
#include "llvm/Frontend/OpenMP/OMP.h.inc"
#include "llvm/Frontend/OpenMP/OMPConstants.h"
#include <optional>
@@ -1180,6 +1182,11 @@ class SemaOpenMP : public SemaBase {
SourceLocation LParenLoc,
SourceLocation EndLoc);
+ /// Called on well-formed 'ompx_name' clause.
+ OMPClause *ActOnOpenMPOmpxNameClause(Expr *Name, SourceLocation StartLoc,
+ SourceLocation LParenLoc,
+ SourceLocation EndLoc);
+
/// Data used for processing a list of variables in OpenMP clauses.
struct OpenMPVarListDataTy final {
Expr *DepModOrTailExpr = nullptr;
@@ -1501,6 +1508,14 @@ class SemaOpenMP : public SemaBase {
private:
void *VarDataSharingAttributesStack;
+ /// User-provided target kernel names from 'ompx_name' clauses in this
+ /// translation unit, keyed to their first source location.
+ llvm::StringMap<SourceLocation> OMPKernelNames;
+
+ /// Source locations for duplicate kernel names that have already been
+ /// diagnosed. This prevents repeated diagnostics during template transforms.
+ llvm::DenseSet<unsigned> DiagnosedOMPKernelNameLocs;
+
/// Number of nested '#pragma omp declare target' directives.
SmallVector<DeclareTargetContextInfo, 4> DeclareTargetNesting;
diff --git a/clang/lib/AST/OpenMPClause.cpp b/clang/lib/AST/OpenMPClause.cpp
index ed00e80144c25..d1c2e1cdeedf9 100644
--- a/clang/lib/AST/OpenMPClause.cpp
+++ b/clang/lib/AST/OpenMPClause.cpp
@@ -2100,6 +2100,12 @@ void OMPClausePrinter::VisitOMPAllocatorClause(OMPAllocatorClause *Node) {
OS << ")";
}
+void OMPClausePrinter::VisitOMPXNameClause(OMPXNameClause *Node) {
+ OS << "ompx_name(";
+ Node->getName()->printPretty(OS, nullptr, Policy, 0);
+ OS << ")";
+}
+
void OMPClausePrinter::VisitOMPCollapseClause(OMPCollapseClause *Node) {
OS << "collapse(";
Node->getNumForLoops()->printPretty(OS, nullptr, Policy, 0);
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index 22e2cc56bc700..ff679d416e1a9 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -533,6 +533,11 @@ void OMPClauseProfiler::VisitOMPAllocatorClause(const OMPAllocatorClause *C) {
Profiler->VisitStmt(C->getAllocator());
}
+void OMPClauseProfiler::VisitOMPXNameClause(const OMPXNameClause *C) {
+ if (C->getName())
+ Profiler->VisitStmt(C->getName());
+}
+
void OMPClauseProfiler::VisitOMPCollapseClause(const OMPCollapseClause *C) {
if (C->getNumForLoops())
Profiler->VisitStmt(C->getNumForLoops());
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index eb2f92cdbf972..5589025d701bb 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -6388,6 +6388,9 @@ void CGOpenMPRuntime::emitTargetOutlinedFunctionHelper(
llvm::TargetRegionEntryInfo EntryInfo =
getEntryInfoFromPresumedLoc(CGM, OMPBuilder, D.getBeginLoc(), ParentName);
+ if (auto *C = D.getSingleClause<OMPXNameClause>())
+ if (auto *S = dyn_cast<StringLiteral>(C->getName()->IgnoreParenImpCasts()))
+ EntryInfo.UserProvidedName = S->getString().str();
CodeGenFunction CGF(CGM, true);
llvm::OpenMPIRBuilder::FunctionGenCallback &&GenerateOutlinedFunction =
diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
index 58d6df302d1a7..3cc2c1d6f168f 100644
--- a/clang/lib/Parse/ParseOpenMP.cpp
+++ b/clang/lib/Parse/ParseOpenMP.cpp
@@ -3251,6 +3251,7 @@ OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind,
case OMPC_align:
case OMPC_message:
case OMPC_ompx_dyn_cgroup_mem:
+ case OMPC_ompx_name:
case OMPC_dyn_groupprivate:
case OMPC_transparent:
// OpenMP [2.5, Restrictions]
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 247a4a6ad9271..088f9460831e1 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -16923,6 +16923,9 @@ OMPClause *SemaOpenMP::ActOnOpenMPSingleExprClause(OpenMPClauseKind Kind,
case OMPC_message:
Res = ActOnOpenMPMessageClause(Expr, StartLoc, LParenLoc, EndLoc);
break;
+ case OMPC_ompx_name:
+ Res = ActOnOpenMPOmpxNameClause(Expr, StartLoc, LParenLoc, EndLoc);
+ break;
case OMPC_align:
Res = ActOnOpenMPAlignClause(Expr, StartLoc, LParenLoc, EndLoc);
break;
@@ -18030,6 +18033,39 @@ OMPClause *SemaOpenMP::ActOnOpenMPMessageClause(Expr *ME,
ME, HelperValStmt, CaptureRegion, StartLoc, LParenLoc, EndLoc);
}
+OMPClause *SemaOpenMP::ActOnOpenMPOmpxNameClause(Expr *Name,
+ SourceLocation StartLoc,
+ SourceLocation LParenLoc,
+ SourceLocation EndLoc) {
+ if (!Name) {
+ Diag(StartLoc, diag::err_ompx_name_argument_not_string);
+ return nullptr;
+ }
+
+ if (!Name->isTypeDependent() && !Name->isValueDependent()) {
+ if (auto *PE = dyn_cast<PredefinedExpr>(Name->IgnoreParenCasts()))
+ Name = PE->getFunctionName();
+ auto *SL = dyn_cast<StringLiteral>(Name->IgnoreParenCasts());
+ if (!SL) {
+ Diag(Name->getExprLoc(), diag::err_ompx_name_argument_not_string);
+ return nullptr;
+ }
+
+ StringRef KernelName = SL->getString();
+ SourceLocation NameLoc = SL->getBeginLoc();
+ auto It = OMPKernelNames.find(KernelName);
+ if (It == OMPKernelNames.end()) {
+ OMPKernelNames[KernelName] = NameLoc;
+ } else if (DiagnosedOMPKernelNameLocs.insert(NameLoc.getRawEncoding())
+ .second) {
+ Diag(NameLoc, diag::err_ompx_name_duplicate) << KernelName;
+ Diag(It->second, diag::note_ompx_name_previous);
+ }
+ }
+ return new (getASTContext())
+ OMPXNameClause(Name, StartLoc, LParenLoc, EndLoc);
+}
+
OMPClause *SemaOpenMP::ActOnOpenMPOrderClause(
OpenMPOrderClauseModifier Modifier, OpenMPOrderClauseKind Kind,
SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation MLoc,
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index 53107c827006d..251538a485566 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -1823,6 +1823,17 @@ class TreeTransform {
EndLoc);
}
+ /// Build a new OpenMP 'ompx_name' clause.
+ ///
+ /// By default, performs semantic analysis to build the new OpenMP clause.
+ /// Subclasses may override this routine to provide different behavior.
+ OMPClause *RebuildOMPXNameClause(Expr *Name, SourceLocation StartLoc,
+ SourceLocation LParenLoc,
+ SourceLocation EndLoc) {
+ return getSema().OpenMP().ActOnOpenMPOmpxNameClause(Name, StartLoc,
+ LParenLoc, EndLoc);
+ }
+
/// Build a new OpenMP 'collapse' clause.
///
/// By default, performs semantic analysis to build the new OpenMP clause.
@@ -10612,6 +10623,15 @@ TreeTransform<Derived>::TransformOMPAllocatorClause(OMPAllocatorClause *C) {
E.get(), C->getBeginLoc(), C->getLParenLoc(), C->getEndLoc());
}
+template <typename Derived>
+OMPClause *TreeTransform<Derived>::TransformOMPXNameClause(OMPXNameClause *C) {
+ ExprResult E = getDerived().TransformExpr(C->getName());
+ if (E.isInvalid())
+ return nullptr;
+ return getDerived().RebuildOMPXNameClause(E.get(), C->getBeginLoc(),
+ C->getLParenLoc(), C->getEndLoc());
+}
+
template <typename Derived>
OMPClause *
TreeTransform<Derived>::TransformOMPSimdlenClause(OMPSimdlenClause *C) {
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index f8a6a38bb9b5c..33071aa17f797 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -11501,6 +11501,9 @@ OMPClause *OMPClauseReader::readClause() {
case llvm::omp::OMPC_allocator:
C = new (Context) OMPAllocatorClause();
break;
+ case llvm::omp::OMPC_ompx_name:
+ C = new (Context) OMPXNameClause();
+ break;
case llvm::omp::OMPC_collapse:
C = new (Context) OMPCollapseClause();
break;
@@ -11937,6 +11940,11 @@ void OMPClauseReader::VisitOMPAllocatorClause(OMPAllocatorClause *C) {
C->setLParenLoc(Record.readSourceLocation());
}
+void OMPClauseReader::VisitOMPXNameClause(OMPXNameClause *C) {
+ C->setName(Record.readExpr());
+ C->setLParenLoc(Record.readSourceLocation());
+}
+
void OMPClauseReader::VisitOMPCollapseClause(OMPCollapseClause *C) {
C->setNumForLoops(Record.readSubExpr());
C->setLParenLoc(Record.readSourceLocation());
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index 21dda6f3733e4..7d80fbe8310e8 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -8119,6 +8119,11 @@ void OMPClauseWriter::VisitOMPAllocatorClause(OMPAllocatorClause *C) {
Record.AddSourceLocation(C->getLParenLoc());
}
+void OMPClauseWriter::VisitOMPXNameClause(OMPXNameClause *C) {
+ Record.AddStmt(C->getName());
+ Record.AddSourceLocation(C->getLParenLoc());
+}
+
void OMPClauseWriter::VisitOMPCollapseClause(OMPCollapseClause *C) {
Record.AddStmt(C->getNumForLoops());
Record.AddSourceLocation(C->getLParenLoc());
diff --git a/clang/test/OpenMP/ompx_name_codegen.cpp b/clang/test/OpenMP/ompx_name_codegen.cpp
new file mode 100644
index 0000000000000..60bc31c84df02
--- /dev/null
+++ b/clang/test/OpenMP/ompx_name_codegen.cpp
@@ -0,0 +1,53 @@
+// Test for ompx_name clause code generation
+//
+// This test verifies that the ompx_name clause correctly sets the kernel name.
+//
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s
+
+// expected-no-diagnostics
+
+#define TO_STR(x) #x
+
+// CHECK: define {{.*}} @my_custom_kernel(
+void test_target() {
+ #pragma omp target ompx_name("my_custom_kernel")
+ {
+ }
+}
+
+// CHECK: define {{.*}} @another_kernel(
+void test_target_parallel() {
+ #pragma omp target parallel ompx_name("another_kernel")
+ {
+ }
+}
+
+// CHECK: define {{.*}} @teams_kernel_name(
+void test_target_teams() {
+ #pragma omp target teams ompx_name("teams_kernel_" "name")
+ {
+ }
+}
+
+// CHECK: define {{.*}} @simd_kernel_name(
+void test_target_simd() {
+ #pragma omp target simd ompx_name("simd_kernel_name")
+ for (int i = 0; i < 10; i++)
+ ;
+}
+
+// CHECK: define {{.*}} @parallel_for_kernel_3(
+void test_target_parallel_for() {
+ #pragma omp target parallel for ompx_name("parallel_for_kernel_" TO_STR(3))
+ for (int i = 0; i < 10; i++)
+ ;
+}
+
+// Verify default kernel name generation without ompx_name
+// CHECK: define {{.*}} @__omp_offloading_{{[0-9a-f]+}}_{{[0-9a-f]+}}_{{.*}}_l{{[0-9]+}}(
+void test_default_name() {
+ #pragma omp target
+ {
+ }
+}
diff --git a/clang/test/OpenMP/ompx_name_messages_errors.cpp b/clang/test/OpenMP/ompx_name_messages_errors.cpp
new file mode 100644
index 0000000000000..d6668a46932fe
--- /dev/null
+++ b/clang/test/OpenMP/ompx_name_messages_errors.cpp
@@ -0,0 +1,62 @@
+// Test for ompx_name clause error checking
+// RUN: %clang_cc1 -std=c++20 -verify -fopenmp %s
+
+static void foo() {
+}
+
+void bar() {
+ int x = 5;
+
+ // expected-error at +1 {{argument to 'ompx_name' clause must be a string literal}}
+ #pragma omp target ompx_name(x)
+ {
+ }
+
+ // expected-error at +1 {{argument to 'ompx_name' clause must be a string literal}}
+ #pragma omp target ompx_name(123)
+ {
+ }
+
+ // This should work - string literal
+ #pragma omp target ompx_name("valid_name")
+ {
+ }
+
+// expected-note at +1 {{previous use of this kernel name is here}}
+#pragma omp target ompx_name("baz")
+ foo();
+
+// expected-error at +1 {{OpenMP target kernel name 'baz' is used more than once in this translation unit}}
+#pragma omp target ompx_name("baz")
+ foo();
+
+#pragma omp target ompx_name(foo) // expected-error {{argument to 'ompx_name' clause must be a string literal}}
+ foo();
+
+#pragma omp target ompx_name("foo", "bar") // expected-error {{expected ')'}} expected-note {{to match this '('}}
+ foo();
+}
+
+consteval const char* getStr() {
+ return "foobar3";
+}
+
+void foobar() {
+// CHECK: define {{.*}} @foobar3(
+ #pragma omp target ompx_name(getStr()) // expected-error {{argument to 'ompx_name' clause must be a string literal}}
+ {}
+}
+
+template<typename T>
+void TTT() {
+// expected-note at +2 {{previous use of this kernel name is here}}
+// expected-error at +1 {{OpenMP target kernel name 'template' is used more than once in this translation unit}}
+ #pragma omp target ompx_name("template")
+ {}
+}
+
+void test2() {
+// expected-note at +1 {{in instantiation of function template specialization 'TTT<int>' requested here}}
+ TTT<int>();
+ TTT<float>();
+}
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index ac2fad38a1348..f96fffd9077ef 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2792,6 +2792,7 @@ void OMPClauseEnqueue::VisitOMPDoacrossClause(const OMPDoacrossClause *C) {
void OMPClauseEnqueue::VisitOMPXAttributeClause(const OMPXAttributeClause *C) {
}
void OMPClauseEnqueue::VisitOMPXBareClause(const OMPXBareClause *C) {}
+void OMPClauseEnqueue::VisitOMPXNameClause(const OMPXNameClause *C) {}
} // namespace
diff --git a/flang/include/flang/Lower/OpenMP/Clauses.h b/flang/include/flang/Lower/OpenMP/Clauses.h
index f334374280c73..cd0a7c3ff5b06 100644
--- a/flang/include/flang/Lower/OpenMP/Clauses.h
+++ b/flang/include/flang/Lower/OpenMP/Clauses.h
@@ -284,6 +284,7 @@ using NumTeams = tomp::clause::NumTeamsT<TypeTy, IdTy, ExprTy>;
using NumThreads = tomp::clause::NumThreadsT<TypeTy, IdTy, ExprTy>;
using OmpxAttribute = tomp::clause::OmpxAttributeT<TypeTy, IdTy, ExprTy>;
using OmpxBare = tomp::clause::OmpxBareT<TypeTy, IdTy, ExprTy>;
+using OmpxName = tomp::clause::OmpxNameT<TypeTy, IdTy, ExprTy>;
using OmpxDynCgroupMem = tomp::clause::OmpxDynCgroupMemT<TypeTy, IdTy, ExprTy>;
using Order = tomp::clause::OrderT<TypeTy, IdTy, ExprTy>;
using Ordered = tomp::clause::OrderedT<TypeTy, IdTy, ExprTy>;
diff --git a/flang/lib/Semantics/check-omp-structure.cpp b/flang/lib/Semantics/check-omp-structure.cpp
index 7c531ae0046ae..098e478729b7d 100644
--- a/flang/lib/Semantics/check-omp-structure.cpp
+++ b/flang/lib/Semantics/check-omp-structure.cpp
@@ -6022,6 +6022,7 @@ CHECK_SIMPLE_CLAUSE(Notinbranch, OMPC_notinbranch)
CHECK_SIMPLE_CLAUSE(Novariants, OMPC_novariants)
CHECK_SIMPLE_CLAUSE(NumTasks, OMPC_num_tasks)
CHECK_SIMPLE_CLAUSE(OmpxAttribute, OMPC_ompx_attribute)
+CHECK_SIMPLE_CLAUSE(OmpxName, OMPC_ompx_name)
CHECK_SIMPLE_CLAUSE(Order, OMPC_order)
CHECK_SIMPLE_CLAUSE(Otherwise, OMPC_otherwise)
CHECK_SIMPLE_CLAUSE(Partial, OMPC_partial)
diff --git a/llvm/include/llvm/Frontend/OpenMP/ClauseT.h b/llvm/include/llvm/Frontend/OpenMP/ClauseT.h
index b2f809e9b51ee..e505ae3f64be8 100644
--- a/llvm/include/llvm/Frontend/OpenMP/ClauseT.h
+++ b/llvm/include/llvm/Frontend/OpenMP/ClauseT.h
@@ -1038,6 +1038,11 @@ struct OmpxBareT {
using EmptyTrait = std::true_type;
};
+template <typename T, typename I, typename E> //
+struct OmpxNameT {
+ using EmptyTrait = std::true_type;
+};
+
template <typename T, typename I, typename E> //
struct OmpxDynCgroupMemT {
using WrapperTrait = std::true_type;
@@ -1389,7 +1394,7 @@ struct WriteT {
template <typename T, typename I, typename E>
using ExtensionClausesT =
std::variant<OmpxAttributeT<T, I, E>, OmpxBareT<T, I, E>,
- OmpxDynCgroupMemT<T, I, E>>;
+ OmpxDynCgroupMemT<T, I, E>, OmpxNameT<T, I, E>>;
template <typename T, typename I, typename E>
using EmptyClausesT = std::variant<
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMP.td b/llvm/include/llvm/Frontend/OpenMP/OMP.td
index e1e66df72dfc5..08585979ed318 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMP.td
+++ b/llvm/include/llvm/Frontend/OpenMP/OMP.td
@@ -426,6 +426,9 @@ def OMPC_OMPX_DynCGroupMem : Clause<[Spelling<"ompx_dyn_cgroup_mem">]> {
let clangClass = "OMPXDynCGroupMemClause";
let flangClass = "ScalarIntExpr";
}
+def OMPC_OMPX_Name : Clause<[Spelling<"ompx_name">]> {
+ let clangClass = "OMPXNameClause";
+}
def OMP_ORDER_concurrent : EnumVal<"concurrent",1,1> {}
def OMP_ORDER_unknown : EnumVal<"unknown",2,0> { let isDefault = 1; }
def OMPC_Order : Clause<[Spelling<"order">]> {
@@ -1224,6 +1227,7 @@ def OMP_Target : Directive<[Spelling<"target">]> {
VersionedClause<OMPC_If>,
VersionedClause<OMPC_NoWait>,
VersionedClause<OMPC_OMPX_Bare>,
+ VersionedClause<OMPC_OMPX_Name>,
VersionedClause<OMPC_OMPX_DynCGroupMem>,
VersionedClause<OMPC_Replayable, 60>,
VersionedClause<OMPC_ThreadLimit, 51>,
@@ -1707,6 +1711,7 @@ def OMP_target_loop : Directive<[Spelling<"target loop">]> {
VersionedClause<OMPC_Device>,
VersionedClause<OMPC_DefaultMap>,
VersionedClause<OMPC_NoWait>,
+ VersionedClause<OMPC_OMPX_Name>,
];
let leafConstructs = [OMP_Target, OMP_loop];
let category = CA_Executable;
@@ -2195,6 +2200,7 @@ def OMP_TargetParallel : Directive<[Spelling<"target parallel">]> {
VersionedClause<OMPC_ProcBind>,
VersionedClause<OMPC_Severity, 60>,
VersionedClause<OMPC_ThreadLimit, 51>,
+ VersionedClause<OMPC_OMPX_Name>,
];
let leafConstructs = [OMP_Target, OMP_Parallel];
let category = CA_Executable;
@@ -2305,6 +2311,7 @@ def OMP_TargetParallelFor : Directive<[Spelling<"target parallel for">]> {
let allowedOnceClauses = [
VersionedClause<OMPC_OMPX_DynCGroupMem>,
VersionedClause<OMPC_ThreadLimit, 51>,
+ VersionedClause<OMPC_OMPX_Name>,
];
let leafConstructs = [OMP_Target, OMP_Parallel, OMP_For];
let category = CA_Executable;
@@ -2348,6 +2355,7 @@ def OMP_TargetParallelForSimd
let allowedOnceClauses = [
VersionedClause<OMPC_OMPX_DynCGroupMem>,
VersionedClause<OMPC_ThreadLimit, 51>,
+ VersionedClause<OMPC_OMPX_Name>,
];
let leafConstructs = [OMP_Target, OMP_Parallel, OMP_For, OMP_Simd];
let category = CA_Executable;
@@ -2384,6 +2392,7 @@ def OMP_target_parallel_loop : Directive<[Spelling<"target parallel loop">]> {
VersionedClause<OMPC_ProcBind>,
VersionedClause<OMPC_Severity, 60>,
VersionedClause<OMPC_ThreadLimit, 51>,
+ VersionedClause<OMPC_OMPX_Name>,
];
let leafConstructs = [OMP_Target, OMP_Parallel, OMP_loop];
let category = CA_Executable;
@@ -2423,6 +2432,7 @@ def OMP_TargetSimd : Directive<[Spelling<"target simd">]> {
VersionedClause<OMPC_Severity, 60>,
VersionedClause<OMPC_SimdLen>,
VersionedClause<OMPC_ThreadLimit, 51>,
+ VersionedClause<OMPC_OMPX_Name>,
];
let leafConstructs = [OMP_Target, OMP_Simd];
let category = CA_Executable;
@@ -2451,6 +2461,7 @@ def OMP_TargetTeams : Directive<[Spelling<"target teams">]> {
VersionedClause<OMPC_NumTeams>,
VersionedClause<OMPC_OMPX_DynCGroupMem>,
VersionedClause<OMPC_OMPX_Bare>,
+ VersionedClause<OMPC_OMPX_Name>,
VersionedClause<OMPC_ThreadLimit>,
];
let leafConstructs = [OMP_Target, OMP_Teams];
@@ -2485,6 +2496,7 @@ def OMP_TargetTeamsDistribute
VersionedClause<OMPC_OMPX_DynCGroupMem>,
VersionedClause<OMPC_Order, 50>,
VersionedClause<OMPC_ThreadLimit>,
+ VersionedClause<OMPC_OMPX_Name>,
];
let leafConstructs = [OMP_Target, OMP_Teams, OMP_Distribute];
let category = CA_Executable;
@@ -2605,6 +2617,7 @@ def OMP_TargetTeamsDistributeParallelFor
];
let allowedOnceClauses = [
VersionedClause<OMPC_OMPX_DynCGroupMem>,
+ VersionedClause<OMPC_OMPX_Name>,
];
let leafConstructs =
[OMP_Target, OMP_Teams, OMP_Distribute, OMP_Parallel, OMP_For];
@@ -2650,6 +2663,7 @@ def OMP_TargetTeamsDistributeParallelForSimd
];
let allowedOnceClauses = [
VersionedClause<OMPC_OMPX_DynCGroupMem>,
+ VersionedClause<OMPC_OMPX_Name>,
];
let leafConstructs =
[OMP_Target, OMP_Teams, OMP_Distribute, OMP_Parallel, OMP_For, OMP_Simd];
@@ -2690,6 +2704,7 @@ def OMP_TargetTeamsDistributeSimd
VersionedClause<OMPC_SafeLen>,
VersionedClause<OMPC_SimdLen>,
VersionedClause<OMPC_ThreadLimit>,
+ VersionedClause<OMPC_OMPX_Name>,
];
let leafConstructs = [OMP_Target, OMP_Teams, OMP_Distribute, OMP_Simd];
let category = CA_Executable;
@@ -2717,6 +2732,7 @@ def OMP_TargetTeamsWorkdistribute : Directive<[Spelling<"target teams workdistri
VersionedClause<OMPC_NumTeams>,
VersionedClause<OMPC_OMPX_DynCGroupMem>,
VersionedClause<OMPC_OMPX_Bare>,
+ VersionedClause<OMPC_OMPX_Name>,
VersionedClause<OMPC_ThreadLimit>,
];
let leafConstructs = [OMP_Target, OMP_Teams, OMP_Workdistribute];
@@ -2751,6 +2767,7 @@ def OMP_target_teams_loop : Directive<[Spelling<"target teams loop">]> {
VersionedClause<OMPC_OMPX_DynCGroupMem>,
VersionedClause<OMPC_Order>,
VersionedClause<OMPC_ThreadLimit>,
+ VersionedClause<OMPC_OMPX_Name>,
];
let leafConstructs = [OMP_Target, OMP_Teams, OMP_loop];
let category = CA_Executable;
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
index 888870a9dc5c5..c2204459689bd 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
@@ -239,22 +239,25 @@ struct TargetRegionEntryInfo {
unsigned FileID;
unsigned Line;
unsigned Count;
+ std::string UserProvidedName;
TargetRegionEntryInfo() : DeviceID(0), FileID(0), Line(0), Count(0) {}
TargetRegionEntryInfo(StringRef ParentName, unsigned DeviceID,
- unsigned FileID, unsigned Line, unsigned Count = 0)
+ unsigned FileID, unsigned Line, unsigned Count = 0,
+ StringRef UserProvidedName = "")
: ParentName(ParentName), DeviceID(DeviceID), FileID(FileID), Line(Line),
- Count(Count) {}
+ Count(Count), UserProvidedName(UserProvidedName) {}
LLVM_ABI static void
getTargetRegionEntryFnName(SmallVectorImpl<char> &Name, StringRef ParentName,
unsigned DeviceID, unsigned FileID, unsigned Line,
- unsigned Count);
+ unsigned Count, StringRef UserProvidedName = "");
bool operator<(const TargetRegionEntryInfo &RHS) const {
- return std::make_tuple(ParentName, DeviceID, FileID, Line, Count) <
+ return std::make_tuple(ParentName, DeviceID, FileID, Line, Count,
+ UserProvidedName) <
std::make_tuple(RHS.ParentName, RHS.DeviceID, RHS.FileID, RHS.Line,
- RHS.Count);
+ RHS.Count, RHS.UserProvidedName);
}
};
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
index d1de4c7bd03f8..b36f2c73a58e7 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -11911,7 +11911,8 @@ void OpenMPIRBuilder::createOffloadEntriesAndInfoMetadata(
}
createOffloadEntry(CE->getID(), CE->getAddress(),
/*Size=*/0, CE->getFlags(),
- GlobalValue::ExternalLinkage);
+ GlobalValue::ExternalLinkage,
+ E.second.UserProvidedName);
} else if (const auto *CE = dyn_cast<
OffloadEntriesInfoManager::OffloadEntryInfoDeviceGlobalVar>(
E.first)) {
@@ -11995,7 +11996,12 @@ void OpenMPIRBuilder::createOffloadEntriesAndInfoMetadata(
void TargetRegionEntryInfo::getTargetRegionEntryFnName(
SmallVectorImpl<char> &Name, StringRef ParentName, unsigned DeviceID,
- unsigned FileID, unsigned Line, unsigned Count) {
+ unsigned FileID, unsigned Line, unsigned Count,
+ StringRef UserProvidedName) {
+ if (!UserProvidedName.empty()) {
+ Name.append(UserProvidedName.begin(), UserProvidedName.end());
+ return;
+ }
raw_svector_ostream OS(Name);
OS << KernelNamePrefix << llvm::format("%x", DeviceID)
<< llvm::format("_%x_", FileID) << ParentName << "_l" << Line;
@@ -12008,7 +12014,7 @@ void OffloadEntriesInfoManager::getTargetRegionEntryFnName(
unsigned NewCount = getTargetRegionEntryInfoCount(EntryInfo);
TargetRegionEntryInfo::getTargetRegionEntryFnName(
Name, EntryInfo.ParentName, EntryInfo.DeviceID, EntryInfo.FileID,
- EntryInfo.Line, NewCount);
+ EntryInfo.Line, NewCount, EntryInfo.UserProvidedName);
}
TargetRegionEntryInfo
diff --git a/offload/test/offloading/ompx_name.c b/offload/test/offloading/ompx_name.c
new file mode 100644
index 0000000000000..d734b65f17f09
--- /dev/null
+++ b/offload/test/offloading/ompx_name.c
@@ -0,0 +1,78 @@
+// RUN: %libomptarget-compile-generic
+// RUN: env LIBOMPTARGET_INFO=63 %libomptarget-run-generic 2>&1 | \
+// RUN: %fcheck-generic
+//
+// REQUIRES: gpu
+
+#include <stdio.h>
+
+int main() {
+ int result = 0;
+
+// CHECK: PluginInterface device {{[0-9]}} info: Launching kernel
+// CHECK-SAME: my_custom_kernel
+#pragma omp target ompx_name("my_custom_kernel") map(from : result)
+ {
+ result = 42;
+ }
+
+ if (result != 42) {
+ printf("FAIL: result = %d\n", result);
+ return 1;
+ }
+
+ result = 0;
+
+// CHECK: PluginInterface device {{[0-9]}} info: Launching kernel
+// CHECK-SAME: parallel_kernel_name
+#pragma omp target parallel ompx_name("parallel_kernel_name") \
+ map(tofrom : result)
+ {
+#pragma omp atomic
+ result++;
+ }
+
+ if (result == 0) {
+ printf("FAIL: parallel result = %d\n", result);
+ return 1;
+ }
+
+ result = 0;
+
+// CHECK: PluginInterface device {{[0-9]}} info: Launching kernel
+// CHECK-SAME: teams_kernel_name
+#pragma omp target teams ompx_name("teams_kernel_name") map(tofrom : result)
+ {
+#pragma omp atomic
+ result++;
+ }
+
+ if (result == 0) {
+ printf("FAIL: teams result = %d\n", result);
+ return 1;
+ }
+
+ int data[100];
+ for (int i = 0; i < 100; i++)
+ data[i] = 0;
+
+// CHECK: PluginInterface device {{[0-9]}} info: Launching kernel
+// CHECK-SAME: parallel_for_kernel
+#pragma omp target parallel for ompx_name("parallel_for_kernel") \
+ map(tofrom : data[0 : 100])
+ for (int i = 0; i < 100; i++) {
+ data[i] = i;
+ }
+
+ for (int i = 0; i < 100; i++) {
+ if (data[i] != i) {
+ printf("FAIL: data[%d] = %d\n", i, data[i]);
+ return 1;
+ }
+ }
+
+ // CHECK: PASS
+ printf("PASS\n");
+
+ return 0;
+}
diff --git a/offload/test/offloading/ompx_name_duplicate_link.c b/offload/test/offloading/ompx_name_duplicate_link.c
new file mode 100644
index 0000000000000..444416a903deb
--- /dev/null
+++ b/offload/test/offloading/ompx_name_duplicate_link.c
@@ -0,0 +1,40 @@
+// RUN: %libomptarget-compile-generic -DFIRST -c -o %t.first.o
+// RUN: %libomptarget-compile-generic -DSECOND -c -o %t.second.o
+// RUN: not %clang-generic %t.second.o %t.first.o -o %t 2>&1 | \
+// %fcheck-plain-generic %s
+//
+// REQUIRES: gpu
+//
+// CHECK: multiple definition
+
+#include <stdio.h>
+
+#ifdef FIRST
+void first(void) {
+ int x = 0;
+#pragma omp target ompx_name("duplicate_link_kernel") map(tofrom : x)
+ {
+ x = 1;
+ }
+ printf("x: %i\n", x);
+}
+#endif
+
+#ifdef SECOND
+void second(void) {
+ int x = 0;
+#pragma omp target ompx_name("duplicate_link_kernel") map(tofrom : x)
+ {
+ x = 2;
+ }
+ printf("x: %i\n", x);
+}
+
+void first(void);
+
+int main(void) {
+ first();
+ second();
+ return 0;
+}
+#endif
More information about the flang-commits
mailing list