Skip to content

Commit

Permalink
OpenMP: Don't use target regions in library function test
Browse files Browse the repository at this point in the history 
Use pragma omp declare target to reduce the IR noise around the call
we're actually trying to test.
  • Loading branch information
@arsenm
arsenm committed Jun 30, 2023
1 parent 8f9eee3 commit 917eddf
Showing 1 changed file with 46 additions and 78 deletions.
124 changes: 46 additions & 78 deletions clang/test/Headers/amdgcn_openmp_device_math_c.c
View file
Original file line number Diff line number Diff line change
@@ -1,67 +1,42 @@
// 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[.].+[.|,]"
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
// RUN: %clang_cc1 -internal-isystem %S/Inputs/include -x c -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -triple powerpc64le-unknown-unknown -D__OFFLOAD_ARCH_gfx90a__ -emit-llvm-bc %s -o %t-host.bc
// RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -x c -fopenmp -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -o - | FileCheck %s --check-prefixes=CHECK
// REQUIRES: amdgpu-registered-target

#include <stdlib.h>
#pragma omp begin declare target

void test_math_int(int x) {
#pragma omp target
{
int l1 = abs(x);
}
}

void test_math_long(long x) {
#pragma omp target
{
long l1 = labs(x);
}
}

void test_math_long_long(long long x) {
#pragma omp target
{
long long l1 = llabs(x);
}
}
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_test_math_int_l9
// CHECK-SAME: (i64 noundef [[X:%.*]]) #[[ATTR0:[0-9]+]] {
// CHECK-LABEL: @test_math_int(
// CHECK-NEXT: entry:
// CHECK-NEXT: [[RETVAL_I:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK-NEXT: [[__X_ADDR_I:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK-NEXT: [[__SGN_I:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK-NEXT: [[X_ADDR:%.*]] = alloca i64, align 8, addrspace(5)
// CHECK-NEXT: [[X_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK-NEXT: [[L1:%.*]] = alloca i32, align 4, addrspace(5)
// CHECK-NEXT: [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X_ADDR]] to ptr
// CHECK-NEXT: [[L1_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[L1]] to ptr
// CHECK-NEXT: store i64 [[X]], ptr [[X_ADDR_ASCAST]], align 8
// CHECK-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr addrspacecast (ptr addrspace(1) @[[GLOB1:[0-9]+]] to ptr), i8 1, i1 true)
// CHECK-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
// CHECK-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK: user_code.entry:
// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[X_ADDR_ASCAST]], align 4
// CHECK-NEXT: store i32 [[X:%.*]], ptr [[X_ADDR_ASCAST]], align 4
// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[X_ADDR_ASCAST]], align 4
// CHECK-NEXT: [[RETVAL_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I]] to ptr
// CHECK-NEXT: [[__X_ADDR_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I]] to ptr
// CHECK-NEXT: [[__SGN_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[__SGN_I]] to ptr
// CHECK-NEXT: store i32 [[TMP1]], ptr [[__X_ADDR_ASCAST_I]], align 4
// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[__X_ADDR_ASCAST_I]], align 4
// CHECK-NEXT: [[SHR_I:%.*]] = ashr i32 [[TMP2]], 31
// CHECK-NEXT: store i32 [[TMP0]], ptr [[__X_ADDR_ASCAST_I]], align 4
// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[__X_ADDR_ASCAST_I]], align 4
// CHECK-NEXT: [[SHR_I:%.*]] = ashr i32 [[TMP1]], 31
// CHECK-NEXT: store i32 [[SHR_I]], ptr [[__SGN_ASCAST_I]], align 4
// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr [[__X_ADDR_ASCAST_I]], align 4
// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[__X_ADDR_ASCAST_I]], align 4
// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr [[__SGN_ASCAST_I]], align 4
// CHECK-NEXT: [[XOR_I:%.*]] = xor i32 [[TMP2]], [[TMP3]]
// CHECK-NEXT: [[TMP4:%.*]] = load i32, ptr [[__SGN_ASCAST_I]], align 4
// CHECK-NEXT: [[XOR_I:%.*]] = xor i32 [[TMP3]], [[TMP4]]
// CHECK-NEXT: [[TMP5:%.*]] = load i32, ptr [[__SGN_ASCAST_I]], align 4
// CHECK-NEXT: [[SUB_I:%.*]] = sub nsw i32 [[XOR_I]], [[TMP5]]
// CHECK-NEXT: [[SUB_I:%.*]] = sub nsw i32 [[XOR_I]], [[TMP4]]
// CHECK-NEXT: store i32 [[SUB_I]], ptr [[L1_ASCAST]], align 4
// CHECK-NEXT: call void @__kmpc_target_deinit(ptr addrspacecast (ptr addrspace(1) @[[GLOB1]] to ptr), i8 1)
// CHECK-NEXT: ret void
// CHECK: worker.exit:
// CHECK-NEXT: ret void
//
//
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_test_math_long_l16
// CHECK-SAME: (i64 noundef [[X:%.*]]) #[[ATTR0]] {
void test_math_int(int x) {
int l1 = abs(x);
}

// CHECK-LABEL: @test_math_long(
// CHECK-NEXT: entry:
// CHECK-NEXT: [[RETVAL_I:%.*]] = alloca i64, align 8, addrspace(5)
// CHECK-NEXT: [[__X_ADDR_I:%.*]] = alloca i64, align 8, addrspace(5)
Expand All @@ -70,33 +45,28 @@ void test_math_long_long(long long x) {
// CHECK-NEXT: [[L1:%.*]] = alloca i64, align 8, addrspace(5)
// CHECK-NEXT: [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X_ADDR]] to ptr
// CHECK-NEXT: [[L1_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[L1]] to ptr
// CHECK-NEXT: store i64 [[X]], ptr [[X_ADDR_ASCAST]], align 8
// CHECK-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr addrspacecast (ptr addrspace(1) @[[GLOB1]] to ptr), i8 1, i1 true)
// CHECK-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
// CHECK-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK: user_code.entry:
// CHECK-NEXT: [[TMP1:%.*]] = load i64, ptr [[X_ADDR_ASCAST]], align 8
// CHECK-NEXT: store i64 [[X:%.*]], ptr [[X_ADDR_ASCAST]], align 8
// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[X_ADDR_ASCAST]], align 8
// CHECK-NEXT: [[RETVAL_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I]] to ptr
// CHECK-NEXT: [[__X_ADDR_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I]] to ptr
// CHECK-NEXT: [[__SGN_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[__SGN_I]] to ptr
// CHECK-NEXT: store i64 [[TMP1]], ptr [[__X_ADDR_ASCAST_I]], align 8
// CHECK-NEXT: [[TMP2:%.*]] = load i64, ptr [[__X_ADDR_ASCAST_I]], align 8
// CHECK-NEXT: [[SHR_I:%.*]] = ashr i64 [[TMP2]], 63
// CHECK-NEXT: store i64 [[TMP0]], ptr [[__X_ADDR_ASCAST_I]], align 8
// CHECK-NEXT: [[TMP1:%.*]] = load i64, ptr [[__X_ADDR_ASCAST_I]], align 8
// CHECK-NEXT: [[SHR_I:%.*]] = ashr i64 [[TMP1]], 63
// CHECK-NEXT: store i64 [[SHR_I]], ptr [[__SGN_ASCAST_I]], align 8
// CHECK-NEXT: [[TMP3:%.*]] = load i64, ptr [[__X_ADDR_ASCAST_I]], align 8
// CHECK-NEXT: [[TMP2:%.*]] = load i64, ptr [[__X_ADDR_ASCAST_I]], align 8
// CHECK-NEXT: [[TMP3:%.*]] = load i64, ptr [[__SGN_ASCAST_I]], align 8
// CHECK-NEXT: [[XOR_I:%.*]] = xor i64 [[TMP2]], [[TMP3]]
// CHECK-NEXT: [[TMP4:%.*]] = load i64, ptr [[__SGN_ASCAST_I]], align 8
// CHECK-NEXT: [[XOR_I:%.*]] = xor i64 [[TMP3]], [[TMP4]]
// CHECK-NEXT: [[TMP5:%.*]] = load i64, ptr [[__SGN_ASCAST_I]], align 8
// CHECK-NEXT: [[SUB_I:%.*]] = sub nsw i64 [[XOR_I]], [[TMP5]]
// CHECK-NEXT: [[SUB_I:%.*]] = sub nsw i64 [[XOR_I]], [[TMP4]]
// CHECK-NEXT: store i64 [[SUB_I]], ptr [[L1_ASCAST]], align 8
// CHECK-NEXT: call void @__kmpc_target_deinit(ptr addrspacecast (ptr addrspace(1) @[[GLOB1]] to ptr), i8 1)
// CHECK-NEXT: ret void
// CHECK: worker.exit:
// CHECK-NEXT: ret void
//
//
// CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_test_math_long_long_l23
// CHECK-SAME: (i64 noundef [[X:%.*]]) #[[ATTR0]] {
void test_math_long(long x) {
long l1 = labs(x);
}

// CHECK-LABEL: @test_math_long_long(
// CHECK-NEXT: entry:
// CHECK-NEXT: [[RETVAL_I:%.*]] = alloca i64, align 8, addrspace(5)
// CHECK-NEXT: [[__X_ADDR_I:%.*]] = alloca i64, align 8, addrspace(5)
Expand All @@ -105,27 +75,25 @@ void test_math_long_long(long long x) {
// CHECK-NEXT: [[L1:%.*]] = alloca i64, align 8, addrspace(5)
// CHECK-NEXT: [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X_ADDR]] to ptr
// CHECK-NEXT: [[L1_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[L1]] to ptr
// CHECK-NEXT: store i64 [[X]], ptr [[X_ADDR_ASCAST]], align 8
// CHECK-NEXT: [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr addrspacecast (ptr addrspace(1) @[[GLOB1]] to ptr), i8 1, i1 true)
// CHECK-NEXT: [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
// CHECK-NEXT: br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
// CHECK: user_code.entry:
// CHECK-NEXT: [[TMP1:%.*]] = load i64, ptr [[X_ADDR_ASCAST]], align 8
// CHECK-NEXT: store i64 [[X:%.*]], ptr [[X_ADDR_ASCAST]], align 8
// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[X_ADDR_ASCAST]], align 8
// CHECK-NEXT: [[RETVAL_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I]] to ptr
// CHECK-NEXT: [[__X_ADDR_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I]] to ptr
// CHECK-NEXT: [[__SGN_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[__SGN_I]] to ptr
// CHECK-NEXT: store i64 [[TMP1]], ptr [[__X_ADDR_ASCAST_I]], align 8
// CHECK-NEXT: [[TMP2:%.*]] = load i64, ptr [[__X_ADDR_ASCAST_I]], align 8
// CHECK-NEXT: [[SHR_I:%.*]] = ashr i64 [[TMP2]], 63
// CHECK-NEXT: store i64 [[TMP0]], ptr [[__X_ADDR_ASCAST_I]], align 8
// CHECK-NEXT: [[TMP1:%.*]] = load i64, ptr [[__X_ADDR_ASCAST_I]], align 8
// CHECK-NEXT: [[SHR_I:%.*]] = ashr i64 [[TMP1]], 63
// CHECK-NEXT: store i64 [[SHR_I]], ptr [[__SGN_ASCAST_I]], align 8
// CHECK-NEXT: [[TMP3:%.*]] = load i64, ptr [[__X_ADDR_ASCAST_I]], align 8
// CHECK-NEXT: [[TMP2:%.*]] = load i64, ptr [[__X_ADDR_ASCAST_I]], align 8
// CHECK-NEXT: [[TMP3:%.*]] = load i64, ptr [[__SGN_ASCAST_I]], align 8
// CHECK-NEXT: [[XOR_I:%.*]] = xor i64 [[TMP2]], [[TMP3]]
// CHECK-NEXT: [[TMP4:%.*]] = load i64, ptr [[__SGN_ASCAST_I]], align 8
// CHECK-NEXT: [[XOR_I:%.*]] = xor i64 [[TMP3]], [[TMP4]]
// CHECK-NEXT: [[TMP5:%.*]] = load i64, ptr [[__SGN_ASCAST_I]], align 8
// CHECK-NEXT: [[SUB_I:%.*]] = sub nsw i64 [[XOR_I]], [[TMP5]]
// CHECK-NEXT: [[SUB_I:%.*]] = sub nsw i64 [[XOR_I]], [[TMP4]]
// CHECK-NEXT: store i64 [[SUB_I]], ptr [[L1_ASCAST]], align 8
// CHECK-NEXT: call void @__kmpc_target_deinit(ptr addrspacecast (ptr addrspace(1) @[[GLOB1]] to ptr), i8 1)
// CHECK-NEXT: ret void
// CHECK: worker.exit:
// CHECK-NEXT: ret void
//
void test_math_long_long(long long x) {
long long l1 = llabs(x);
}

#pragma omp end declare target

0 comments on commit 917eddf

Please sign in to comment.