clang/test/CodeGenCUDA/builtins-amdgcn.cu

   1 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
   2 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx906 -x hip \
   3 // RUN:  -aux-triple x86_64-unknown-linux-gnu -fcuda-is-device -emit-llvm %s \
   4 // RUN:  -o - | FileCheck %s
   5
   6 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx906 -x hip \
   7 // RUN:  -aux-triple x86_64-pc-windows-msvc -fcuda-is-device -emit-llvm %s \
   8 // RUN:  -o - | FileCheck %s
   9
  10 #include "Inputs/cuda.h"
  11
  12 // CHECK-LABEL: @_Z16use_dispatch_ptrPi(
  13 // CHECK-NEXT:  entry:
  14 // CHECK-NEXT:    [[OUT:%.*]] = alloca i32*, align 8, addrspace(5)
  15 // CHECK-NEXT:    [[OUT_ADDR:%.*]] = alloca i32*, align 8, addrspace(5)
  16 // CHECK-NEXT:    [[DISPATCH_PTR:%.*]] = alloca i32*, align 8, addrspace(5)
  17 // CHECK-NEXT:    [[OUT_ASCAST:%.*]] = addrspacecast i32* addrspace(5)* [[OUT]] to i32**
  18 // CHECK-NEXT:    [[OUT_ADDR_ASCAST:%.*]] = addrspacecast i32* addrspace(5)* [[OUT_ADDR]] to i32**
  19 // CHECK-NEXT:    [[DISPATCH_PTR_ASCAST:%.*]] = addrspacecast i32* addrspace(5)* [[DISPATCH_PTR]] to i32**
  20 // CHECK-NEXT:    [[TMP0:%.*]] = addrspacecast i32 addrspace(1)* [[OUT_COERCE:%.*]] to i32*
  21 // CHECK-NEXT:    store i32* [[TMP0]], i32** [[OUT_ASCAST]], align 8
  22 // CHECK-NEXT:    [[OUT1:%.*]] = load i32*, i32** [[OUT_ASCAST]], align 8
  23 // CHECK-NEXT:    store i32* [[OUT1]], i32** [[OUT_ADDR_ASCAST]], align 8
  24 // CHECK-NEXT:    [[TMP1:%.*]] = call align 4 dereferenceable(64) i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr()
  25 // CHECK-NEXT:    [[TMP2:%.*]] = addrspacecast i8 addrspace(4)* [[TMP1]] to i32*
  26 // CHECK-NEXT:    store i32* [[TMP2]], i32** [[DISPATCH_PTR_ASCAST]], align 8
  27 // CHECK-NEXT:    [[TMP3:%.*]] = load i32*, i32** [[DISPATCH_PTR_ASCAST]], align 8
  28 // CHECK-NEXT:    [[TMP4:%.*]] = load i32, i32* [[TMP3]], align 4
  29 // CHECK-NEXT:    [[TMP5:%.*]] = load i32*, i32** [[OUT_ADDR_ASCAST]], align 8
  30 // CHECK-NEXT:    store i32 [[TMP4]], i32* [[TMP5]], align 4
  31 // CHECK-NEXT:    ret void
  32 //
  33 __global__ void use_dispatch_ptr(int* out) {
  34   const int* dispatch_ptr = (const int*)__builtin_amdgcn_dispatch_ptr();
  35   *out = *dispatch_ptr;
  36 }
  37
  38 __global__
  39     // CHECK-LABEL: @_Z12test_ds_fmaxf(
  40     // CHECK-NEXT:  entry:
  41     // CHECK-NEXT:    [[SRC_ADDR:%.*]] = alloca float, align 4, addrspace(5)
  42     // CHECK-NEXT:    [[X:%.*]] = alloca float, align 4, addrspace(5)
  43     // CHECK-NEXT:    [[SRC_ADDR_ASCAST:%.*]] = addrspacecast float addrspace(5)* [[SRC_ADDR]] to float*
  44     // CHECK-NEXT:    [[X_ASCAST:%.*]] = addrspacecast float addrspace(5)* [[X]] to float*
  45     // CHECK-NEXT:    store float [[SRC:%.*]], float* [[SRC_ADDR_ASCAST]], align 4
  46     // CHECK-NEXT:    [[TMP0:%.*]] = load float, float* [[SRC_ADDR_ASCAST]], align 4
  47     // CHECK-NEXT:    [[TMP1:%.*]] = call contract float @llvm.amdgcn.ds.fmax.f32(float addrspace(3)* @_ZZ12test_ds_fmaxfE6shared, float [[TMP0]], i32 0, i32 0, i1 false)
  48     // CHECK-NEXT:    store volatile float [[TMP1]], float* [[X_ASCAST]], align 4
  49     // CHECK-NEXT:    ret void
  50     //
  51     void
  52     test_ds_fmax(float src) {
  53   __shared__ float shared;
  54   volatile float x = __builtin_amdgcn_ds_fmaxf(&shared, src, 0, 0, false);
  55 }
  56
  57 // CHECK-LABEL: @_Z12test_ds_faddf(
  58 // CHECK-NEXT:  entry:
  59 // CHECK-NEXT:    [[SRC_ADDR:%.*]] = alloca float, align 4, addrspace(5)
  60 // CHECK-NEXT:    [[X:%.*]] = alloca float, align 4, addrspace(5)
  61 // CHECK-NEXT:    [[SRC_ADDR_ASCAST:%.*]] = addrspacecast float addrspace(5)* [[SRC_ADDR]] to float*
  62 // CHECK-NEXT:    [[X_ASCAST:%.*]] = addrspacecast float addrspace(5)* [[X]] to float*
  63 // CHECK-NEXT:    store float [[SRC:%.*]], float* [[SRC_ADDR_ASCAST]], align 4
  64 // CHECK-NEXT:    [[TMP0:%.*]] = load float, float* [[SRC_ADDR_ASCAST]], align 4
  65 // CHECK-NEXT:    [[TMP1:%.*]] = call contract float @llvm.amdgcn.ds.fadd.f32(float addrspace(3)* @_ZZ12test_ds_faddfE6shared, float [[TMP0]], i32 0, i32 0, i1 false)
  66 // CHECK-NEXT:    store volatile float [[TMP1]], float* [[X_ASCAST]], align 4
  67 // CHECK-NEXT:    ret void
  68 //
  69 __global__ void test_ds_fadd(float src) {
  70   __shared__ float shared;
  71   volatile float x = __builtin_amdgcn_ds_faddf(&shared, src, 0, 0, false);
  72 }
  73
  74 // CHECK-LABEL: @_Z12test_ds_fminfPf(
  75 // CHECK-NEXT:  entry:
  76 // CHECK-NEXT:    [[SHARED:%.*]] = alloca float*, align 8, addrspace(5)
  77 // CHECK-NEXT:    [[SRC_ADDR:%.*]] = alloca float, align 4, addrspace(5)
  78 // CHECK-NEXT:    [[SHARED_ADDR:%.*]] = alloca float*, align 8, addrspace(5)
  79 // CHECK-NEXT:    [[X:%.*]] = alloca float, align 4, addrspace(5)
  80 // CHECK-NEXT:    [[SHARED_ASCAST:%.*]] = addrspacecast float* addrspace(5)* [[SHARED]] to float**
  81 // CHECK-NEXT:    [[SRC_ADDR_ASCAST:%.*]] = addrspacecast float addrspace(5)* [[SRC_ADDR]] to float*
  82 // CHECK-NEXT:    [[SHARED_ADDR_ASCAST:%.*]] = addrspacecast float* addrspace(5)* [[SHARED_ADDR]] to float**
  83 // CHECK-NEXT:    [[X_ASCAST:%.*]] = addrspacecast float addrspace(5)* [[X]] to float*
  84 // CHECK-NEXT:    [[TMP0:%.*]] = addrspacecast float addrspace(1)* [[SHARED_COERCE:%.*]] to float*
  85 // CHECK-NEXT:    store float* [[TMP0]], float** [[SHARED_ASCAST]], align 8
  86 // CHECK-NEXT:    [[SHARED1:%.*]] = load float*, float** [[SHARED_ASCAST]], align 8
  87 // CHECK-NEXT:    store float [[SRC:%.*]], float* [[SRC_ADDR_ASCAST]], align 4
  88 // CHECK-NEXT:    store float* [[SHARED1]], float** [[SHARED_ADDR_ASCAST]], align 8
  89 // CHECK-NEXT:    [[TMP1:%.*]] = load float*, float** [[SHARED_ADDR_ASCAST]], align 8
  90 // CHECK-NEXT:    [[TMP2:%.*]] = addrspacecast float* [[TMP1]] to float addrspace(3)*
  91 // CHECK-NEXT:    [[TMP3:%.*]] = load float, float* [[SRC_ADDR_ASCAST]], align 4
  92 // CHECK-NEXT:    [[TMP4:%.*]] = call contract float @llvm.amdgcn.ds.fmin.f32(float addrspace(3)* [[TMP2]], float [[TMP3]], i32 0, i32 0, i1 false)
  93 // CHECK-NEXT:    store volatile float [[TMP4]], float* [[X_ASCAST]], align 4
  94 // CHECK-NEXT:    ret void
  95 //
  96 __global__ void test_ds_fmin(float src, float *shared) {
  97   volatile float x = __builtin_amdgcn_ds_fminf(shared, src, 0, 0, false);
  98 }
  99
 100 // CHECK-LABEL: @_Z33test_ret_builtin_nondef_addrspacev(
 101 // CHECK-NEXT:  entry:
 102 // CHECK-NEXT:    [[X:%.*]] = alloca i8*, align 8, addrspace(5)
 103 // CHECK-NEXT:    [[X_ASCAST:%.*]] = addrspacecast i8* addrspace(5)* [[X]] to i8**
 104 // CHECK-NEXT:    [[TMP0:%.*]] = call align 4 dereferenceable(64) i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr()
 105 // CHECK-NEXT:    [[TMP1:%.*]] = addrspacecast i8 addrspace(4)* [[TMP0]] to i8*
 106 // CHECK-NEXT:    store i8* [[TMP1]], i8** [[X_ASCAST]], align 8
 107 // CHECK-NEXT:    ret void
 108 //
 109 __device__ void test_ret_builtin_nondef_addrspace() {
 110   void *x = __builtin_amdgcn_dispatch_ptr();
 111 }
 112
 113 // CHECK-LABEL: @_Z6endpgmv(
 114 // CHECK-NEXT:  entry:
 115 // CHECK-NEXT:    call void @llvm.amdgcn.endpgm()
 116 // CHECK-NEXT:    ret void
 117 //
 118 __global__ void endpgm() {
 119   __builtin_amdgcn_endpgm();
 120 }
 121
 122 // Check the 64 bit argument is correctly passed to the intrinsic without truncation or assertion.
 123
 124 // CHECK-LABEL: @_Z14test_uicmp_i64Pyyy(
 125 // CHECK-NEXT:  entry:
 126 // CHECK-NEXT:    [[OUT:%.*]] = alloca i64*, align 8, addrspace(5)
 127 // CHECK-NEXT:    [[OUT_ADDR:%.*]] = alloca i64*, align 8, addrspace(5)
 128 // CHECK-NEXT:    [[A_ADDR:%.*]] = alloca i64, align 8, addrspace(5)
 129 // CHECK-NEXT:    [[B_ADDR:%.*]] = alloca i64, align 8, addrspace(5)
 130 // CHECK-NEXT:    [[OUT_ASCAST:%.*]] = addrspacecast i64* addrspace(5)* [[OUT]] to i64**
 131 // CHECK-NEXT:    [[OUT_ADDR_ASCAST:%.*]] = addrspacecast i64* addrspace(5)* [[OUT_ADDR]] to i64**
 132 // CHECK-NEXT:    [[A_ADDR_ASCAST:%.*]] = addrspacecast i64 addrspace(5)* [[A_ADDR]] to i64*
 133 // CHECK-NEXT:    [[B_ADDR_ASCAST:%.*]] = addrspacecast i64 addrspace(5)* [[B_ADDR]] to i64*
 134 // CHECK-NEXT:    [[TMP0:%.*]] = addrspacecast i64 addrspace(1)* [[OUT_COERCE:%.*]] to i64*
 135 // CHECK-NEXT:    store i64* [[TMP0]], i64** [[OUT_ASCAST]], align 8
 136 // CHECK-NEXT:    [[OUT1:%.*]] = load i64*, i64** [[OUT_ASCAST]], align 8
 137 // CHECK-NEXT:    store i64* [[OUT1]], i64** [[OUT_ADDR_ASCAST]], align 8
 138 // CHECK-NEXT:    store i64 [[A:%.*]], i64* [[A_ADDR_ASCAST]], align 8
 139 // CHECK-NEXT:    store i64 [[B:%.*]], i64* [[B_ADDR_ASCAST]], align 8
 140 // CHECK-NEXT:    [[TMP1:%.*]] = load i64, i64* [[A_ADDR_ASCAST]], align 8
 141 // CHECK-NEXT:    [[TMP2:%.*]] = load i64, i64* [[B_ADDR_ASCAST]], align 8
 142 // CHECK-NEXT:    [[TMP3:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i64(i64 [[TMP1]], i64 [[TMP2]], i32 35)
 143 // CHECK-NEXT:    [[TMP4:%.*]] = load i64*, i64** [[OUT_ADDR_ASCAST]], align 8
 144 // CHECK-NEXT:    store i64 [[TMP3]], i64* [[TMP4]], align 8
 145 // CHECK-NEXT:    ret void
 146 //
 147 __global__ void test_uicmp_i64(unsigned long long *out, unsigned long long a, unsigned long long b)
 148 {
 149   *out = __builtin_amdgcn_uicmpl(a, b, 30+5);
 150 }
 151
 152 // Check the 64 bit return value is correctly returned without truncation or assertion.
 153
 154 // CHECK-LABEL: @_Z14test_s_memtimePy(
 155 // CHECK-NEXT:  entry:
 156 // CHECK-NEXT:    [[OUT:%.*]] = alloca i64*, align 8, addrspace(5)
 157 // CHECK-NEXT:    [[OUT_ADDR:%.*]] = alloca i64*, align 8, addrspace(5)
 158 // CHECK-NEXT:    [[OUT_ASCAST:%.*]] = addrspacecast i64* addrspace(5)* [[OUT]] to i64**
 159 // CHECK-NEXT:    [[OUT_ADDR_ASCAST:%.*]] = addrspacecast i64* addrspace(5)* [[OUT_ADDR]] to i64**
 160 // CHECK-NEXT:    [[TMP0:%.*]] = addrspacecast i64 addrspace(1)* [[OUT_COERCE:%.*]] to i64*
 161 // CHECK-NEXT:    store i64* [[TMP0]], i64** [[OUT_ASCAST]], align 8
 162 // CHECK-NEXT:    [[OUT1:%.*]] = load i64*, i64** [[OUT_ASCAST]], align 8
 163 // CHECK-NEXT:    store i64* [[OUT1]], i64** [[OUT_ADDR_ASCAST]], align 8
 164 // CHECK-NEXT:    [[TMP1:%.*]] = call i64 @llvm.amdgcn.s.memtime()
 165 // CHECK-NEXT:    [[TMP2:%.*]] = load i64*, i64** [[OUT_ADDR_ASCAST]], align 8
 166 // CHECK-NEXT:    store i64 [[TMP1]], i64* [[TMP2]], align 8
 167 // CHECK-NEXT:    ret void
 168 //
 169 __global__ void test_s_memtime(unsigned long long* out)
 170 {
 171   *out = __builtin_amdgcn_s_memtime();
 172 }
 173
 174 // Check a generic pointer can be passed as a shared pointer and a generic pointer.
 175 __device__ void func(float *x);
 176
 177 // CHECK-LABEL: @_Z17test_ds_fmin_funcfPf(
 178 // CHECK-NEXT:  entry:
 179 // CHECK-NEXT:    [[SHARED:%.*]] = alloca float*, align 8, addrspace(5)
 180 // CHECK-NEXT:    [[SRC_ADDR:%.*]] = alloca float, align 4, addrspace(5)
 181 // CHECK-NEXT:    [[SHARED_ADDR:%.*]] = alloca float*, align 8, addrspace(5)
 182 // CHECK-NEXT:    [[X:%.*]] = alloca float, align 4, addrspace(5)
 183 // CHECK-NEXT:    [[SHARED_ASCAST:%.*]] = addrspacecast float* addrspace(5)* [[SHARED]] to float**
 184 // CHECK-NEXT:    [[SRC_ADDR_ASCAST:%.*]] = addrspacecast float addrspace(5)* [[SRC_ADDR]] to float*
 185 // CHECK-NEXT:    [[SHARED_ADDR_ASCAST:%.*]] = addrspacecast float* addrspace(5)* [[SHARED_ADDR]] to float**
 186 // CHECK-NEXT:    [[X_ASCAST:%.*]] = addrspacecast float addrspace(5)* [[X]] to float*
 187 // CHECK-NEXT:    [[TMP0:%.*]] = addrspacecast float addrspace(1)* [[SHARED_COERCE:%.*]] to float*
 188 // CHECK-NEXT:    store float* [[TMP0]], float** [[SHARED_ASCAST]], align 8
 189 // CHECK-NEXT:    [[SHARED1:%.*]] = load float*, float** [[SHARED_ASCAST]], align 8
 190 // CHECK-NEXT:    store float [[SRC:%.*]], float* [[SRC_ADDR_ASCAST]], align 4
 191 // CHECK-NEXT:    store float* [[SHARED1]], float** [[SHARED_ADDR_ASCAST]], align 8
 192 // CHECK-NEXT:    [[TMP1:%.*]] = load float*, float** [[SHARED_ADDR_ASCAST]], align 8
 193 // CHECK-NEXT:    [[TMP2:%.*]] = addrspacecast float* [[TMP1]] to float addrspace(3)*
 194 // CHECK-NEXT:    [[TMP3:%.*]] = load float, float* [[SRC_ADDR_ASCAST]], align 4
 195 // CHECK-NEXT:    [[TMP4:%.*]] = call contract float @llvm.amdgcn.ds.fmin.f32(float addrspace(3)* [[TMP2]], float [[TMP3]], i32 0, i32 0, i1 false)
 196 // CHECK-NEXT:    store volatile float [[TMP4]], float* [[X_ASCAST]], align 4
 197 // CHECK-NEXT:    [[TMP5:%.*]] = load float*, float** [[SHARED_ADDR_ASCAST]], align 8
 198 // CHECK-NEXT:    call void @_Z4funcPf(float* noundef [[TMP5]]) #[[ATTR8:[0-9]+]]
 199 // CHECK-NEXT:    ret void
 200 //
 201 __global__ void test_ds_fmin_func(float src, float *__restrict shared) {
 202   volatile float x = __builtin_amdgcn_ds_fminf(shared, src, 0, 0, false);
 203   func(shared);
 204 }
 205
 206 // CHECK-LABEL: @_Z14test_is_sharedPf(
 207 // CHECK-NEXT:  entry:
 208 // CHECK-NEXT:    [[X:%.*]] = alloca float*, align 8, addrspace(5)
 209 // CHECK-NEXT:    [[X_ADDR:%.*]] = alloca float*, align 8, addrspace(5)
 210 // CHECK-NEXT:    [[RET:%.*]] = alloca i8, align 1, addrspace(5)
 211 // CHECK-NEXT:    [[X_ASCAST:%.*]] = addrspacecast float* addrspace(5)* [[X]] to float**
 212 // CHECK-NEXT:    [[X_ADDR_ASCAST:%.*]] = addrspacecast float* addrspace(5)* [[X_ADDR]] to float**
 213 // CHECK-NEXT:    [[RET_ASCAST:%.*]] = addrspacecast i8 addrspace(5)* [[RET]] to i8*
 214 // CHECK-NEXT:    [[TMP0:%.*]] = addrspacecast float addrspace(1)* [[X_COERCE:%.*]] to float*
 215 // CHECK-NEXT:    store float* [[TMP0]], float** [[X_ASCAST]], align 8
 216 // CHECK-NEXT:    [[X1:%.*]] = load float*, float** [[X_ASCAST]], align 8
 217 // CHECK-NEXT:    store float* [[X1]], float** [[X_ADDR_ASCAST]], align 8
 218 // CHECK-NEXT:    [[TMP1:%.*]] = load float*, float** [[X_ADDR_ASCAST]], align 8
 219 // CHECK-NEXT:    [[TMP2:%.*]] = bitcast float* [[TMP1]] to i8*
 220 // CHECK-NEXT:    [[TMP3:%.*]] = call i1 @llvm.amdgcn.is.shared(i8* [[TMP2]])
 221 // CHECK-NEXT:    [[FROMBOOL:%.*]] = zext i1 [[TMP3]] to i8
 222 // CHECK-NEXT:    store i8 [[FROMBOOL]], i8* [[RET_ASCAST]], align 1
 223 // CHECK-NEXT:    ret void
 224 //
 225 __global__ void test_is_shared(float *x){
 226   bool ret = __builtin_amdgcn_is_shared(x);
 227 }