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 ptr, align 8, addrspace(5)
  15 // CHECK-NEXT:    [[OUT_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
  16 // CHECK-NEXT:    [[DISPATCH_PTR:%.*]] = alloca ptr, align 8, addrspace(5)
  17 // CHECK-NEXT:    [[OUT_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT]] to ptr
  18 // CHECK-NEXT:    [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr
  19 // CHECK-NEXT:    [[DISPATCH_PTR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DISPATCH_PTR]] to ptr
  20 // CHECK-NEXT:    [[TMP0:%.*]] = addrspacecast ptr addrspace(1) [[OUT_COERCE:%.*]] to ptr
  21 // CHECK-NEXT:    store ptr [[TMP0]], ptr [[OUT_ASCAST]], align 8
  22 // CHECK-NEXT:    [[OUT1:%.*]] = load ptr, ptr [[OUT_ASCAST]], align 8
  23 // CHECK-NEXT:    store ptr [[OUT1]], ptr [[OUT_ADDR_ASCAST]], align 8
  24 // CHECK-NEXT:    [[TMP1:%.*]] = call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
  25 // CHECK-NEXT:    [[TMP2:%.*]] = addrspacecast ptr addrspace(4) [[TMP1]] to ptr
  26 // CHECK-NEXT:    store ptr [[TMP2]], ptr [[DISPATCH_PTR_ASCAST]], align 8
  27 // CHECK-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DISPATCH_PTR_ASCAST]], align 8
  28 // CHECK-NEXT:    [[TMP4:%.*]] = load i32, ptr [[TMP3]], align 4
  29 // CHECK-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[OUT_ADDR_ASCAST]], align 8
  30 // CHECK-NEXT:    store i32 [[TMP4]], ptr [[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 // CHECK-LABEL: @_Z13use_queue_ptrPi(
  39 // CHECK-NEXT:  entry:
  40 // CHECK-NEXT:    [[OUT:%.*]] = alloca ptr, align 8, addrspace(5)
  41 // CHECK-NEXT:    [[OUT_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
  42 // CHECK-NEXT:    [[QUEUE_PTR:%.*]] = alloca ptr, align 8, addrspace(5)
  43 // CHECK-NEXT:    [[OUT_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT]] to ptr
  44 // CHECK-NEXT:    [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr
  45 // CHECK-NEXT:    [[QUEUE_PTR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[QUEUE_PTR]] to ptr
  46 // CHECK-NEXT:    [[TMP0:%.*]] = addrspacecast ptr addrspace(1) [[OUT_COERCE:%.*]] to ptr
  47 // CHECK-NEXT:    store ptr [[TMP0]], ptr [[OUT_ASCAST]], align 8
  48 // CHECK-NEXT:    [[OUT1:%.*]] = load ptr, ptr [[OUT_ASCAST]], align 8
  49 // CHECK-NEXT:    store ptr [[OUT1]], ptr [[OUT_ADDR_ASCAST]], align 8
  50 // CHECK-NEXT:    [[TMP1:%.*]] = call ptr addrspace(4) @llvm.amdgcn.queue.ptr()
  51 // CHECK-NEXT:    [[TMP2:%.*]] = addrspacecast ptr addrspace(4) [[TMP1]] to ptr
  52 // CHECK-NEXT:    store ptr [[TMP2]], ptr [[QUEUE_PTR_ASCAST]], align 8
  53 // CHECK-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[QUEUE_PTR_ASCAST]], align 8
  54 // CHECK-NEXT:    [[TMP4:%.*]] = load i32, ptr [[TMP3]], align 4
  55 // CHECK-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[OUT_ADDR_ASCAST]], align 8
  56 // CHECK-NEXT:    store i32 [[TMP4]], ptr [[TMP5]], align 4
  57 // CHECK-NEXT:    ret void
  58 //
  59 __global__ void use_queue_ptr(int* out) {
  60   const int* queue_ptr = (const int*)__builtin_amdgcn_queue_ptr();
  61   *out = *queue_ptr;
  62 }
  63
  64 // CHECK-LABEL: @_Z19use_implicitarg_ptrPi(
  65 // CHECK-NEXT:  entry:
  66 // CHECK-NEXT:    [[OUT:%.*]] = alloca ptr, align 8, addrspace(5)
  67 // CHECK-NEXT:    [[OUT_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
  68 // CHECK-NEXT:    [[IMPLICITARG_PTR:%.*]] = alloca ptr, align 8, addrspace(5)
  69 // CHECK-NEXT:    [[OUT_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT]] to ptr
  70 // CHECK-NEXT:    [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr
  71 // CHECK-NEXT:    [[IMPLICITARG_PTR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[IMPLICITARG_PTR]] to ptr
  72 // CHECK-NEXT:    [[TMP0:%.*]] = addrspacecast ptr addrspace(1) [[OUT_COERCE:%.*]] to ptr
  73 // CHECK-NEXT:    store ptr [[TMP0]], ptr [[OUT_ASCAST]], align 8
  74 // CHECK-NEXT:    [[OUT1:%.*]] = load ptr, ptr [[OUT_ASCAST]], align 8
  75 // CHECK-NEXT:    store ptr [[OUT1]], ptr [[OUT_ADDR_ASCAST]], align 8
  76 // CHECK-NEXT:    [[TMP1:%.*]] = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
  77 // CHECK-NEXT:    [[TMP2:%.*]] = addrspacecast ptr addrspace(4) [[TMP1]] to ptr
  78 // CHECK-NEXT:    store ptr [[TMP2]], ptr [[IMPLICITARG_PTR_ASCAST]], align 8
  79 // CHECK-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[IMPLICITARG_PTR_ASCAST]], align 8
  80 // CHECK-NEXT:    [[TMP4:%.*]] = load i32, ptr [[TMP3]], align 4
  81 // CHECK-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[OUT_ADDR_ASCAST]], align 8
  82 // CHECK-NEXT:    store i32 [[TMP4]], ptr [[TMP5]], align 4
  83 // CHECK-NEXT:    ret void
  84 //
  85 __global__ void use_implicitarg_ptr(int* out) {
  86   const int* implicitarg_ptr = (const int*)__builtin_amdgcn_implicitarg_ptr();
  87   *out = *implicitarg_ptr;
  88 }
  89
  90 __global__
  91     //
  92     void
  93 // CHECK-LABEL: @_Z12test_ds_fmaxf(
  94 // CHECK-NEXT:  entry:
  95 // CHECK-NEXT:    [[SRC_ADDR:%.*]] = alloca float, align 4, addrspace(5)
  96 // CHECK-NEXT:    [[X:%.*]] = alloca float, align 4, addrspace(5)
  97 // CHECK-NEXT:    [[SRC_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SRC_ADDR]] to ptr
  98 // CHECK-NEXT:    [[X_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X]] to ptr
  99 // CHECK-NEXT:    store float [[SRC:%.*]], ptr [[SRC_ADDR_ASCAST]], align 4
 100 // CHECK-NEXT:    [[TMP0:%.*]] = load float, ptr [[SRC_ADDR_ASCAST]], align 4
 101 // CHECK-NEXT:    [[TMP1:%.*]] = call contract float @llvm.amdgcn.ds.fmax.f32(ptr addrspace(3) @_ZZ12test_ds_fmaxfE6shared, float [[TMP0]], i32 0, i32 0, i1 false)
 102 // CHECK-NEXT:    store volatile float [[TMP1]], ptr [[X_ASCAST]], align 4
 103 // CHECK-NEXT:    ret void
 104 //
 105     test_ds_fmax(float src) {
 106   __shared__ float shared;
 107   volatile float x = __builtin_amdgcn_ds_fmaxf(&shared, src, 0, 0, false);
 108 }
 109
 110 // CHECK-LABEL: @_Z12test_ds_faddf(
 111 // CHECK-NEXT:  entry:
 112 // CHECK-NEXT:    [[SRC_ADDR:%.*]] = alloca float, align 4, addrspace(5)
 113 // CHECK-NEXT:    [[X:%.*]] = alloca float, align 4, addrspace(5)
 114 // CHECK-NEXT:    [[SRC_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SRC_ADDR]] to ptr
 115 // CHECK-NEXT:    [[X_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X]] to ptr
 116 // CHECK-NEXT:    store float [[SRC:%.*]], ptr [[SRC_ADDR_ASCAST]], align 4
 117 // CHECK-NEXT:    [[TMP0:%.*]] = load float, ptr [[SRC_ADDR_ASCAST]], align 4
 118 // CHECK-NEXT:    [[TMP1:%.*]] = call contract float @llvm.amdgcn.ds.fadd.f32(ptr addrspace(3) @_ZZ12test_ds_faddfE6shared, float [[TMP0]], i32 0, i32 0, i1 false)
 119 // CHECK-NEXT:    store volatile float [[TMP1]], ptr [[X_ASCAST]], align 4
 120 // CHECK-NEXT:    ret void
 121 //
 122 __global__ void test_ds_fadd(float src) {
 123   __shared__ float shared;
 124   volatile float x = __builtin_amdgcn_ds_faddf(&shared, src, 0, 0, false);
 125 }
 126
 127 // CHECK-LABEL: @_Z12test_ds_fminfPf(
 128 // CHECK-NEXT:  entry:
 129 // CHECK-NEXT:    [[SHARED:%.*]] = alloca ptr, align 8, addrspace(5)
 130 // CHECK-NEXT:    [[SRC_ADDR:%.*]] = alloca float, align 4, addrspace(5)
 131 // CHECK-NEXT:    [[SHARED_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
 132 // CHECK-NEXT:    [[X:%.*]] = alloca float, align 4, addrspace(5)
 133 // CHECK-NEXT:    [[SHARED_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SHARED]] to ptr
 134 // CHECK-NEXT:    [[SRC_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SRC_ADDR]] to ptr
 135 // CHECK-NEXT:    [[SHARED_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SHARED_ADDR]] to ptr
 136 // CHECK-NEXT:    [[X_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X]] to ptr
 137 // CHECK-NEXT:    [[TMP0:%.*]] = addrspacecast ptr addrspace(1) [[SHARED_COERCE:%.*]] to ptr
 138 // CHECK-NEXT:    store ptr [[TMP0]], ptr [[SHARED_ASCAST]], align 8
 139 // CHECK-NEXT:    [[SHARED1:%.*]] = load ptr, ptr [[SHARED_ASCAST]], align 8
 140 // CHECK-NEXT:    store float [[SRC:%.*]], ptr [[SRC_ADDR_ASCAST]], align 4
 141 // CHECK-NEXT:    store ptr [[SHARED1]], ptr [[SHARED_ADDR_ASCAST]], align 8
 142 // CHECK-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[SHARED_ADDR_ASCAST]], align 8
 143 // CHECK-NEXT:    [[TMP2:%.*]] = addrspacecast ptr [[TMP1]] to ptr addrspace(3)
 144 // CHECK-NEXT:    [[TMP3:%.*]] = load float, ptr [[SRC_ADDR_ASCAST]], align 4
 145 // CHECK-NEXT:    [[TMP4:%.*]] = call contract float @llvm.amdgcn.ds.fmin.f32(ptr addrspace(3) [[TMP2]], float [[TMP3]], i32 0, i32 0, i1 false)
 146 // CHECK-NEXT:    store volatile float [[TMP4]], ptr [[X_ASCAST]], align 4
 147 // CHECK-NEXT:    ret void
 148 //
 149 __global__ void test_ds_fmin(float src, float *shared) {
 150   volatile float x = __builtin_amdgcn_ds_fminf(shared, src, 0, 0, false);
 151 }
 152
 153 // CHECK-LABEL: @_Z33test_ret_builtin_nondef_addrspacev(
 154 // CHECK-NEXT:  entry:
 155 // CHECK-NEXT:    [[X:%.*]] = alloca ptr, align 8, addrspace(5)
 156 // CHECK-NEXT:    [[X_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X]] to ptr
 157 // CHECK-NEXT:    [[TMP0:%.*]] = call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
 158 // CHECK-NEXT:    [[TMP1:%.*]] = addrspacecast ptr addrspace(4) [[TMP0]] to ptr
 159 // CHECK-NEXT:    store ptr [[TMP1]], ptr [[X_ASCAST]], align 8
 160 // CHECK-NEXT:    ret void
 161 //
 162 __device__ void test_ret_builtin_nondef_addrspace() {
 163   void *x = __builtin_amdgcn_dispatch_ptr();
 164 }
 165
 166 // CHECK-LABEL: @_Z6endpgmv(
 167 // CHECK-NEXT:  entry:
 168 // CHECK-NEXT:    call void @llvm.amdgcn.endpgm()
 169 // CHECK-NEXT:    ret void
 170 //
 171 __global__ void endpgm() {
 172   __builtin_amdgcn_endpgm();
 173 }
 174
 175 // Check the 64 bit argument is correctly passed to the intrinsic without truncation or assertion.
 176
 177 // CHECK-LABEL: @_Z14test_uicmp_i64Pyyy(
 178 // CHECK-NEXT:  entry:
 179 // CHECK-NEXT:    [[OUT:%.*]] = alloca ptr, align 8, addrspace(5)
 180 // CHECK-NEXT:    [[OUT_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
 181 // CHECK-NEXT:    [[A_ADDR:%.*]] = alloca i64, align 8, addrspace(5)
 182 // CHECK-NEXT:    [[B_ADDR:%.*]] = alloca i64, align 8, addrspace(5)
 183 // CHECK-NEXT:    [[OUT_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT]] to ptr
 184 // CHECK-NEXT:    [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr
 185 // CHECK-NEXT:    [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
 186 // CHECK-NEXT:    [[B_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[B_ADDR]] to ptr
 187 // CHECK-NEXT:    [[TMP0:%.*]] = addrspacecast ptr addrspace(1) [[OUT_COERCE:%.*]] to ptr
 188 // CHECK-NEXT:    store ptr [[TMP0]], ptr [[OUT_ASCAST]], align 8
 189 // CHECK-NEXT:    [[OUT1:%.*]] = load ptr, ptr [[OUT_ASCAST]], align 8
 190 // CHECK-NEXT:    store ptr [[OUT1]], ptr [[OUT_ADDR_ASCAST]], align 8
 191 // CHECK-NEXT:    store i64 [[A:%.*]], ptr [[A_ADDR_ASCAST]], align 8
 192 // CHECK-NEXT:    store i64 [[B:%.*]], ptr [[B_ADDR_ASCAST]], align 8
 193 // CHECK-NEXT:    [[TMP1:%.*]] = load i64, ptr [[A_ADDR_ASCAST]], align 8
 194 // CHECK-NEXT:    [[TMP2:%.*]] = load i64, ptr [[B_ADDR_ASCAST]], align 8
 195 // CHECK-NEXT:    [[TMP3:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i64(i64 [[TMP1]], i64 [[TMP2]], i32 35)
 196 // CHECK-NEXT:    [[TMP4:%.*]] = load ptr, ptr [[OUT_ADDR_ASCAST]], align 8
 197 // CHECK-NEXT:    store i64 [[TMP3]], ptr [[TMP4]], align 8
 198 // CHECK-NEXT:    ret void
 199 //
 200 __global__ void test_uicmp_i64(unsigned long long *out, unsigned long long a, unsigned long long b)
 201 {
 202   *out = __builtin_amdgcn_uicmpl(a, b, 30+5);
 203 }
 204
 205 // Check the 64 bit return value is correctly returned without truncation or assertion.
 206
 207 // CHECK-LABEL: @_Z14test_s_memtimePy(
 208 // CHECK-NEXT:  entry:
 209 // CHECK-NEXT:    [[OUT:%.*]] = alloca ptr, align 8, addrspace(5)
 210 // CHECK-NEXT:    [[OUT_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
 211 // CHECK-NEXT:    [[OUT_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT]] to ptr
 212 // CHECK-NEXT:    [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr
 213 // CHECK-NEXT:    [[TMP0:%.*]] = addrspacecast ptr addrspace(1) [[OUT_COERCE:%.*]] to ptr
 214 // CHECK-NEXT:    store ptr [[TMP0]], ptr [[OUT_ASCAST]], align 8
 215 // CHECK-NEXT:    [[OUT1:%.*]] = load ptr, ptr [[OUT_ASCAST]], align 8
 216 // CHECK-NEXT:    store ptr [[OUT1]], ptr [[OUT_ADDR_ASCAST]], align 8
 217 // CHECK-NEXT:    [[TMP1:%.*]] = call i64 @llvm.amdgcn.s.memtime()
 218 // CHECK-NEXT:    [[TMP2:%.*]] = load ptr, ptr [[OUT_ADDR_ASCAST]], align 8
 219 // CHECK-NEXT:    store i64 [[TMP1]], ptr [[TMP2]], align 8
 220 // CHECK-NEXT:    ret void
 221 //
 222 __global__ void test_s_memtime(unsigned long long* out)
 223 {
 224   *out = __builtin_amdgcn_s_memtime();
 225 }
 226
 227 // Check a generic pointer can be passed as a shared pointer and a generic pointer.
 228 __device__ void func(float *x);
 229
 230 // CHECK-LABEL: @_Z17test_ds_fmin_funcfPf(
 231 // CHECK-NEXT:  entry:
 232 // CHECK-NEXT:    [[SHARED:%.*]] = alloca ptr, align 8, addrspace(5)
 233 // CHECK-NEXT:    [[SRC_ADDR:%.*]] = alloca float, align 4, addrspace(5)
 234 // CHECK-NEXT:    [[SHARED_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
 235 // CHECK-NEXT:    [[X:%.*]] = alloca float, align 4, addrspace(5)
 236 // CHECK-NEXT:    [[SHARED_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SHARED]] to ptr
 237 // CHECK-NEXT:    [[SRC_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SRC_ADDR]] to ptr
 238 // CHECK-NEXT:    [[SHARED_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SHARED_ADDR]] to ptr
 239 // CHECK-NEXT:    [[X_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X]] to ptr
 240 // CHECK-NEXT:    [[TMP0:%.*]] = addrspacecast ptr addrspace(1) [[SHARED_COERCE:%.*]] to ptr
 241 // CHECK-NEXT:    store ptr [[TMP0]], ptr [[SHARED_ASCAST]], align 8
 242 // CHECK-NEXT:    [[SHARED1:%.*]] = load ptr, ptr [[SHARED_ASCAST]], align 8
 243 // CHECK-NEXT:    store float [[SRC:%.*]], ptr [[SRC_ADDR_ASCAST]], align 4
 244 // CHECK-NEXT:    store ptr [[SHARED1]], ptr [[SHARED_ADDR_ASCAST]], align 8
 245 // CHECK-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[SHARED_ADDR_ASCAST]], align 8
 246 // CHECK-NEXT:    [[TMP2:%.*]] = addrspacecast ptr [[TMP1]] to ptr addrspace(3)
 247 // CHECK-NEXT:    [[TMP3:%.*]] = load float, ptr [[SRC_ADDR_ASCAST]], align 4
 248 // CHECK-NEXT:    [[TMP4:%.*]] = call contract float @llvm.amdgcn.ds.fmin.f32(ptr addrspace(3) [[TMP2]], float [[TMP3]], i32 0, i32 0, i1 false)
 249 // CHECK-NEXT:    store volatile float [[TMP4]], ptr [[X_ASCAST]], align 4
 250 // CHECK-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[SHARED_ADDR_ASCAST]], align 8
 251 // CHECK-NEXT:    call void @_Z4funcPf(ptr noundef [[TMP5]]) #[[ATTR8:[0-9]+]]
 252 // CHECK-NEXT:    ret void
 253 //
 254 __global__ void test_ds_fmin_func(float src, float *__restrict shared) {
 255   volatile float x = __builtin_amdgcn_ds_fminf(shared, src, 0, 0, false);
 256   func(shared);
 257 }
 258
 259 // CHECK-LABEL: @_Z14test_is_sharedPf(
 260 // CHECK-NEXT:  entry:
 261 // CHECK-NEXT:    [[X:%.*]] = alloca ptr, align 8, addrspace(5)
 262 // CHECK-NEXT:    [[X_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
 263 // CHECK-NEXT:    [[RET:%.*]] = alloca i8, align 1, addrspace(5)
 264 // CHECK-NEXT:    [[X_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X]] to ptr
 265 // CHECK-NEXT:    [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X_ADDR]] to ptr
 266 // CHECK-NEXT:    [[RET_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RET]] to ptr
 267 // CHECK-NEXT:    [[TMP0:%.*]] = addrspacecast ptr addrspace(1) [[X_COERCE:%.*]] to ptr
 268 // CHECK-NEXT:    store ptr [[TMP0]], ptr [[X_ASCAST]], align 8
 269 // CHECK-NEXT:    [[X1:%.*]] = load ptr, ptr [[X_ASCAST]], align 8
 270 // CHECK-NEXT:    store ptr [[X1]], ptr [[X_ADDR_ASCAST]], align 8
 271 // CHECK-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[X_ADDR_ASCAST]], align 8
 272 // CHECK-NEXT:    [[TMP2:%.*]] = call i1 @llvm.amdgcn.is.shared(ptr [[TMP1]])
 273 // CHECK-NEXT:    [[FROMBOOL:%.*]] = zext i1 [[TMP2]] to i8
 274 // CHECK-NEXT:    store i8 [[FROMBOOL]], ptr [[RET_ASCAST]], align 1
 275 // CHECK-NEXT:    ret void
 276 //
 277 __global__ void test_is_shared(float *x){
 278   bool ret = __builtin_amdgcn_is_shared(x);
 279 }
 280
 281 // CHECK-LABEL: @_Z15test_is_privatePi(
 282 // CHECK-NEXT:  entry:
 283 // CHECK-NEXT:    [[X:%.*]] = alloca ptr, align 8, addrspace(5)
 284 // CHECK-NEXT:    [[X_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
 285 // CHECK-NEXT:    [[RET:%.*]] = alloca i8, align 1, addrspace(5)
 286 // CHECK-NEXT:    [[X_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X]] to ptr
 287 // CHECK-NEXT:    [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X_ADDR]] to ptr
 288 // CHECK-NEXT:    [[RET_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RET]] to ptr
 289 // CHECK-NEXT:    [[TMP0:%.*]] = addrspacecast ptr addrspace(1) [[X_COERCE:%.*]] to ptr
 290 // CHECK-NEXT:    store ptr [[TMP0]], ptr [[X_ASCAST]], align 8
 291 // CHECK-NEXT:    [[X1:%.*]] = load ptr, ptr [[X_ASCAST]], align 8
 292 // CHECK-NEXT:    store ptr [[X1]], ptr [[X_ADDR_ASCAST]], align 8
 293 // CHECK-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[X_ADDR_ASCAST]], align 8
 294 // CHECK-NEXT:    [[TMP2:%.*]] = call i1 @llvm.amdgcn.is.private(ptr [[TMP1]])
 295 // CHECK-NEXT:    [[FROMBOOL:%.*]] = zext i1 [[TMP2]] to i8
 296 // CHECK-NEXT:    store i8 [[FROMBOOL]], ptr [[RET_ASCAST]], align 1
 297 // CHECK-NEXT:    ret void
 298 //
 299 __global__ void test_is_private(int *x){
 300   bool ret = __builtin_amdgcn_is_private(x);
 301 }