llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-offload
Author: David Pagan (ddpagan)
<details>
<summary>Changes</summary>
Implement code generation for the OpenMP 6.0 declare_target 'local' clause,
which creates device-only variables with per-device static storage.
A 'local' variable exists in the device image with its static initializer and
is always accessed directly by device code. This is the same as 'to'/'enter'
without unified shared memory, except that no offload entry is registered.
Using 'device_type(nohost)' with 'local' is not yet supported. Sema generates a
warning and converts it to 'device_type(any)'.
Testing:
- Updated tests:
clang/test/OpenMP/declare_target_messages.cpp
clang/test/OpenMP/declare_target_ast_print.cpp
- New tests:
clang/test/OpenMP/declare_target_local_codegen.cpp
clang/test/OpenMP/declare_target_local_usm_codegen.cpp
offload/test/offloading/declare_target_local.cpp
---
Patch is 47.57 KiB, truncated to 20.00 KiB below, full version:
https://github.com/llvm/llvm-project/pull/196431.diff
10 Files Affected:
- (modified) clang/include/clang/Basic/DiagnosticSemaKinds.td (+4-5)
- (modified) clang/lib/CodeGen/CGExpr.cpp (+9-10)
- (modified) clang/lib/CodeGen/CGOpenMPRuntime.cpp (+16-10)
- (modified) clang/lib/CodeGen/CodeGenModule.cpp (+5-6)
- (modified) clang/lib/Sema/SemaOpenMP.cpp (+10-5)
- (modified) clang/test/OpenMP/declare_target_ast_print.cpp (+6-6)
- (added) clang/test/OpenMP/declare_target_local_codegen.cpp (+430)
- (added) clang/test/OpenMP/declare_target_local_usm_codegen.cpp (+52)
- (modified) clang/test/OpenMP/declare_target_messages.cpp (+5-4)
- (added) offload/test/offloading/declare_target_local.cpp (+40)
``````````diff
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td
b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index c69b2ce3648f8..2e4102d9e1741 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12174,6 +12174,10 @@ def err_omp_declare_target_var_in_both_clauses
def err_omp_declare_target_local_host_only
: Error<"'local' clause is incompatible with 'device_type(host)'; "
"local variables exist only on the device">;
+def warn_omp_declare_target_local_nohost
+ : Warning<"'device_type(nohost)' is not yet supported with 'local' clause;
"
+ "treating as 'device_type(any)'">,
+ InGroup<OpenMPTarget>;
def warn_omp_not_in_target_context : Warning<
"declaration is not declared in any declare target region">,
InGroup<OpenMPTarget>;
@@ -12596,11 +12600,6 @@ def err_omp_declare_target_has_local_vars : Error<
def warn_omp_declare_target_after_first_use : Warning<
"declaration marked as declare target after first use, it may lead to
incorrect results">,
InGroup<OpenMPTarget>;
-def warn_omp_declare_target_local_not_implemented
- : Warning<"'local' clause on 'declare_target' directive is not yet fully "
- "implemented; "
- "variable will be treated as 'enter'">,
- InGroup<OpenMPTarget>;
def err_omp_declare_variant_incompat_attributes : Error<
"'#pragma omp declare variant' is not compatible with any target-specific
attributes">;
def warn_omp_declare_variant_score_not_constant
diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp
index 9107553652688..5764b59e538ae 100644
--- a/clang/lib/CodeGen/CGExpr.cpp
+++ b/clang/lib/CodeGen/CGExpr.cpp
@@ -3360,19 +3360,18 @@ static Address
emitDeclTargetVarDeclLValue(CodeGenFunction &CGF,
const VarDecl *VD, QualType T) {
std::optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
- // Return an invalid address if variable is MT_To (or MT_Enter starting with
- // OpenMP 5.2, or MT_Local in OpenMP 6.0) and unified memory is not enabled.
- // For all other cases: MT_Link and MT_To (or MT_Enter/MT_Local) with unified
- // memory, return a valid address.
- if (!Res || ((*Res == OMPDeclareTargetDeclAttr::MT_To ||
- *Res == OMPDeclareTargetDeclAttr::MT_Enter ||
- *Res == OMPDeclareTargetDeclAttr::MT_Local) &&
- !CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory()))
+ // Always return an invalid address for MT_Local, and also for
+ // MT_To/MT_Enter when unified memory is not enabled. These use direct
+ // access (global exists in device image). Otherwise, return a valid
+ // address.
+ if (!Res || *Res == OMPDeclareTargetDeclAttr::MT_Local ||
+ ((*Res == OMPDeclareTargetDeclAttr::MT_To ||
+ *Res == OMPDeclareTargetDeclAttr::MT_Enter) &&
+ !CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory()))
return Address::invalid();
assert(((*Res == OMPDeclareTargetDeclAttr::MT_Link) ||
((*Res == OMPDeclareTargetDeclAttr::MT_To ||
- *Res == OMPDeclareTargetDeclAttr::MT_Enter ||
- *Res == OMPDeclareTargetDeclAttr::MT_Local) &&
+ *Res == OMPDeclareTargetDeclAttr::MT_Enter) &&
CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory())) &&
"Expected link clause OR to clause with unified memory enabled.");
QualType PtrTy = CGF.getContext().getPointerType(VD->getType());
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index a99a257c14a2a..7cdc206aea0c4 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -1529,12 +1529,14 @@ convertCaptureClause(const VarDecl *VD) {
return llvm::OffloadEntriesInfoManager::OMPTargetGlobalVarEntryTo;
break;
case OMPDeclareTargetDeclAttr::MapTypeTy::MT_Enter:
- case OMPDeclareTargetDeclAttr::MapTypeTy::MT_Local:
return llvm::OffloadEntriesInfoManager::OMPTargetGlobalVarEntryEnter;
- break;
case OMPDeclareTargetDeclAttr::MapTypeTy::MT_Link:
return llvm::OffloadEntriesInfoManager::OMPTargetGlobalVarEntryLink;
break;
+ case OMPDeclareTargetDeclAttr::MapTypeTy::MT_Local:
+ // MT_Local variables don't need offload entry (device-local).
+ llvm_unreachable("MT_Local should not reach convertCaptureClause");
+ break;
default:
return llvm::OffloadEntriesInfoManager::OMPTargetGlobalVarEntryNone;
break;
@@ -7983,8 +7985,7 @@ class MappableExprsHandler {
OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD)) {
if ((*Res == OMPDeclareTargetDeclAttr::MT_Link) ||
((*Res == OMPDeclareTargetDeclAttr::MT_To ||
- *Res == OMPDeclareTargetDeclAttr::MT_Enter ||
- *Res == OMPDeclareTargetDeclAttr::MT_Local) &&
+ *Res == OMPDeclareTargetDeclAttr::MT_Enter) &&
CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory())) {
RequiresReference = true;
BP = CGF.CGM.getOpenMPRuntime().getAddrOfDeclareTargetVar(VD);
@@ -11326,8 +11327,7 @@ bool
CGOpenMPRuntime::emitTargetGlobalVariable(GlobalDecl GD) {
cast<VarDecl>(GD.getDecl()));
if (!Res || *Res == OMPDeclareTargetDeclAttr::MT_Link ||
((*Res == OMPDeclareTargetDeclAttr::MT_To ||
- *Res == OMPDeclareTargetDeclAttr::MT_Enter ||
- *Res == OMPDeclareTargetDeclAttr::MT_Local) &&
+ *Res == OMPDeclareTargetDeclAttr::MT_Enter) &&
HasRequiresUnifiedSharedMemory)) {
DeferredGlobalVariables.insert(cast<VarDecl>(GD.getDecl()));
return true;
@@ -11350,6 +11350,11 @@ void
CGOpenMPRuntime::registerTargetGlobalVariable(const VarDecl *VD,
VD->hasExternalStorage())
return;
+ // MT_Local variables use direct access with no host-device mapping.
+ // No offload entry needed — the device global keeps its own initializer.
+ if (Res && *Res == OMPDeclareTargetDeclAttr::MT_Local)
+ return;
+
if (!Res) {
if (CGM.getLangOpts().OpenMPIsTargetDevice) {
// Register non-target variables being emitted in device code (debug info
@@ -11396,10 +11401,11 @@ void CGOpenMPRuntime::emitDeferredTargetDecls() const
{
OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
if (!Res)
continue;
- if ((*Res == OMPDeclareTargetDeclAttr::MT_To ||
- *Res == OMPDeclareTargetDeclAttr::MT_Enter ||
- *Res == OMPDeclareTargetDeclAttr::MT_Local) &&
- !HasRequiresUnifiedSharedMemory) {
+ // MT_Local and MT_To/MT_Enter without USM are always emitted.
+ if (*Res == OMPDeclareTargetDeclAttr::MT_Local ||
+ ((*Res == OMPDeclareTargetDeclAttr::MT_To ||
+ *Res == OMPDeclareTargetDeclAttr::MT_Enter) &&
+ !HasRequiresUnifiedSharedMemory)) {
CGM.EmitGlobal(VD);
} else {
assert((*Res == OMPDeclareTargetDeclAttr::MT_Link ||
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp
b/clang/lib/CodeGen/CodeGenModule.cpp
index e88fed1296667..5b26332765fab 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -4465,16 +4465,15 @@ void CodeGenModule::EmitGlobal(GlobalDecl GD) {
bool UnifiedMemoryEnabled =
getOpenMPRuntime().hasRequiresUnifiedSharedMemory();
- if ((*Res == OMPDeclareTargetDeclAttr::MT_To ||
- *Res == OMPDeclareTargetDeclAttr::MT_Enter ||
- *Res == OMPDeclareTargetDeclAttr::MT_Local) &&
- !UnifiedMemoryEnabled) {
+ if (*Res == OMPDeclareTargetDeclAttr::MT_Local ||
+ ((*Res == OMPDeclareTargetDeclAttr::MT_To ||
+ *Res == OMPDeclareTargetDeclAttr::MT_Enter) &&
+ !UnifiedMemoryEnabled)) {
(void)GetAddrOfGlobalVar(VD);
} else {
assert(((*Res == OMPDeclareTargetDeclAttr::MT_Link) ||
((*Res == OMPDeclareTargetDeclAttr::MT_To ||
- *Res == OMPDeclareTargetDeclAttr::MT_Enter ||
- *Res == OMPDeclareTargetDeclAttr::MT_Local) &&
+ *Res == OMPDeclareTargetDeclAttr::MT_Enter) &&
UnifiedMemoryEnabled)) &&
"Link clause or to clause with unified memory expected.");
(void)getOpenMPRuntime().getAddrOfDeclareTargetVar(VD);
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 53ded7a5e177e..d6f6bc919a31b 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -24859,13 +24859,18 @@ void SemaOpenMP::ActOnOpenMPDeclareTargetName(
if (!IndirectE)
IsIndirect = true;
}
- // FIXME: 'local' clause is not yet implemented in CodeGen. For now, it is
- // treated as 'enter'. For host compilation, 'local' is a no-op.
+ // FIXME: 'local' with 'device_type(nohost)' is not yet fully supported
+ // in codegen. Treat as 'device_type(any)' for now. The variable will
+ // exist on both host and device, but the host copy is unused.
+ auto DT = DTCI.DT;
if (MT == OMPDeclareTargetDeclAttr::MT_Local &&
- getLangOpts().OpenMPIsTargetDevice)
- Diag(Loc, diag::warn_omp_declare_target_local_not_implemented);
+ DT == OMPDeclareTargetDeclAttr::DT_NoHost) {
+ Diag(Loc, diag::warn_omp_declare_target_local_nohost);
+ DT = OMPDeclareTargetDeclAttr::DT_Any;
+ }
+
auto *A = OMPDeclareTargetDeclAttr::CreateImplicit(
- getASTContext(), MT, DTCI.DT, IndirectE, IsIndirect, Level,
+ getASTContext(), MT, DT, IndirectE, IsIndirect, Level,
SourceRange(Loc, Loc));
ND->addAttr(A);
if (ASTMutationListener *ML = getASTContext().getASTMutationListener())
diff --git a/clang/test/OpenMP/declare_target_ast_print.cpp
b/clang/test/OpenMP/declare_target_ast_print.cpp
index 3ebe261cf79f0..7b63c15dd455e 100644
--- a/clang/test/OpenMP/declare_target_ast_print.cpp
+++ b/clang/test/OpenMP/declare_target_ast_print.cpp
@@ -4,7 +4,7 @@
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -I %S/Inputs
-ast-print %s | FileCheck %s --check-prefix=CHECK --check-prefix=OMP50
// RUN: %clang_cc1 -verify -fopenmp -I %S/Inputs -ast-print %s | FileCheck %s
--check-prefix=CHECK --check-prefix=OMP51
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=52 -I %S/Inputs
-ast-print %s | FileCheck %s --check-prefix=CHECK --check-prefix=OMP52
-// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=60 -I %S/Inputs
-ast-print %s | FileCheck %s --check-prefix=CHECK --check-prefix=OMP60
+// RUN: %clang_cc1 -verify=omp60 -fopenmp -fopenmp-version=60 -I %S/Inputs
-ast-print %s | FileCheck %s --check-prefix=CHECK --check-prefix=OMP60
// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -I %S/Inputs
-emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -std=c++11 -include-pch %t -I
%S/Inputs -verify %s -ast-print | FileCheck %s --check-prefix=CHECK
--check-prefix=OMP50
@@ -133,18 +133,18 @@ int l1;
// OMP60: #pragma omp end declare target
int l2;
-#pragma omp declare target device_type(nohost) local(l2)
-// OMP60: #pragma omp declare target device_type(nohost) local
+#pragma omp declare target device_type(nohost) local(l2) // omp60-warning
{{'device_type(nohost)' is not yet supported with 'local' clause; treating as
'device_type(any)'}}
+// OMP60: #pragma omp declare target local
// OMP60: int l2;
// OMP60: #pragma omp end declare target
int l3;
int a = 0;
-#pragma omp declare target local(l3) device_type(nohost) local(a)
-// OMP60: #pragma omp declare target device_type(nohost) local
+#pragma omp declare target local(l3) device_type(nohost) local(a) //
omp60-warning 2 {{'device_type(nohost)' is not yet supported with 'local'
clause; treating as 'device_type(any)'}}
+// OMP60: #pragma omp declare target local
// OMP60: int l3;
// OMP60: #pragma omp end declare target
-// OMP60: #pragma omp declare target device_type(nohost) local
+// OMP60: #pragma omp declare target local
// OMP60: int a = 0;
// OMP60: #pragma omp end declare target
diff --git a/clang/test/OpenMP/declare_target_local_codegen.cpp
b/clang/test/OpenMP/declare_target_local_codegen.cpp
new file mode 100644
index 0000000000000..b82e8b3bba9ff
--- /dev/null
+++ b/clang/test/OpenMP/declare_target_local_codegen.cpp
@@ -0,0 +1,430 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
UTC_ARGS: --function-signature --include-generated-funcs --replace-value-regex
"__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]"
"pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _
+// RUN: %clang_cc1 -verify=omp60 -fopenmp -fopenmp-version=60 -x c++ -triple
powerpc64le-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o
- | FileCheck %s --check-prefix=HOST
+// RUN: %clang_cc1 -verify=omp60 -fopenmp -fopenmp-version=60 -x c++ -triple
powerpc64le-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s
-o %t-host.bc
+// RUN: %clang_cc1 -verify=omp60 -fopenmp -fopenmp-version=60 -x c++ -triple
amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-target-device
-fvisibility=protected -fopenmp-host-ir-file-path %t-host.bc -o - | FileCheck
%s --check-prefix=DEVICE
+// RUN: %clang_cc1 -verify=omp60 -fopenmp -fopenmp-version=60 -x c++ -triple
amdgcn-amd-amdhsa %s -fopenmp-is-target-device -fvisibility=protected
-fopenmp-host-ir-file-path %t-host.bc -emit-pch -o %t
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=60 -x c++ -triple
amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-target-device
-fvisibility=protected -fopenmp-host-ir-file-path %t-host.bc -include-pch %t -o
- | FileCheck %s --check-prefix=DEVICE
+
+#ifndef HEADER
+#define HEADER
+
+// ---------------------------------------------------------------------------
+// Explicit local clause (default device_type is 'any')
+// ---------------------------------------------------------------------------
+int local_scalar;
+#pragma omp declare target local(local_scalar)
+
+int local_array[64];
+#pragma omp declare target local(local_array)
+
+// ---------------------------------------------------------------------------
+// local + device_type(nohost)
+// ---------------------------------------------------------------------------
+int local_nohost_var;
+#pragma omp declare target local(local_nohost_var) device_type(nohost) //
omp60-warning {{'device_type(nohost)' is not yet supported with 'local' clause;
treating as 'device_type(any)'}}
+
+double local_nohost_arr[32];
+#pragma omp declare target local(local_nohost_arr) device_type(nohost) //
omp60-warning {{'device_type(nohost)' is not yet supported with 'local' clause;
treating as 'device_type(any)'}}
+
+// ---------------------------------------------------------------------------
+// Template with local variable
+// ---------------------------------------------------------------------------
+template <typename T>
+struct LocalStorage {
+ static T value;
+};
+
+template <typename T>
+T LocalStorage<T>::value;
+
+#pragma omp declare target local(LocalStorage<int>::value)
+#pragma omp declare target local(LocalStorage<double>::value)
+
+#pragma omp begin declare target
+template <typename T>
+T read_local_storage() {
+ return LocalStorage<T>::value;
+}
+#pragma omp end declare target
+
+// ---------------------------------------------------------------------------
+// Non-template static data member with local
+// ---------------------------------------------------------------------------
+struct PlainStruct {
+ static int s_member;
+};
+int PlainStruct::s_member;
+#pragma omp declare target local(PlainStruct::s_member)
+
+// ---------------------------------------------------------------------------
+// Initialized local variable
+// ---------------------------------------------------------------------------
+int local_init_var = 42;
+#pragma omp declare target local(local_init_var)
+
+// ---------------------------------------------------------------------------
+// Use local variables in a target region
+// ---------------------------------------------------------------------------
+int use_local_vars() {
+ int result = 0;
+ #pragma omp target map(from: result)
+ {
+ local_scalar = 42;
+ local_array[0] = 1;
+ LocalStorage<int>::value = 100;
+ result = local_scalar + local_array[0]
+ + read_local_storage<int>();
+ }
+ return result;
+}
+
+// ---------------------------------------------------------------------------
+// Use nohost local variables in a target region
+// ---------------------------------------------------------------------------
+int use_nohost_local_vars() {
+ int result = 0;
+ #pragma omp target map(from: result)
+ {
+ local_nohost_var = 7;
+ result = local_nohost_var;
+ }
+ return result;
+}
+
+// ---------------------------------------------------------------------------
+// Use static data member, initialized var, and static local in target region
+// ---------------------------------------------------------------------------
+int use_new_local_vars() {
+ int result = 0;
+ #pragma omp target map(from: result)
+ {
+ PlainStruct::s_member = 55;
+ local_init_var = 77;
+ result = PlainStruct::s_member + local_init_var;
+ }
+ return result;
+}
+
+#endif
+// HOST-LABEL: define {{[^@]+}}@_Z14use_local_varsv
+// HOST-SAME: () #[[ATTR0:[0-9]+]] {
+// HOST-NEXT: entry:
+// HOST-NEXT: [[RESULT:%.*]] = alloca i32, align 4
+// HOST-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [2 x ptr], align 8
+// HOST-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [2 x ptr], align 8
+// HOST-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [2 x ptr], align 8
+// HOST-NEXT: [[KERNEL_ARGS:%.*]] = alloca
[[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8
+// HOST-NEXT: store i32 0, ptr [[RESULT]], align 4
+// HOST-NEXT: [[TMP0:%.*]] = getelementptr inbounds [2 x ptr], ptr
[[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// HOST-NEXT: store ptr [[RESULT]], ptr [[TMP0]], align 8
+// HOST-NEXT: [[TMP1:%.*]] = getelementptr inbounds [2 x ptr], ptr
[[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// HOST-NEXT: store ptr [[RESULT]], ptr [[TMP1]], align 8
+// HOST-NEXT: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr
[[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
+// HOST-NEXT: store ptr null, ptr [[TMP2]], align 8
+// HOST-NEXT: [[TMP3:%.*]] = getelementptr inbounds [2 x ptr], ptr
[[DOTOFFLOAD_BASEPTRS]], i32 0, i32 1
+// HOST-NEXT: store ptr null, ptr [[TMP3]], align 8
+// HOST-NEXT: [[TMP4:%.*]] = getelementptr inbounds [2 x ptr], ptr
[[DOTOFFLOAD_PTRS]], i32 0, i32 1
+// HOST-NEXT: store ptr null, ptr [[TMP4]], align 8
+// HOST-NEXT: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr
[[DOTOFFLOAD_MAPPERS]], i64 0, i64 1
+// HOST-NEXT: store ptr null, ptr [[TMP5]], align 8
+// HOST-NEXT: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr
[[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
+// HOST-NEXT: [[TMP7:%.*]] = getelementptr inbounds [2 x ptr], ptr
[[DOTOFFLOAD_PTRS]], i32 0, i32 0
+// HOST-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0
+// HOST-NEXT: store i32 4, ptr [[TMP8]], align 4
+// HOST-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1
+// HOST-NEXT: store i32 2, ptr [[TMP9]], align 4
+// HOST-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2
+// HOST-NEXT: store ptr [[TMP6]], ptr [[TMP10]], align 8
+// HOST-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3
+// HOST-NEXT: store ptr [[TMP7]], ptr [[TMP11]], align 8
+// HOST-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4
+// HOST-NEXT: store ptr @.offload_sizes, ptr [[TMP12]], align 8
+// HOST-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5
+// HOST-NEXT: store ptr @.offload_maptypes, ptr [[TMP13]], align 8
+// HOST-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw
[[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6
+// HOST-NEXT: store ptr null, ptr [[TMP14]], align 8
+// HOST-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRU...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/196431
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
