diff options
| author | Gheorghe-Teodor Bercea <gheorghe-teod.bercea@ibm.com> | 2017-09-26 18:12:12 +0000 |
|---|---|---|
| committer | Gheorghe-Teodor Bercea <gheorghe-teod.bercea@ibm.com> | 2017-09-26 18:12:12 +0000 |
| commit | b379ba6a62efca54313c1a91ebc5da3dbe25f3c0 (patch) | |
| tree | 81b9b31ac6adac7a28f1bf2c72d199151f123726 /clang/test/OpenMP/target_map_codegen.cpp | |
| parent | 9f1a390f72ea5ee64f3e7859133983e878835c0a (diff) | |
| download | bcm5719-llvm-b379ba6a62efca54313c1a91ebc5da3dbe25f3c0.tar.gz bcm5719-llvm-b379ba6a62efca54313c1a91ebc5da3dbe25f3c0.zip | |
[OpenMP] Add an additional test for D34888
Summary: Test for checking if the mapping is performed correctly. This is a test initially included in Patch https://reviews.llvm.org/D29905
Reviewers: Hahnfeld, carlo.bertolli, caomhin
Reviewed By: Hahnfeld
Subscribers: tra, cfe-commits
Differential Revision: https://reviews.llvm.org/D38040
llvm-svn: 314228
Diffstat (limited to 'clang/test/OpenMP/target_map_codegen.cpp')
| -rw-r--r-- | clang/test/OpenMP/target_map_codegen.cpp | 30 |
1 files changed, 30 insertions, 0 deletions
diff --git a/clang/test/OpenMP/target_map_codegen.cpp b/clang/test/OpenMP/target_map_codegen.cpp index 460a02ef2dd..7fde98d13cc 100644 --- a/clang/test/OpenMP/target_map_codegen.cpp +++ b/clang/test/OpenMP/target_map_codegen.cpp @@ -4554,3 +4554,33 @@ void explicit_maps_member_pointer_references(SSA *sap) { } #endif #endif + +///==========================================================================/// +// RUN: %clang -DCK30 -std=c++11 -fopenmp -S -emit-llvm -fopenmp-targets=nvptx64-nvidia-cuda \ +// RUN: -nocudalib %s -o - 2>&1 | FileCheck -check-prefix=CK30 %s + +#ifdef CK30 + +void target_maps_parallel_integer(int a){ + int ParamToKernel = a; +#pragma omp target map(tofrom: ParamToKernel) + { + ParamToKernel += 1; + } +} + +// CK30: {{.*}}void @__omp_offloading_{{.*}}(i32* dereferenceable(4) %ParamToKernel) + +// CK30: {{.*}}void {{.*}}target_maps_parallel_integer{{.*}} { + +// CK30: [[GEPOBP:%.+]] = getelementptr inbounds {{.*}}offload_baseptrs +// CK30: [[GEPOBPBIT:%.+]] = bitcast i8** [[GEPOBP]] +// CK30: store i32* %ParamToKernel, i32** [[GEPOBPBIT]] +// CK30: [[GEPOP:%.+]] = getelementptr inbounds {{.*}}offload_ptrs +// CK30: [[GEPOPBIT:%.+]] = bitcast i8** [[GEPOP]] +// CK30: store i32* %ParamToKernel, i32** [[GEPOPBIT]] +// CK30: [[GEPOBPARG:%.+]] = getelementptr inbounds {{.*}}offload_baseptrs +// CK30: [[GEPOPARG:%.+]] = getelementptr inbounds {{.*}}offload_ptrs +// CK30: call {{.*}}tgt_target({{.*}}i8** [[GEPOBPARG]], i8** [[GEPOPARG]] + +#endif |
