clang/test/Headers/amdgcn_openmp_device_math_c.c

   1 // 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[.].+[.|,]"
   2 // 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
   3 // 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
   4 // REQUIRES: amdgpu-registered-target
   5
   6 #include <stdlib.h>
   7
   8 void test_math_int(int x) {
   9 #pragma omp target
  10   {
  11     int l1 = abs(x);
  12   }
  13 }
  14
  15 void test_math_long(long x) {
  16 #pragma omp target
  17   {
  18     long l1 = labs(x);
  19   }
  20 }
  21
  22 void test_math_long_long(long long x) {
  23 #pragma omp target
  24   {
  25     long long l1 = llabs(x);
  26   }
  27 }
  28 // CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_test_math_int_l9
  29 // CHECK-SAME: (i64 noundef [[X:%.*]]) #[[ATTR0:[0-9]+]] {
  30 // CHECK-NEXT:  entry:
  31 // CHECK-NEXT:    [[RETVAL_I:%.*]] = alloca i32, align 4, addrspace(5)
  32 // CHECK-NEXT:    [[__X_ADDR_I:%.*]] = alloca i32, align 4, addrspace(5)
  33 // CHECK-NEXT:    [[__SGN_I:%.*]] = alloca i32, align 4, addrspace(5)
  34 // CHECK-NEXT:    [[X_ADDR:%.*]] = alloca i64, align 8, addrspace(5)
  35 // CHECK-NEXT:    [[L1:%.*]] = alloca i32, align 4, addrspace(5)
  36 // CHECK-NEXT:    [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X_ADDR]] to ptr
  37 // CHECK-NEXT:    [[L1_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[L1]] to ptr
  38 // CHECK-NEXT:    store i64 [[X]], ptr [[X_ADDR_ASCAST]], align 8
  39 // CHECK-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr addrspacecast (ptr addrspace(1) @[[GLOB1:[0-9]+]] to ptr), i8 1, i1 true)
  40 // CHECK-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
  41 // CHECK-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
  42 // CHECK:       user_code.entry:
  43 // CHECK-NEXT:    [[TMP1:%.*]] = load i32, ptr [[X_ADDR_ASCAST]], align 4
  44 // CHECK-NEXT:    [[RETVAL_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I]] to ptr
  45 // CHECK-NEXT:    [[__X_ADDR_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I]] to ptr
  46 // CHECK-NEXT:    [[__SGN_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[__SGN_I]] to ptr
  47 // CHECK-NEXT:    store i32 [[TMP1]], ptr [[__X_ADDR_ASCAST_I]], align 4
  48 // CHECK-NEXT:    [[TMP2:%.*]] = load i32, ptr [[__X_ADDR_ASCAST_I]], align 4
  49 // CHECK-NEXT:    [[SHR_I:%.*]] = ashr i32 [[TMP2]], 31
  50 // CHECK-NEXT:    store i32 [[SHR_I]], ptr [[__SGN_ASCAST_I]], align 4
  51 // CHECK-NEXT:    [[TMP3:%.*]] = load i32, ptr [[__X_ADDR_ASCAST_I]], align 4
  52 // CHECK-NEXT:    [[TMP4:%.*]] = load i32, ptr [[__SGN_ASCAST_I]], align 4
  53 // CHECK-NEXT:    [[XOR_I:%.*]] = xor i32 [[TMP3]], [[TMP4]]
  54 // CHECK-NEXT:    [[TMP5:%.*]] = load i32, ptr [[__SGN_ASCAST_I]], align 4
  55 // CHECK-NEXT:    [[SUB_I:%.*]] = sub nsw i32 [[XOR_I]], [[TMP5]]
  56 // CHECK-NEXT:    store i32 [[SUB_I]], ptr [[L1_ASCAST]], align 4
  57 // CHECK-NEXT:    call void @__kmpc_target_deinit(ptr addrspacecast (ptr addrspace(1) @[[GLOB1]] to ptr), i8 1)
  58 // CHECK-NEXT:    ret void
  59 // CHECK:       worker.exit:
  60 // CHECK-NEXT:    ret void
  61 //
  62 //
  63 // CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_test_math_long_l16
  64 // CHECK-SAME: (i64 noundef [[X:%.*]]) #[[ATTR0]] {
  65 // CHECK-NEXT:  entry:
  66 // CHECK-NEXT:    [[RETVAL_I:%.*]] = alloca i64, align 8, addrspace(5)
  67 // CHECK-NEXT:    [[__X_ADDR_I:%.*]] = alloca i64, align 8, addrspace(5)
  68 // CHECK-NEXT:    [[__SGN_I:%.*]] = alloca i64, align 8, addrspace(5)
  69 // CHECK-NEXT:    [[X_ADDR:%.*]] = alloca i64, align 8, addrspace(5)
  70 // CHECK-NEXT:    [[L1:%.*]] = alloca i64, align 8, addrspace(5)
  71 // CHECK-NEXT:    [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X_ADDR]] to ptr
  72 // CHECK-NEXT:    [[L1_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[L1]] to ptr
  73 // CHECK-NEXT:    store i64 [[X]], ptr [[X_ADDR_ASCAST]], align 8
  74 // CHECK-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr addrspacecast (ptr addrspace(1) @[[GLOB1]] to ptr), i8 1, i1 true)
  75 // CHECK-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
  76 // CHECK-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
  77 // CHECK:       user_code.entry:
  78 // CHECK-NEXT:    [[TMP1:%.*]] = load i64, ptr [[X_ADDR_ASCAST]], align 8
  79 // CHECK-NEXT:    [[RETVAL_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I]] to ptr
  80 // CHECK-NEXT:    [[__X_ADDR_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I]] to ptr
  81 // CHECK-NEXT:    [[__SGN_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[__SGN_I]] to ptr
  82 // CHECK-NEXT:    store i64 [[TMP1]], ptr [[__X_ADDR_ASCAST_I]], align 8
  83 // CHECK-NEXT:    [[TMP2:%.*]] = load i64, ptr [[__X_ADDR_ASCAST_I]], align 8
  84 // CHECK-NEXT:    [[SHR_I:%.*]] = ashr i64 [[TMP2]], 63
  85 // CHECK-NEXT:    store i64 [[SHR_I]], ptr [[__SGN_ASCAST_I]], align 8
  86 // CHECK-NEXT:    [[TMP3:%.*]] = load i64, ptr [[__X_ADDR_ASCAST_I]], align 8
  87 // CHECK-NEXT:    [[TMP4:%.*]] = load i64, ptr [[__SGN_ASCAST_I]], align 8
  88 // CHECK-NEXT:    [[XOR_I:%.*]] = xor i64 [[TMP3]], [[TMP4]]
  89 // CHECK-NEXT:    [[TMP5:%.*]] = load i64, ptr [[__SGN_ASCAST_I]], align 8
  90 // CHECK-NEXT:    [[SUB_I:%.*]] = sub nsw i64 [[XOR_I]], [[TMP5]]
  91 // CHECK-NEXT:    store i64 [[SUB_I]], ptr [[L1_ASCAST]], align 8
  92 // CHECK-NEXT:    call void @__kmpc_target_deinit(ptr addrspacecast (ptr addrspace(1) @[[GLOB1]] to ptr), i8 1)
  93 // CHECK-NEXT:    ret void
  94 // CHECK:       worker.exit:
  95 // CHECK-NEXT:    ret void
  96 //
  97 //
  98 // CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_test_math_long_long_l23
  99 // CHECK-SAME: (i64 noundef [[X:%.*]]) #[[ATTR0]] {
 100 // CHECK-NEXT:  entry:
 101 // CHECK-NEXT:    [[RETVAL_I:%.*]] = alloca i64, align 8, addrspace(5)
 102 // CHECK-NEXT:    [[__X_ADDR_I:%.*]] = alloca i64, align 8, addrspace(5)
 103 // CHECK-NEXT:    [[__SGN_I:%.*]] = alloca i64, align 8, addrspace(5)
 104 // CHECK-NEXT:    [[X_ADDR:%.*]] = alloca i64, align 8, addrspace(5)
 105 // CHECK-NEXT:    [[L1:%.*]] = alloca i64, align 8, addrspace(5)
 106 // CHECK-NEXT:    [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X_ADDR]] to ptr
 107 // CHECK-NEXT:    [[L1_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[L1]] to ptr
 108 // CHECK-NEXT:    store i64 [[X]], ptr [[X_ADDR_ASCAST]], align 8
 109 // CHECK-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr addrspacecast (ptr addrspace(1) @[[GLOB1]] to ptr), i8 1, i1 true)
 110 // CHECK-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
 111 // CHECK-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
 112 // CHECK:       user_code.entry:
 113 // CHECK-NEXT:    [[TMP1:%.*]] = load i64, ptr [[X_ADDR_ASCAST]], align 8
 114 // CHECK-NEXT:    [[RETVAL_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL_I]] to ptr
 115 // CHECK-NEXT:    [[__X_ADDR_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[__X_ADDR_I]] to ptr
 116 // CHECK-NEXT:    [[__SGN_ASCAST_I:%.*]] = addrspacecast ptr addrspace(5) [[__SGN_I]] to ptr
 117 // CHECK-NEXT:    store i64 [[TMP1]], ptr [[__X_ADDR_ASCAST_I]], align 8
 118 // CHECK-NEXT:    [[TMP2:%.*]] = load i64, ptr [[__X_ADDR_ASCAST_I]], align 8
 119 // CHECK-NEXT:    [[SHR_I:%.*]] = ashr i64 [[TMP2]], 63
 120 // CHECK-NEXT:    store i64 [[SHR_I]], ptr [[__SGN_ASCAST_I]], align 8
 121 // CHECK-NEXT:    [[TMP3:%.*]] = load i64, ptr [[__X_ADDR_ASCAST_I]], align 8
 122 // CHECK-NEXT:    [[TMP4:%.*]] = load i64, ptr [[__SGN_ASCAST_I]], align 8
 123 // CHECK-NEXT:    [[XOR_I:%.*]] = xor i64 [[TMP3]], [[TMP4]]
 124 // CHECK-NEXT:    [[TMP5:%.*]] = load i64, ptr [[__SGN_ASCAST_I]], align 8
 125 // CHECK-NEXT:    [[SUB_I:%.*]] = sub nsw i64 [[XOR_I]], [[TMP5]]
 126 // CHECK-NEXT:    store i64 [[SUB_I]], ptr [[L1_ASCAST]], align 8
 127 // CHECK-NEXT:    call void @__kmpc_target_deinit(ptr addrspacecast (ptr addrspace(1) @[[GLOB1]] to ptr), i8 1)
 128 // CHECK-NEXT:    ret void
 129 // CHECK:       worker.exit:
 130 // CHECK-NEXT:    ret void
 131 //