Skip to content

[NFC][OpenMP] Add tests for mapping pointers and their dereferences. #146934

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 3 commits into from
Jul 8, 2025
Merged

[NFC][OpenMP] Add tests for mapping pointers and their dereferences. #146934

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
145160b
[NFC][OpenMP] Add tests for mapping pointers and their dereferences.
abhinavgaba Jul 3, 2025
15e7d50
Minor clang-format fix.
abhinavgaba Jul 3, 2025
a9b9127
Add another target construct in one test.
abhinavgaba Jul 7, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
393 changes: 237 additions & 156 deletions clang/test/OpenMP/target_map_both_pointer_pointee_codegen.cpp
View file
Open in desktop

Large diffs are not rendered by default.

249 changes: 249 additions & 0 deletions clang/test/OpenMP/target_map_both_pointer_pointee_codegen_global.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,249 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --check-globals --filter-out-after "getelem.*kernel" --filter-out "= alloca.*" --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _ --global-value-regex "\.offload_.*" --global-hex-value-regex ".offload_maptypes.*"
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s
// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s

// expected-no-diagnostics
#ifndef HEADER
#define HEADER

int *ptr;

void f1() {
// &ptr, &ptr, sizeof(ptr), TO | PARAM
#pragma omp target map(ptr)
ptr[1] = 5;
}

void f2() {
// &ptr, &ptr[2], sizeof(ptr[2]), TO | FROM | PARAM | PTR_AND_OBJ
#pragma omp target map(ptr[2])
ptr[1] = 6;
}

void f3() {
// &ptr, &ptr[0], sizeof(ptr[0:2]), TO | FROM | PARAM | PTR_AND_OBJ
#pragma omp target map(ptr, ptr[0:2])
ptr[1] = 7;
}

void f4() {
// &ptr, &ptr[2], sizeof(ptr[2]), TO | FROM | PARAM | PTR_AND_OBJ
#pragma omp target map(ptr, ptr[2])
ptr[2] = 8;
}

void f5() {
// &ptr, &ptr[2], sizeof(ptr[2]), TO | FROM | PARAM | PTR_AND_OBJ
#pragma omp target map(ptr[2], ptr)
ptr[2] = 9;
}

void f6() {
// &ptr, &ptr[2], sizeof(ptr[2]), TO | FROM | PARAM | PTR_AND_OBJ
// TODO: PARAM should not be needed here.
#pragma omp target data map(ptr, ptr[2])
ptr[2] = 10;
}

void f7() {
// &ptr, &ptr[2], sizeof(ptr[2]), TO | FROM | PARAM | PTR_AND_OBJ
// TODO: PARAM should not be needed here.
#pragma omp target data map(ptr[2], ptr)
ptr[2] = 11;
}
#endif
//.
// CHECK: @.offload_sizes = private unnamed_addr constant [1 x i64] [i64 8]
// CHECK: @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 [[#0x23]]]
// CHECK: @.offload_sizes.1 = private unnamed_addr constant [1 x i64] [i64 4]
// CHECK: @.offload_maptypes.2 = private unnamed_addr constant [1 x i64] [i64 [[#0x33]]]
// CHECK: @.offload_sizes.3 = private unnamed_addr constant [1 x i64] [i64 8]
// CHECK: @.offload_maptypes.4 = private unnamed_addr constant [1 x i64] [i64 [[#0x33]]]
// CHECK: @.offload_sizes.5 = private unnamed_addr constant [1 x i64] [i64 4]
// CHECK: @.offload_maptypes.6 = private unnamed_addr constant [1 x i64] [i64 [[#0x33]]]
// CHECK: @.offload_sizes.7 = private unnamed_addr constant [1 x i64] [i64 4]
// CHECK: @.offload_maptypes.8 = private unnamed_addr constant [1 x i64] [i64 [[#0x33]]]
// CHECK: @.offload_sizes.9 = private unnamed_addr constant [1 x i64] [i64 4]
// CHECK: @.offload_maptypes.10 = private unnamed_addr constant [1 x i64] [i64 [[#0x33]]]
// CHECK: @.offload_sizes.11 = private unnamed_addr constant [1 x i64] [i64 4]
// CHECK: @.offload_maptypes.12 = private unnamed_addr constant [1 x i64] [i64 [[#0x33]]]
//.
// CHECK-LABEL: define {{[^@]+}}@_Z2f1v
// CHECK-SAME: () #[[ATTR0:[0-9]+]] {
// CHECK: entry:
// CHECK: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0
// CHECK: store ptr @ptr, ptr [[TMP0]], align 8
// CHECK: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0
// CHECK: store ptr @ptr, ptr [[TMP1]], align 8
// CHECK: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0
// CHECK: store ptr null, ptr [[TMP2]], align 8
// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0
//
//
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f1v_l14
// CHECK-SAME: (ptr noundef nonnull align 8 dereferenceable(8) [[PTR:%.*]]) #[[ATTR1:[0-9]+]] {
// CHECK: entry:
// CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 8
// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8, !nonnull [[META13:![0-9]+]], !align [[META14:![0-9]+]]
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is there a way to scrub this metadata from the CHECKs? It adds unnecessary noise/diffs here. I couldn't find one in update_cc_test_checks. There's an option to scrub attributes in update_test_checks, but even that doesn't have one for metadata.

If not, @alexey-bataev, please help merge.

// CHECK: [[TMP1:%.*]] = load ptr, ptr [[TMP0]], align 8
// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 1
// CHECK: store i32 5, ptr [[ARRAYIDX]], align 4
// CHECK: ret void
//
//
// CHECK-LABEL: define {{[^@]+}}@_Z2f2v
// CHECK-SAME: () #[[ATTR0]] {
// CHECK: entry:
// CHECK: [[TMP0:%.*]] = load ptr, ptr @ptr, align 8
// CHECK: [[TMP1:%.*]] = load ptr, ptr @ptr, align 8
// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 2
// CHECK: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0
// CHECK: store ptr @ptr, ptr [[TMP2]], align 8
// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0
// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP3]], align 8
// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0
// CHECK: store ptr null, ptr [[TMP4]], align 8
// CHECK: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0
//
//
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f2v_l20
// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR1]] {
// CHECK: entry:
// CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 8
// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8
// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 1
// CHECK: store i32 6, ptr [[ARRAYIDX]], align 4
// CHECK: ret void
//
//
// CHECK-LABEL: define {{[^@]+}}@_Z2f3v
// CHECK-SAME: () #[[ATTR0]] {
// CHECK: entry:
// CHECK: [[TMP0:%.*]] = load ptr, ptr @ptr, align 8
// CHECK: [[TMP1:%.*]] = load ptr, ptr @ptr, align 8
// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds nuw i32, ptr [[TMP1]], i64 0
// CHECK: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0
// CHECK: store ptr @ptr, ptr [[TMP2]], align 8
// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0
// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP3]], align 8
// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0
// CHECK: store ptr null, ptr [[TMP4]], align 8
// CHECK: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0
//
//
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f3v_l26
// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR1]] {
// CHECK: entry:
// CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 8
// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8
// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 1
// CHECK: store i32 7, ptr [[ARRAYIDX]], align 4
// CHECK: ret void
//
//
// CHECK-LABEL: define {{[^@]+}}@_Z2f4v
// CHECK-SAME: () #[[ATTR0]] {
// CHECK: entry:
// CHECK: [[TMP0:%.*]] = load ptr, ptr @ptr, align 8
// CHECK: [[TMP1:%.*]] = load ptr, ptr @ptr, align 8
// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 2
// CHECK: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0
// CHECK: store ptr @ptr, ptr [[TMP2]], align 8
// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0
// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP3]], align 8
// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0
// CHECK: store ptr null, ptr [[TMP4]], align 8
// CHECK: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0
//
//
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f4v_l32
// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR1]] {
// CHECK: entry:
// CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 8
// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8
// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 2
// CHECK: store i32 8, ptr [[ARRAYIDX]], align 4
// CHECK: ret void
//
//
// CHECK-LABEL: define {{[^@]+}}@_Z2f5v
// CHECK-SAME: () #[[ATTR0]] {
// CHECK: entry:
// CHECK: [[TMP0:%.*]] = load ptr, ptr @ptr, align 8
// CHECK: [[TMP1:%.*]] = load ptr, ptr @ptr, align 8
// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 2
// CHECK: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0
// CHECK: store ptr @ptr, ptr [[TMP2]], align 8
// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0
// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP3]], align 8
// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0
// CHECK: store ptr null, ptr [[TMP4]], align 8
// CHECK: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK: [[TMP6:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], ptr [[KERNEL_ARGS:%.*]], i32 0, i32 0
//
//
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z2f5v_l38
// CHECK-SAME: (ptr noundef [[PTR:%.*]]) #[[ATTR1]] {
// CHECK: entry:
// CHECK: store ptr [[PTR]], ptr [[PTR_ADDR:%.*]], align 8
// CHECK: [[TMP0:%.*]] = load ptr, ptr [[PTR_ADDR]], align 8
// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 2
// CHECK: store i32 9, ptr [[ARRAYIDX]], align 4
// CHECK: ret void
//
//
// CHECK-LABEL: define {{[^@]+}}@_Z2f6v
// CHECK-SAME: () #[[ATTR0]] {
// CHECK: entry:
// CHECK: [[TMP0:%.*]] = load ptr, ptr @ptr, align 8
// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 2
// CHECK: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0
// CHECK: store ptr @ptr, ptr [[TMP1]], align 8
// CHECK: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0
// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP2]], align 8
// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0
// CHECK: store ptr null, ptr [[TMP3]], align 8
// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK: call void @__tgt_target_data_begin_mapper(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 1, ptr [[TMP4]], ptr [[TMP5]], ptr @.offload_sizes.9, ptr @.offload_maptypes.10, ptr null, ptr null)
// CHECK: [[TMP6:%.*]] = load ptr, ptr @ptr, align 8
// CHECK: [[ARRAYIDX1:%.*]] = getelementptr inbounds i32, ptr [[TMP6]], i64 2
// CHECK: store i32 10, ptr [[ARRAYIDX1]], align 4
// CHECK: [[TMP7:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK: [[TMP8:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK: call void @__tgt_target_data_end_mapper(ptr @[[GLOB1]], i64 -1, i32 1, ptr [[TMP7]], ptr [[TMP8]], ptr @.offload_sizes.9, ptr @.offload_maptypes.10, ptr null, ptr null)
// CHECK: ret void
//
//
// CHECK-LABEL: define {{[^@]+}}@_Z2f7v
// CHECK-SAME: () #[[ATTR0]] {
// CHECK: entry:
// CHECK: [[TMP0:%.*]] = load ptr, ptr @ptr, align 8
// CHECK: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP0]], i64 2
// CHECK: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS:%.*]], i32 0, i32 0
// CHECK: store ptr @ptr, ptr [[TMP1]], align 8
// CHECK: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS:%.*]], i32 0, i32 0
// CHECK: store ptr [[ARRAYIDX]], ptr [[TMP2]], align 8
// CHECK: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS:%.*]], i64 0, i64 0
// CHECK: store ptr null, ptr [[TMP3]], align 8
// CHECK: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK: call void @__tgt_target_data_begin_mapper(ptr @[[GLOB1]], i64 -1, i32 1, ptr [[TMP4]], ptr [[TMP5]], ptr @.offload_sizes.11, ptr @.offload_maptypes.12, ptr null, ptr null)
// CHECK: [[TMP6:%.*]] = load ptr, ptr @ptr, align 8
// CHECK: [[ARRAYIDX1:%.*]] = getelementptr inbounds i32, ptr [[TMP6]], i64 2
// CHECK: store i32 11, ptr [[ARRAYIDX1]], align 4
// CHECK: [[TMP7:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
// CHECK: [[TMP8:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
// CHECK: call void @__tgt_target_data_end_mapper(ptr @[[GLOB1]], i64 -1, i32 1, ptr [[TMP7]], ptr [[TMP8]], ptr @.offload_sizes.11, ptr @.offload_maptypes.12, ptr null, ptr null)
// CHECK: ret void
//
Loading
Loading