jdenny created this revision.
jdenny added reviewers: grokos, ABataev, jdoerfert.
Herald added subscribers: llvm-commits, openmp-commits, cfe-commits, sstefan1,
guansong, yaxunl.
Herald added projects: clang, OpenMP, LLVM.
Without this patch, the following example fails but shouldn't
according to my read of OpenMP TR8:
#pragma omp target enter data map(alloc:i)
#pragma omp target data map(present, alloc: i)
{
#pragma omp target exit data map(delete:i)
} // fails presence check here
OpenMP TR8 sec. 2.22.7.1 "map Clause", p. 321, L23-26 states:
> If the map clause appears on a target, target data, target enter
> data or target exit data construct with a present map-type-modifier
> then on entry to the region if the corresponding list item does not
> appear in the device data environment an error occurs and the
> program terminates.
I see no corresponding statement about the exit from a region. Thus,
the `present` modifier should:
1. Check for presence upon entry into a `target exit data` construct.
2. Should not check for presence upon exit from a `target data` region, as in
the above example.
The problem is that Clang calls the same set of
`__tgt_target_data_end_*` functions for these two cases, making them
indistinguishable in the runtime where the presence check is
implemented. To fix that, this patch changes Clang to generate calls
to a new set of runtime functions, `__tgt_target_exit_data_*`, for the
case of `target exit data`.
For symmetry, this patch makes a similar change for `target enter
data`, but that change isn't required for the above fix.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D84422
Files:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/test/OpenMP/declare_mapper_codegen.cpp
clang/test/OpenMP/target_enter_data_codegen.cpp
clang/test/OpenMP/target_enter_data_depend_codegen.cpp
clang/test/OpenMP/target_exit_data_codegen.cpp
clang/test/OpenMP/target_exit_data_depend_codegen.cpp
clang/test/OpenMP/target_map_member_expr_array_section_codegen.cpp
llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
openmp/libomptarget/include/omptarget.h
openmp/libomptarget/src/exports
openmp/libomptarget/src/interface.cpp
openmp/libomptarget/src/omptarget.cpp
openmp/libomptarget/src/private.h
openmp/libomptarget/test/mapping/present/target_data_at_exit.c
Index: openmp/libomptarget/test/mapping/present/target_data_at_exit.c
===================================================================
--- /dev/null
+++ openmp/libomptarget/test/mapping/present/target_data_at_exit.c
@@ -0,0 +1,37 @@
+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 \
+// RUN: | %fcheck-aarch64-unknown-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu -fopenmp-version=51
+// RUN: %libomptarget-run-x86_64-pc-linux-gnu 2>&1 \
+// RUN: | %fcheck-x86_64-pc-linux-gnu
+
+#include <stdio.h>
+
+int main() {
+ int i;
+
+#pragma omp target enter data map(alloc:i)
+
+ // i isn't present at the end of the target data region, but the "present"
+ // modifier is only checked at the beginning of a region.
+#pragma omp target data map(present, alloc: i)
+ {
+#pragma omp target exit data map(delete:i)
+ }
+
+ // CHECK-NOT: Libomptarget
+ // CHECK: success
+ // CHECK-NOT: Libomptarget
+ fprintf(stderr, "success\n");
+
+ return 0;
+}
Index: openmp/libomptarget/src/private.h
===================================================================
--- openmp/libomptarget/src/private.h
+++ openmp/libomptarget/src/private.h
@@ -24,8 +24,8 @@
extern int target_data_end(DeviceTy &Device, int32_t arg_num, void **args_base,
void **args, int64_t *arg_sizes, int64_t *arg_types,
- void **arg_mappers,
- __tgt_async_info *async_info_ptr);
+ void **arg_mappers, __tgt_async_info *async_info_ptr,
+ bool for_exit_data);
extern int target_data_update(DeviceTy &Device, int32_t arg_num,
void **args_base, void **args,
Index: openmp/libomptarget/src/omptarget.cpp
===================================================================
--- openmp/libomptarget/src/omptarget.cpp
+++ openmp/libomptarget/src/omptarget.cpp
@@ -421,10 +421,30 @@
return OFFLOAD_SUCCESS;
}
+static int target_data_end_not_for_exit_data(DeviceTy &Device, int32_t arg_num,
+ void **args_base, void **args,
+ int64_t *arg_sizes,
+ int64_t *arg_types,
+ void **arg_mappers,
+ __tgt_async_info *async_info_ptr) {
+ return target_data_end(Device, arg_num, args_base, args, arg_sizes, arg_types,
+ arg_mappers, async_info_ptr, /*for_exit_data=*/false);
+}
+
+static int target_data_end_for_exit_data(DeviceTy &Device, int32_t arg_num,
+ void **args_base, void **args,
+ int64_t *arg_sizes, int64_t *arg_types,
+ void **arg_mappers,
+ __tgt_async_info *async_info_ptr) {
+ return target_data_end(Device, arg_num, args_base, args, arg_sizes, arg_types,
+ arg_mappers, async_info_ptr, /*for_exit_data=*/true);
+}
+
/// Internal function to undo the mapping and retrieve the data from the device.
int target_data_end(DeviceTy &Device, int32_t arg_num, void **args_base,
void **args, int64_t *arg_sizes, int64_t *arg_types,
- void **arg_mappers, __tgt_async_info *async_info_ptr) {
+ void **arg_mappers, __tgt_async_info *async_info_ptr,
+ bool for_exit_data) {
// process each input.
for (int32_t i = arg_num - 1; i >= 0; --i) {
// Ignore private variables and arrays - there is no mapping for them.
@@ -439,8 +459,11 @@
// with new arguments.
DP("Calling target_data_mapper for the %dth argument\n", i);
- int rc = target_data_mapper(Device, args_base[i], args[i], arg_sizes[i],
- arg_types[i], arg_mappers[i], target_data_end);
+ int rc =
+ target_data_mapper(Device, args_base[i], args[i], arg_sizes[i],
+ arg_types[i], arg_mappers[i],
+ for_exit_data ? target_data_end_for_exit_data
+ : target_data_end_not_for_exit_data);
if (rc != OFFLOAD_SUCCESS) {
DP("Call to target_data_end via target_data_mapper for custom mapper"
@@ -483,9 +506,9 @@
if (!TgtPtrBegin && (data_size || HasPresentModifier)) {
DP("Mapping does not exist (%s)\n",
(HasPresentModifier ? "'present' map type modifier" : "ignored"));
- if (HasPresentModifier) {
- // FIXME: This should not be an error on exit from "omp target data",
- // but it should be an error upon entering an "omp target exit data".
+ // If HasPresentModifier, complain if data is not present upon entering an
+ // "omp target exit data" but not upon exiting an "omp target data".
+ if (HasPresentModifier && for_exit_data) {
MESSAGE("device mapping required by 'present' map type modifier does "
"not exist for host address " DPxMOD " (%ld bytes)",
DPxPTR(HstPtrBegin), data_size);
@@ -938,7 +961,8 @@
// Move data from device.
int rt = target_data_end(Device, arg_num, args_base, args, arg_sizes,
- arg_types, arg_mappers, &AsyncInfo);
+ arg_types, arg_mappers, &AsyncInfo,
+ /*for_exit_data=*/false);
if (rt != OFFLOAD_SUCCESS) {
DP("Call to target_data_end failed, abort targe.\n");
return OFFLOAD_FAIL;
Index: openmp/libomptarget/src/interface.cpp
===================================================================
--- openmp/libomptarget/src/interface.cpp
+++ openmp/libomptarget/src/interface.cpp
@@ -172,9 +172,10 @@
arg_types, nullptr);
}
-EXTERN void __tgt_target_data_end_mapper(int64_t device_id, int32_t arg_num,
- void **args_base, void **args, int64_t *arg_sizes, int64_t *arg_types,
- void **arg_mappers) {
+static void target_data_end_mapper(int64_t device_id, int32_t arg_num,
+ void **args_base, void **args,
+ int64_t *arg_sizes, int64_t *arg_types,
+ void **arg_mappers, bool for_exit_data) {
if (IsOffloadDisabled()) return;
DP("Entering data end region with %d mappings\n", arg_num);
@@ -208,19 +209,74 @@
#endif
int rc = target_data_end(Device, arg_num, args_base, args, arg_sizes,
- arg_types, arg_mappers, nullptr);
+ arg_types, arg_mappers, nullptr, for_exit_data);
HandleTargetOutcome(rc == OFFLOAD_SUCCESS);
}
+static void target_data_end_nowait_mapper(
+ int64_t device_id, int32_t arg_num, void **args_base, void **args,
+ int64_t *arg_sizes, int64_t *arg_types, void **arg_mappers, int32_t depNum,
+ void *depList, int32_t noAliasDepNum, void *noAliasDepList,
+ bool for_exit_data) {
+ if (depNum + noAliasDepNum > 0)
+ __kmpc_omp_taskwait(NULL, __kmpc_global_thread_num(NULL));
+
+ target_data_end_mapper(device_id, arg_num, args_base, args, arg_sizes,
+ arg_types, arg_mappers, for_exit_data);
+}
+
+EXTERN void __tgt_target_data_end_mapper(int64_t device_id, int32_t arg_num,
+ void **args_base, void **args,
+ int64_t *arg_sizes, int64_t *arg_types,
+ void **arg_mappers) {
+ target_data_end_mapper(device_id, arg_num, args_base, args, arg_sizes,
+ arg_types, arg_mappers, /*for_exit_data=*/false);
+}
+
EXTERN void __tgt_target_data_end_nowait_mapper(int64_t device_id,
int32_t arg_num, void **args_base, void **args, int64_t *arg_sizes,
int64_t *arg_types, void **arg_mappers, int32_t depNum, void *depList,
int32_t noAliasDepNum, void *noAliasDepList) {
- if (depNum + noAliasDepNum > 0)
- __kmpc_omp_taskwait(NULL, __kmpc_global_thread_num(NULL));
+ target_data_end_nowait_mapper(device_id, arg_num, args_base, args, arg_sizes,
+ arg_types, arg_mappers, depNum, depList,
+ noAliasDepNum, noAliasDepList,
+ /*for_exit_data=*/false);
+}
- __tgt_target_data_end_mapper(device_id, arg_num, args_base, args, arg_sizes,
- arg_types, arg_mappers);
+EXTERN void __tgt_target_enter_data_mapper(int64_t device_id, int32_t arg_num,
+ void **args_base, void **args,
+ int64_t *arg_sizes,
+ int64_t *arg_types,
+ void **arg_mappers) {
+ return __tgt_target_data_begin_mapper(device_id, arg_num, args_base, args,
+ arg_sizes, arg_types, arg_mappers);
+}
+
+EXTERN void __tgt_target_enter_data_nowait_mapper(
+ int64_t device_id, int32_t arg_num, void **args_base, void **args,
+ int64_t *arg_sizes, int64_t *arg_types, void **arg_mappers, int32_t depNum,
+ void *depList, int32_t noAliasDepNum, void *noAliasDepList) {
+ return __tgt_target_data_begin_nowait_mapper(
+ device_id, arg_num, args_base, args, arg_sizes, arg_types, arg_mappers,
+ depNum, depList, noAliasDepNum, noAliasDepList);
+}
+
+EXTERN void __tgt_target_exit_data_mapper(int64_t device_id, int32_t arg_num,
+ void **args_base, void **args,
+ int64_t *arg_sizes,
+ int64_t *arg_types,
+ void **arg_mappers) {
+ return target_data_end_mapper(device_id, arg_num, args_base, args, arg_sizes,
+ arg_types, arg_mappers, /*for_exit_data=*/true);
+}
+
+EXTERN void __tgt_target_exit_data_nowait_mapper(
+ int64_t device_id, int32_t arg_num, void **args_base, void **args,
+ int64_t *arg_sizes, int64_t *arg_types, void **arg_mappers, int32_t depNum,
+ void *depList, int32_t noAliasDepNum, void *noAliasDepList) {
+ return target_data_end_nowait_mapper(
+ device_id, arg_num, args_base, args, arg_sizes, arg_types, arg_mappers,
+ depNum, depList, noAliasDepNum, noAliasDepList, /*for_exit_data=*/true);
}
EXTERN void __tgt_target_data_update(int64_t device_id, int32_t arg_num,
Index: openmp/libomptarget/src/exports
===================================================================
--- openmp/libomptarget/src/exports
+++ openmp/libomptarget/src/exports
@@ -15,11 +15,15 @@
__tgt_target_teams_nowait;
__tgt_target_data_begin_mapper;
__tgt_target_data_end_mapper;
+ __tgt_target_enter_data_mapper;
+ __tgt_target_exit_data_mapper;
__tgt_target_data_update_mapper;
__tgt_target_mapper;
__tgt_target_teams_mapper;
__tgt_target_data_begin_nowait_mapper;
__tgt_target_data_end_nowait_mapper;
+ __tgt_target_enter_data_nowait_mapper;
+ __tgt_target_exit_data_nowait_mapper;
__tgt_target_data_update_nowait_mapper;
__tgt_target_nowait_mapper;
__tgt_target_teams_nowait_mapper;
Index: openmp/libomptarget/include/omptarget.h
===================================================================
--- openmp/libomptarget/include/omptarget.h
+++ openmp/libomptarget/include/omptarget.h
@@ -192,6 +192,24 @@
void *depList, int32_t noAliasDepNum,
void *noAliasDepList);
+void __tgt_target_enter_data_mapper(int64_t device_id, int32_t arg_num,
+ void **args_base, void **args,
+ int64_t *arg_sizes, int64_t *arg_types,
+ void **arg_mappers);
+void __tgt_target_enter_data_nowait_mapper(
+ int64_t device_id, int32_t arg_num, void **args_base, void **args,
+ int64_t *arg_sizes, int64_t *arg_types, void **arg_mappers, int32_t depNum,
+ void *depList, int32_t noAliasDepNum, void *noAliasDepList);
+
+void __tgt_target_exit_data_mapper(int64_t device_id, int32_t arg_num,
+ void **args_base, void **args,
+ int64_t *arg_sizes, int64_t *arg_types,
+ void **arg_mappers);
+void __tgt_target_exit_data_nowait_mapper(
+ int64_t device_id, int32_t arg_num, void **args_base, void **args,
+ int64_t *arg_sizes, int64_t *arg_types, void **arg_mappers, int32_t depNum,
+ void *depList, int32_t noAliasDepNum, void *noAliasDepList);
+
/// passes data to/from the target
void __tgt_target_data_update(int64_t device_id, int32_t arg_num,
void **args_base, void **args, int64_t *arg_sizes,
Index: llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
===================================================================
--- llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
+++ llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
@@ -574,6 +574,14 @@
VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr)
__OMP_RTL(__tgt_target_data_end_nowait_mapper, false, Void, Int64, Int32,
VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr)
+__OMP_RTL(__tgt_target_enter_data_mapper, false, Void, Int64, Int32,
+ VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr)
+__OMP_RTL(__tgt_target_enter_data_nowait_mapper, false, Void, Int64, Int32,
+ VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr)
+__OMP_RTL(__tgt_target_exit_data_mapper, false, Void, Int64, Int32,
+ VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr)
+__OMP_RTL(__tgt_target_exit_data_nowait_mapper, false, Void, Int64, Int32,
+ VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr)
__OMP_RTL(__tgt_target_data_update_mapper, false, Void, Int64, Int32,
VoidPtrPtr, VoidPtrPtr, Int64Ptr, Int64Ptr, VoidPtrPtr)
__OMP_RTL(__tgt_target_data_update_nowait_mapper, false, Void, Int64, Int32,
@@ -996,6 +1004,12 @@
__OMP_RTL_ATTRS(__tgt_target_data_end_mapper, ForkAttrs, AttributeSet(), {})
__OMP_RTL_ATTRS(__tgt_target_data_end_nowait_mapper, ForkAttrs,
AttributeSet(), {})
+__OMP_RTL_ATTRS(__tgt_target_enter_data_mapper, ForkAttrs, AttributeSet(), {})
+__OMP_RTL_ATTRS(__tgt_target_enter_data_nowait_mapper, ForkAttrs,
+ AttributeSet(), {})
+__OMP_RTL_ATTRS(__tgt_target_exit_data_mapper, ForkAttrs, AttributeSet(), {})
+__OMP_RTL_ATTRS(__tgt_target_exit_data_nowait_mapper, ForkAttrs,
+ AttributeSet(), {})
__OMP_RTL_ATTRS(__tgt_target_data_update_mapper, ForkAttrs, AttributeSet(), {})
__OMP_RTL_ATTRS(__tgt_target_data_update_nowait_mapper, ForkAttrs,
AttributeSet(), {})
Index: clang/test/OpenMP/target_map_member_expr_array_section_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_map_member_expr_array_section_codegen.cpp
+++ clang/test/OpenMP/target_map_member_expr_array_section_codegen.cpp
@@ -60,7 +60,7 @@
// CHECK: [[BPTR:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BPTRS]], i32 0, i32 0
// CHECK: [[PTR:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRS]], i32 0, i32 0
// CHECK: [[SIZE:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* [[SIZES]], i32 0, i32 0
- // CHECK: call void @__tgt_target_data_begin_mapper(i64 -1, i32 2, i8** [[BPTR]], i8** [[PTR]], i64* [[SIZE]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAP_ENTER]], i32 0, i32 0), i8** null)
+ // CHECK: call void @__tgt_target_enter_data_mapper(i64 -1, i32 2, i8** [[BPTR]], i8** [[PTR]], i64* [[SIZE]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAP_ENTER]], i32 0, i32 0), i8** null)
#pragma omp target enter data map(alloc : s.data[:6])
}
@@ -104,7 +104,7 @@
// CHECK: [[BPTR:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BPTRS]], i32 0, i32 0
// CHECK: [[PTR:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRS]], i32 0, i32 0
// CHECK: [[SIZE:%.+]] = getelementptr inbounds [2 x i64], [2 x i64]* [[SIZES]], i32 0, i32 0
- // CHECK: call void @__tgt_target_data_end_mapper(i64 -1, i32 2, i8** [[BPTR]], i8** [[PTR]], i64* [[SIZE]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAP_EXIT]], i32 0, i32 0), i8** null)
+ // CHECK: call void @__tgt_target_exit_data_mapper(i64 -1, i32 2, i8** [[BPTR]], i8** [[PTR]], i64* [[SIZE]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAP_EXIT]], i32 0, i32 0), i8** null)
#pragma omp target exit data map(delete : s.data[:6])
}
};
Index: clang/test/OpenMP/target_exit_data_depend_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_exit_data_depend_codegen.cpp
+++ clang/test/OpenMP/target_exit_data_depend_codegen.cpp
@@ -373,7 +373,7 @@
}
// CK1: define internal{{.*}} i32 [[TASK_ENTRY0]](i32{{.*}}, %struct.kmp_task_t_with_privates* noalias %1)
-// CK1-DAG: call void @__tgt_target_data_end_nowait_mapper(i64 [[DEV:%[^,]+]], i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** [[GEPM:%.+]])
+// CK1-DAG: call void @__tgt_target_exit_data_nowait_mapper(i64 [[DEV:%[^,]+]], i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** [[GEPM:%.+]])
// CK1-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
// CK1-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}},
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
@@ -389,7 +389,7 @@
// CK1: }
// CK1: define internal{{.*}} i32 [[TASK_ENTRY2]](i32{{.*}}, %struct.kmp_task_t_with_privates{{.+}}* noalias %1)
-// CK1-DAG: call void @__tgt_target_data_end_mapper(i64 4, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}, i8** [[GEPM:%.+]])
+// CK1-DAG: call void @__tgt_target_exit_data_mapper(i64 4, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}, i8** [[GEPM:%.+]])
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
@@ -403,7 +403,7 @@
// CK1: }
// CK1: define internal{{.*}} i32 [[TASK_ENTRY3]](i32{{.*}}, %struct.kmp_task_t_with_privates{{.+}}* noalias %1)
-// CK1-DAG: call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}, i8** [[GEPM:%.+]])
+// CK1-DAG: call void @__tgt_target_exit_data_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}, i8** [[GEPM:%.+]])
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
@@ -413,12 +413,12 @@
// CK1-DAG: [[S]] = load [1 x i64]*, [1 x i64]** [[S_PRIV:%[^,]+]],
// CK1-DAG: [[M]] = load [1 x i8*]*, [1 x i8*]** [[M_PRIV:%[^,]+]],
// CK1-DAG: call void (i8*, ...) %{{.+}}(i8* %{{[^,]+}}, [1 x i8*]** [[BP_PRIV]], [1 x i8*]** [[P_PRIV]], [1 x i64]** [[S_PRIV]], [1 x i8*]** [[M_PRIV]])
-// CK1-NOT: __tgt_target_data_end_mapper
+// CK1-NOT: __tgt_target_exit_data_mapper
// CK1: ret i32 0
// CK1: }
// CK1: define internal{{.*}} i32 [[TASK_ENTRY4]](i32{{.*}}, %struct.kmp_task_t_with_privates{{.+}}* noalias %1)
-// CK1-DAG: call void @__tgt_target_data_end_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE04]]{{.+}}, i8** [[GEPM:%.+]])
+// CK1-DAG: call void @__tgt_target_exit_data_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE04]]{{.+}}, i8** [[GEPM:%.+]])
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
@@ -428,7 +428,7 @@
// CK1-DAG: [[S]] = load [2 x i64]*, [2 x i64]** [[S_PRIV:%[^,]+]],
// CK1-DAG: [[M]] = load [2 x i8*]*, [2 x i8*]** [[M_PRIV:%[^,]+]],
// CK1-DAG: call void (i8*, ...) %{{.+}}(i8* %{{[^,]+}}, [2 x i8*]** [[BP_PRIV]], [2 x i8*]** [[P_PRIV]], [2 x i64]** [[S_PRIV]], [2 x i8*]** [[M_PRIV]])
-// CK1-NOT: __tgt_target_data_end_mapper
+// CK1-NOT: __tgt_target_exit_data_mapper
// CK1: ret i32 0
// CK1: }
Index: clang/test/OpenMP/target_exit_data_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_exit_data_codegen.cpp
+++ clang/test/OpenMP/target_exit_data_codegen.cpp
@@ -50,8 +50,8 @@
float lb[arg];
// Region 00
- // CK1-NOT: __tgt_target_data_begin
- // CK1-DAG: call void @__tgt_target_data_end_nowait_mapper(i64 [[DEV:%[^,]+]], i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+ // CK1-NOT: __tgt_target_enter_data
+ // CK1-DAG: call void @__tgt_target_exit_data_nowait_mapper(i64 [[DEV:%[^,]+]], i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
// CK1-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
// CK1-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}},
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
@@ -74,10 +74,10 @@
{++arg;}
// Region 02
- // CK1-NOT: __tgt_target_data_begin
+ // CK1-NOT: __tgt_target_enter_data
// CK1: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
// CK1: [[IFTHEN]]
- // CK1-DAG: call void @__tgt_target_data_end_mapper(i64 4, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}, i8** null)
+ // CK1-DAG: call void @__tgt_target_exit_data_mapper(i64 4, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}, i8** null)
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
@@ -100,8 +100,8 @@
{++arg;}
// Region 03
- // CK1-NOT: __tgt_target_data_begin
- // CK1-DAG: call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}, i8** null)
+ // CK1-NOT: __tgt_target_enter_data
+ // CK1-DAG: call void @__tgt_target_exit_data_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}, i8** null)
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
@@ -125,8 +125,8 @@
{++arg;}
// Region 04
- // CK1-NOT: __tgt_target_data_begin
- // CK1-DAG: call void @__tgt_target_data_end_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE04]]{{.+}}, i8** null)
+ // CK1-NOT: __tgt_target_enter_data
+ // CK1-DAG: call void @__tgt_target_exit_data_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE04]]{{.+}}, i8** null)
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
@@ -155,8 +155,8 @@
{++arg;}
// Region 05
- // CK1-NOT: __tgt_target_data_begin
- // CK1-DAG: call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}}, i8** null)
+ // CK1-NOT: __tgt_target_enter_data
+ // CK1-DAG: call void @__tgt_target_exit_data_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}}, i8** null)
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
@@ -180,8 +180,8 @@
{++arg;}
// Region 06
- // CK1-NOT: __tgt_target_data_begin
- // CK1-DAG: call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}}, i8** null)
+ // CK1-NOT: __tgt_target_enter_data
+ // CK1-DAG: call void @__tgt_target_exit_data_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}}, i8** null)
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
@@ -242,10 +242,10 @@
}
// Region 00
-// CK2-NOT: __tgt_target_data_begin
+// CK2-NOT: __tgt_target_enter_data
// CK2: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
// CK2: [[IFTHEN]]
-// CK2-DAG: call void @__tgt_target_data_end_mapper(i64 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:.+]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+// CK2-DAG: call void @__tgt_target_exit_data_mapper(i64 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:.+]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
// CK2-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
// CK2-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}},
// CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
@@ -298,9 +298,9 @@
// CK3-LABEL: no_target_devices
void no_target_devices(int arg) {
- // CK3-NOT: tgt_target_data_begin
+ // CK3-NOT: tgt_target_enter_data
// CK3: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
- // CK3-NOT: tgt_target_data_end
+ // CK3-NOT: tgt_target_exit_data
// CK3: ret
#pragma omp target exit data map(from: arg) if(arg) device(4)
{++arg;}
@@ -346,10 +346,10 @@
}
// Region 00
-// CK4-NOT: __tgt_target_data_begin
+// CK4-NOT: __tgt_target_enter_data
// CK4: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
// CK4: [[IFTHEN]]
-// CK4-DAG: call void @__tgt_target_data_end_mapper(i64 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:.+]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+// CK4-DAG: call void @__tgt_target_exit_data_mapper(i64 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:.+]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
// CK4-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
// CK4-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}},
// CK4-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
Index: clang/test/OpenMP/target_enter_data_depend_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_enter_data_depend_codegen.cpp
+++ clang/test/OpenMP/target_enter_data_depend_codegen.cpp
@@ -373,7 +373,7 @@
}
// CK1: define internal{{.*}} i32 [[TASK_ENTRY0]](i32{{.*}}, %struct.kmp_task_t_with_privates* noalias %1)
-// CK1-DAG: call void @__tgt_target_data_begin_nowait_mapper(i64 [[DEV:%[^,]+]], i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** [[GEPM:%.+]])
+// CK1-DAG: call void @__tgt_target_enter_data_nowait_mapper(i64 [[DEV:%[^,]+]], i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** [[GEPM:%.+]])
// CK1-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
// CK1-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}},
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
@@ -389,7 +389,7 @@
// CK1: }
// CK1: define internal{{.*}} i32 [[TASK_ENTRY2]](i32{{.*}}, %struct.kmp_task_t_with_privates{{.+}}* noalias %1)
-// CK1-DAG: call void @__tgt_target_data_begin_mapper(i64 4, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}, i8** [[GEPM:%.+]])
+// CK1-DAG: call void @__tgt_target_enter_data_mapper(i64 4, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}, i8** [[GEPM:%.+]])
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
@@ -403,7 +403,7 @@
// CK1: }
// CK1: define internal{{.*}} i32 [[TASK_ENTRY3]](i32{{.*}}, %struct.kmp_task_t_with_privates{{.+}}* noalias %1)
-// CK1-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}, i8** [[GEPM:%.+]])
+// CK1-DAG: call void @__tgt_target_enter_data_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}, i8** [[GEPM:%.+]])
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
@@ -413,12 +413,12 @@
// CK1-DAG: [[S]] = load [1 x i64]*, [1 x i64]** [[S_PRIV:%[^,]+]],
// CK1-DAG: [[M]] = load [1 x i8*]*, [1 x i8*]** [[M_PRIV:%[^,]+]],
// CK1-DAG: call void (i8*, ...) %{{.+}}(i8* %{{[^,]+}}, [1 x i8*]** [[BP_PRIV]], [1 x i8*]** [[P_PRIV]], [1 x i64]** [[S_PRIV]], [1 x i8*]** [[M_PRIV]])
-// CK1-NOT: __tgt_target_data_end
+// CK1-NOT: __tgt_target_exit_data
// CK1: ret i32 0
// CK1: }
// CK1: define internal{{.*}} i32 [[TASK_ENTRY4]](i32{{.*}}, %struct.kmp_task_t_with_privates{{.+}}* noalias %1)
-// CK1-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE04]]{{.+}}, i8** [[GEPM:%.+]])
+// CK1-DAG: call void @__tgt_target_enter_data_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE04]]{{.+}}, i8** [[GEPM:%.+]])
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
@@ -428,7 +428,7 @@
// CK1-DAG: [[S]] = load [2 x i64]*, [2 x i64]** [[S_PRIV:%[^,]+]],
// CK1-DAG: [[M]] = load [2 x i8*]*, [2 x i8*]** [[M_PRIV:%[^,]+]],
// CK1-DAG: call void (i8*, ...) %{{.+}}(i8* %{{[^,]+}}, [2 x i8*]** [[BP_PRIV]], [2 x i8*]** [[P_PRIV]], [2 x i64]** [[S_PRIV]], [2 x i8*]** [[M_PRIV]])
-// CK1-NOT: __tgt_target_data_end
+// CK1-NOT: __tgt_target_exit_data
// CK1: ret i32 0
// CK1: }
Index: clang/test/OpenMP/target_enter_data_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_enter_data_codegen.cpp
+++ clang/test/OpenMP/target_enter_data_codegen.cpp
@@ -50,7 +50,7 @@
float lb[arg];
// Region 00
- // CK1-DAG: call void @__tgt_target_data_begin_nowait_mapper(i64 [[DEV:%[^,]+]], i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+ // CK1-DAG: call void @__tgt_target_enter_data_nowait_mapper(i64 [[DEV:%[^,]+]], i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
// CK1-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
// CK1-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}},
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
@@ -64,7 +64,7 @@
// CK1-DAG: store [100 x double]* @gc, [100 x double]** [[CP0]]
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
- // CK1-NOT: __tgt_target_data_end
+ // CK1-NOT: __tgt_target_exit_data
#pragma omp target enter data if(1+3-5) device(arg) map(alloc: gc) nowait
{++arg;}
@@ -76,7 +76,7 @@
// Region 02
// CK1: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
// CK1: [[IFTHEN]]
- // CK1-DAG: call void @__tgt_target_data_begin_mapper(i64 4, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}, i8** null)
+ // CK1-DAG: call void @__tgt_target_enter_data_mapper(i64 4, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}, i8** null)
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
@@ -92,7 +92,7 @@
// CK1: br label %[[IFEND]]
// CK1: [[IFEND]]
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
- // CK1-NOT: __tgt_target_data_end
+ // CK1-NOT: __tgt_target_exit_data
#pragma omp target enter data map(to: arg) if(arg) device(4)
{++arg;}
@@ -100,7 +100,7 @@
{++arg;}
// Region 03
- // CK1-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}, i8** null)
+ // CK1-DAG: call void @__tgt_target_enter_data_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}, i8** null)
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
@@ -117,7 +117,7 @@
// CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
// CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
- // CK1-NOT: __tgt_target_data_end
+ // CK1-NOT: __tgt_target_exit_data
#pragma omp target enter data map(always, to: lb)
{++arg;}
@@ -125,7 +125,7 @@
{++arg;}
// Region 04
- // CK1-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE04]]{{.+}}, i8** null)
+ // CK1-DAG: call void @__tgt_target_enter_data_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE04]]{{.+}}, i8** null)
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
@@ -147,7 +147,7 @@
// CK1-DAG: [[SEC11]] = load double*, double** getelementptr inbounds ([[ST]], [[ST]]* @gb, i32 0, i32 1),
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
- // CK1-NOT: __tgt_target_data_end
+ // CK1-NOT: __tgt_target_exit_data
#pragma omp target enter data map(to: gb.b[:3])
{++arg;}
@@ -155,7 +155,7 @@
{++arg;}
// Region 05
- // CK1-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}}, i8** null)
+ // CK1-DAG: call void @__tgt_target_enter_data_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}}, i8** null)
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
@@ -172,7 +172,7 @@
// CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
// CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
- // CK1-NOT: __tgt_target_data_end
+ // CK1-NOT: __tgt_target_exit_data
#pragma omp target enter data map(close, to: lb)
{++arg;}
@@ -180,7 +180,7 @@
{++arg;}
// Region 06
- // CK1-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}}, i8** null)
+ // CK1-DAG: call void @__tgt_target_enter_data_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}}, i8** null)
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
@@ -197,7 +197,7 @@
// CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
// CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
- // CK1-NOT: __tgt_target_data_end
+ // CK1-NOT: __tgt_target_exit_data
#pragma omp target enter data map(always close, to: lb)
{++arg;}
}
@@ -241,7 +241,7 @@
float lb[arg];
// Region 00
- // CK1A-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:32|64]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+ // CK1A-DAG: call void @__tgt_target_enter_data_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:32|64]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
// CK1A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK1A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK1A-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
@@ -258,7 +258,7 @@
// CK1A-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
// CK1A-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
// CK1A: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
- // CK1A-NOT: __tgt_target_data_end
+ // CK1A-NOT: __tgt_target_exit_data
#pragma omp target enter data map(present, to: lb)
{++arg;}
@@ -266,7 +266,7 @@
{++arg;}
// Region 01
- // CK1A-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
+ // CK1A-DAG: call void @__tgt_target_enter_data_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
// CK1A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK1A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK1A-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
@@ -283,7 +283,7 @@
// CK1A-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
// CK1A-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
// CK1A: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
- // CK1A-NOT: __tgt_target_data_end
+ // CK1A-NOT: __tgt_target_exit_data
#pragma omp target enter data map(always close present, to: lb)
{++arg;}
}
@@ -330,7 +330,7 @@
// Region 00
// CK2: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
// CK2: [[IFTHEN]]
-// CK2-DAG: call void @__tgt_target_data_begin_mapper(i64 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:64|32]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+// CK2-DAG: call void @__tgt_target_enter_data_mapper(i64 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:64|32]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
// CK2-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
// CK2-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}},
// CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
@@ -363,7 +363,7 @@
// CK2: br label %[[IFEND]]
// CK2: [[IFEND]]
// CK2: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
-// CK2-NOT: __tgt_target_data_end
+// CK2-NOT: __tgt_target_exit_data
#endif
///==========================================================================///
// RUN: %clang_cc1 -DCK3 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK3 --check-prefix CK3-64
@@ -384,9 +384,9 @@
// CK3-LABEL: no_target_devices
void no_target_devices(int arg) {
- // CK3-NOT: tgt_target_data_begin
+ // CK3-NOT: tgt_target_enter_data
// CK3: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
- // CK3-NOT: tgt_target_data_end
+ // CK3-NOT: tgt_target_exit_data
// CK3: ret
#pragma omp target enter data map(to: arg) if(arg) device(4)
{++arg;}
@@ -430,10 +430,10 @@
// CK4-LABEL: device_side_scan
void device_side_scan(int arg) {
- // CK4: tgt_target_data_begin
+ // CK4: tgt_target_enter_data
// CK4: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
// CK4: ret
- // TCK4-NOT: tgt_target_data_begin
+ // TCK4-NOT: tgt_target_enter_data
#pragma omp target enter data map(to: arg) if(arg) device(4)
{++arg;}
}
@@ -480,7 +480,7 @@
// Region 00
// CK5: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
// CK5: [[IFTHEN]]
-// CK5-DAG: call void @__tgt_target_data_begin_mapper(i64 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:64|32]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
+// CK5-DAG: call void @__tgt_target_enter_data_mapper(i64 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:64|32]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
// CK5-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
// CK5-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}},
// CK5-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
@@ -513,6 +513,6 @@
// CK5: br label %[[IFEND]]
// CK5: [[IFEND]]
// CK5: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
-// CK5-NOT: __tgt_target_data_end
+// CK5-NOT: __tgt_target_exit_data
#endif
#endif
Index: clang/test/OpenMP/declare_mapper_codegen.cpp
===================================================================
--- clang/test/OpenMP/declare_mapper_codegen.cpp
+++ clang/test/OpenMP/declare_mapper_codegen.cpp
@@ -308,7 +308,7 @@
++c.a;
}
- // CK0-DAG: call void @__tgt_target_data_begin_mapper(i64 {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[EDSIZES]]{{.+}}, {{.+}}[[EDTYPES]]{{.+}}, i8** [[MPRGEP:%.+]])
+ // CK0-DAG: call void @__tgt_target_enter_data_mapper(i64 {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[EDSIZES]]{{.+}}, {{.+}}[[EDTYPES]]{{.+}}, i8** [[MPRGEP:%.+]])
// CK0-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
// CK0-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
// CK0-DAG: [[MPRGEP]] = bitcast [1 x i8*]* [[MPR:%[^,]+]] to i8**
@@ -322,7 +322,7 @@
// CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR1]]
#pragma omp target enter data map(mapper(id),to: c)
- // CK0-DAG: call void @__tgt_target_data_begin_nowait_mapper(i64 {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[EDNWSIZES]]{{.+}}, {{.+}}[[EDNWTYPES]]{{.+}}, i8** [[MPRGEP:%.+]])
+ // CK0-DAG: call void @__tgt_target_enter_data_nowait_mapper(i64 {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[EDNWSIZES]]{{.+}}, {{.+}}[[EDNWTYPES]]{{.+}}, i8** [[MPRGEP:%.+]])
// CK0-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
// CK0-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
// CK0-DAG: [[MPRGEP]] = bitcast [1 x i8*]* [[MPR:%[^,]+]] to i8**
@@ -336,7 +336,7 @@
// CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR1]]
#pragma omp target enter data map(mapper(id),to: c) nowait
- // CK0-DAG: call void @__tgt_target_data_end_mapper(i64 {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[EXDSIZES]]{{.+}}, {{.+}}[[EXDTYPES]]{{.+}}, i8** [[MPRGEP:%.+]])
+ // CK0-DAG: call void @__tgt_target_exit_data_mapper(i64 {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[EXDSIZES]]{{.+}}, {{.+}}[[EXDTYPES]]{{.+}}, i8** [[MPRGEP:%.+]])
// CK0-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
// CK0-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
// CK0-DAG: [[MPRGEP]] = bitcast [1 x i8*]* [[MPR:%[^,]+]] to i8**
@@ -350,7 +350,7 @@
// CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR1]]
#pragma omp target exit data map(mapper(id),from: c)
- // CK0-DAG: call void @__tgt_target_data_end_nowait_mapper(i64 {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[EXDNWSIZES]]{{.+}}, {{.+}}[[EXDNWTYPES]]{{.+}}, i8** [[MPRGEP:%.+]])
+ // CK0-DAG: call void @__tgt_target_exit_data_nowait_mapper(i64 {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[EXDNWSIZES]]{{.+}}, {{.+}}[[EXDNWTYPES]]{{.+}}, i8** [[MPRGEP:%.+]])
// CK0-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
// CK0-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
// CK0-DAG: [[MPRGEP]] = bitcast [1 x i8*]* [[MPR:%[^,]+]] to i8**
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -10341,12 +10341,12 @@
RuntimeFunction RTLFn;
switch (D.getDirectiveKind()) {
case OMPD_target_enter_data:
- RTLFn = HasNowait ? OMPRTL___tgt_target_data_begin_nowait_mapper
- : OMPRTL___tgt_target_data_begin_mapper;
+ RTLFn = HasNowait ? OMPRTL___tgt_target_enter_data_nowait_mapper
+ : OMPRTL___tgt_target_enter_data_mapper;
break;
case OMPD_target_exit_data:
- RTLFn = HasNowait ? OMPRTL___tgt_target_data_end_nowait_mapper
- : OMPRTL___tgt_target_data_end_mapper;
+ RTLFn = HasNowait ? OMPRTL___tgt_target_exit_data_nowait_mapper
+ : OMPRTL___tgt_target_exit_data_mapper;
break;
case OMPD_target_update:
RTLFn = HasNowait ? OMPRTL___tgt_target_data_update_nowait_mapper
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
