286 lines · plain
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py2// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip \3// RUN: -aux-triple x86_64-unknown-linux-gnu -fcuda-is-device -emit-llvm %s \4// RUN: -o - | FileCheck %s5 6// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip \7// RUN: -aux-triple x86_64-pc-windows-msvc -fcuda-is-device -emit-llvm %s \8// RUN: -o - | FileCheck %s9 10#include "Inputs/cuda.h"11 12// CHECK-LABEL: @_Z16use_dispatch_ptrPi(13// CHECK-NEXT: entry:14// CHECK-NEXT: [[OUT:%.*]] = alloca ptr addrspace(4), align 815// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(4), align 816// CHECK-NEXT: [[DISPATCH_PTR:%.*]] = alloca ptr addrspace(4), align 817// CHECK-NEXT: [[OUT_ASCAST:%.*]] = addrspacecast ptr [[OUT]] to ptr addrspace(4)18// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr [[OUT_ADDR]] to ptr addrspace(4)19// CHECK-NEXT: [[DISPATCH_PTR_ASCAST:%.*]] = addrspacecast ptr [[DISPATCH_PTR]] to ptr addrspace(4)20// CHECK-NEXT: store ptr addrspace(1) [[OUT_COERCE:%.*]], ptr addrspace(4) [[OUT_ASCAST]], align 821// CHECK-NEXT: [[OUT1:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[OUT_ASCAST]], align 822// CHECK-NEXT: store ptr addrspace(4) [[OUT1]], ptr addrspace(4) [[OUT_ADDR_ASCAST]], align 823// CHECK-NEXT: [[TMP0:%.*]] = call align 4 dereferenceable(64) addrspace(4) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()24// CHECK-NEXT: store ptr addrspace(4) [[TMP0]], ptr addrspace(4) [[DISPATCH_PTR_ASCAST]], align 825// CHECK-NEXT: [[TMP1:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[DISPATCH_PTR_ASCAST]], align 826// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[TMP1]], align 427// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[OUT_ADDR_ASCAST]], align 828// CHECK-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[TMP3]], align 429// CHECK-NEXT: ret void30//31__global__ void use_dispatch_ptr(int* out) {32 const int* dispatch_ptr = (const int*)__builtin_amdgcn_dispatch_ptr();33 *out = *dispatch_ptr;34}35 36// CHECK-LABEL: @_Z13use_queue_ptrPi(37// CHECK-NEXT: entry:38// CHECK-NEXT: [[OUT:%.*]] = alloca ptr addrspace(4), align 839// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(4), align 840// CHECK-NEXT: [[QUEUE_PTR:%.*]] = alloca ptr addrspace(4), align 841// CHECK-NEXT: [[OUT_ASCAST:%.*]] = addrspacecast ptr [[OUT]] to ptr addrspace(4)42// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr [[OUT_ADDR]] to ptr addrspace(4)43// CHECK-NEXT: [[QUEUE_PTR_ASCAST:%.*]] = addrspacecast ptr [[QUEUE_PTR]] to ptr addrspace(4)44// CHECK-NEXT: store ptr addrspace(1) [[OUT_COERCE:%.*]], ptr addrspace(4) [[OUT_ASCAST]], align 845// CHECK-NEXT: [[OUT1:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[OUT_ASCAST]], align 846// CHECK-NEXT: store ptr addrspace(4) [[OUT1]], ptr addrspace(4) [[OUT_ADDR_ASCAST]], align 847// CHECK-NEXT: [[TMP0:%.*]] = call addrspace(4) ptr addrspace(4) @llvm.amdgcn.queue.ptr()48// CHECK-NEXT: store ptr addrspace(4) [[TMP0]], ptr addrspace(4) [[QUEUE_PTR_ASCAST]], align 849// CHECK-NEXT: [[TMP1:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[QUEUE_PTR_ASCAST]], align 850// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[TMP1]], align 451// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[OUT_ADDR_ASCAST]], align 852// CHECK-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[TMP3]], align 453// CHECK-NEXT: ret void54//55__global__ void use_queue_ptr(int* out) {56 const int* queue_ptr = (const int*)__builtin_amdgcn_queue_ptr();57 *out = *queue_ptr;58}59 60// CHECK-LABEL: @_Z19use_implicitarg_ptrPi(61// CHECK-NEXT: entry:62// CHECK-NEXT: [[OUT:%.*]] = alloca ptr addrspace(4), align 863// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(4), align 864// CHECK-NEXT: [[IMPLICITARG_PTR:%.*]] = alloca ptr addrspace(4), align 865// CHECK-NEXT: [[OUT_ASCAST:%.*]] = addrspacecast ptr [[OUT]] to ptr addrspace(4)66// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr [[OUT_ADDR]] to ptr addrspace(4)67// CHECK-NEXT: [[IMPLICITARG_PTR_ASCAST:%.*]] = addrspacecast ptr [[IMPLICITARG_PTR]] to ptr addrspace(4)68// CHECK-NEXT: store ptr addrspace(1) [[OUT_COERCE:%.*]], ptr addrspace(4) [[OUT_ASCAST]], align 869// CHECK-NEXT: [[OUT1:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[OUT_ASCAST]], align 870// CHECK-NEXT: store ptr addrspace(4) [[OUT1]], ptr addrspace(4) [[OUT_ADDR_ASCAST]], align 871// CHECK-NEXT: [[TMP0:%.*]] = call addrspace(4) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()72// CHECK-NEXT: store ptr addrspace(4) [[TMP0]], ptr addrspace(4) [[IMPLICITARG_PTR_ASCAST]], align 873// CHECK-NEXT: [[TMP1:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[IMPLICITARG_PTR_ASCAST]], align 874// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[TMP1]], align 475// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[OUT_ADDR_ASCAST]], align 876// CHECK-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[TMP3]], align 477// CHECK-NEXT: ret void78//79__global__ void use_implicitarg_ptr(int* out) {80 const int* implicitarg_ptr = (const int*)__builtin_amdgcn_implicitarg_ptr();81 *out = *implicitarg_ptr;82}83 84__global__85 //86 void87// CHECK-LABEL: @_Z12test_ds_fmaxf(88// CHECK-NEXT: entry:89// CHECK-NEXT: [[SRC_ADDR:%.*]] = alloca float, align 490// CHECK-NEXT: [[X:%.*]] = alloca float, align 491// CHECK-NEXT: [[SRC_ADDR_ASCAST:%.*]] = addrspacecast ptr [[SRC_ADDR]] to ptr addrspace(4)92// CHECK-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr [[X]] to ptr addrspace(4)93// CHECK-NEXT: store float [[SRC:%.*]], ptr addrspace(4) [[SRC_ADDR_ASCAST]], align 494// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(4) [[SRC_ADDR_ASCAST]], align 495// CHECK-NEXT: [[TMP1:%.*]] = atomicrmw fmax ptr addrspace(3) @_ZZ12test_ds_fmaxfE6shared, float [[TMP0]] monotonic, align 496// CHECK-NEXT: store volatile float [[TMP1]], ptr addrspace(4) [[X_ASCAST]], align 497// CHECK-NEXT: ret void98//99 test_ds_fmax(float src) {100 __shared__ float shared;101 volatile float x = __builtin_amdgcn_ds_fmaxf(&shared, src, 0, 0, false);102}103 104// CHECK-LABEL: @_Z12test_ds_faddf(105// CHECK-NEXT: entry:106// CHECK-NEXT: [[SRC_ADDR:%.*]] = alloca float, align 4107// CHECK-NEXT: [[X:%.*]] = alloca float, align 4108// CHECK-NEXT: [[SRC_ADDR_ASCAST:%.*]] = addrspacecast ptr [[SRC_ADDR]] to ptr addrspace(4)109// CHECK-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr [[X]] to ptr addrspace(4)110// CHECK-NEXT: store float [[SRC:%.*]], ptr addrspace(4) [[SRC_ADDR_ASCAST]], align 4111// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr addrspace(4) [[SRC_ADDR_ASCAST]], align 4112// CHECK-NEXT: [[TMP1:%.*]] = atomicrmw fadd ptr addrspace(3) @_ZZ12test_ds_faddfE6shared, float [[TMP0]] monotonic, align 4113// CHECK-NEXT: store volatile float [[TMP1]], ptr addrspace(4) [[X_ASCAST]], align 4114// CHECK-NEXT: ret void115//116__global__ void test_ds_fadd(float src) {117 __shared__ float shared;118 volatile float x = __builtin_amdgcn_ds_faddf(&shared, src, 0, 0, false);119}120 121// CHECK-LABEL: @_Z12test_ds_fminfPf(122// CHECK-NEXT: entry:123// CHECK-NEXT: [[SHARED:%.*]] = alloca ptr addrspace(4), align 8124// CHECK-NEXT: [[SRC_ADDR:%.*]] = alloca float, align 4125// CHECK-NEXT: [[SHARED_ADDR:%.*]] = alloca ptr addrspace(4), align 8126// CHECK-NEXT: [[X:%.*]] = alloca float, align 4127// CHECK-NEXT: [[SHARED_ASCAST:%.*]] = addrspacecast ptr [[SHARED]] to ptr addrspace(4)128// CHECK-NEXT: [[SRC_ADDR_ASCAST:%.*]] = addrspacecast ptr [[SRC_ADDR]] to ptr addrspace(4)129// CHECK-NEXT: [[SHARED_ADDR_ASCAST:%.*]] = addrspacecast ptr [[SHARED_ADDR]] to ptr addrspace(4)130// CHECK-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr [[X]] to ptr addrspace(4)131// CHECK-NEXT: store ptr addrspace(1) [[SHARED_COERCE:%.*]], ptr addrspace(4) [[SHARED_ASCAST]], align 8132// CHECK-NEXT: [[SHARED1:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[SHARED_ASCAST]], align 8133// CHECK-NEXT: store float [[SRC:%.*]], ptr addrspace(4) [[SRC_ADDR_ASCAST]], align 4134// CHECK-NEXT: store ptr addrspace(4) [[SHARED1]], ptr addrspace(4) [[SHARED_ADDR_ASCAST]], align 8135// CHECK-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[SHARED_ADDR_ASCAST]], align 8136// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr addrspace(4) [[TMP0]] to ptr addrspace(3)137// CHECK-NEXT: [[TMP2:%.*]] = load float, ptr addrspace(4) [[SRC_ADDR_ASCAST]], align 4138// CHECK-NEXT: [[TMP3:%.*]] = atomicrmw fmin ptr addrspace(3) [[TMP1]], float [[TMP2]] monotonic, align 4139// CHECK-NEXT: store volatile float [[TMP3]], ptr addrspace(4) [[X_ASCAST]], align 4140// CHECK-NEXT: ret void141//142__global__ void test_ds_fmin(float src, float *shared) {143 volatile float x = __builtin_amdgcn_ds_fminf(shared, src, 0, 0, false);144}145 146#if 0 // FIXME: returning a pointer to AS4 explicitly is wrong for AMDGPU SPIRV147//148__device__ void test_ret_builtin_nondef_addrspace() {149 void *x = __builtin_amdgcn_dispatch_ptr();150}151#endif152 153// CHECK-LABEL: @_Z6endpgmv(154// CHECK-NEXT: entry:155// CHECK-NEXT: call addrspace(4) void @llvm.amdgcn.endpgm()156// CHECK-NEXT: ret void157//158__global__ void endpgm() {159 __builtin_amdgcn_endpgm();160}161 162// Check the 64 bit argument is correctly passed to the intrinsic without truncation or assertion.163 164// CHECK-LABEL: @_Z14test_uicmp_i64Pyyy(165// CHECK-NEXT: entry:166// CHECK-NEXT: [[OUT:%.*]] = alloca ptr addrspace(4), align 8167// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(4), align 8168// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8169// CHECK-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8170// CHECK-NEXT: [[OUT_ASCAST:%.*]] = addrspacecast ptr [[OUT]] to ptr addrspace(4)171// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr [[OUT_ADDR]] to ptr addrspace(4)172// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr [[A_ADDR]] to ptr addrspace(4)173// CHECK-NEXT: [[B_ADDR_ASCAST:%.*]] = addrspacecast ptr [[B_ADDR]] to ptr addrspace(4)174// CHECK-NEXT: store ptr addrspace(1) [[OUT_COERCE:%.*]], ptr addrspace(4) [[OUT_ASCAST]], align 8175// CHECK-NEXT: [[OUT1:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[OUT_ASCAST]], align 8176// CHECK-NEXT: store ptr addrspace(4) [[OUT1]], ptr addrspace(4) [[OUT_ADDR_ASCAST]], align 8177// CHECK-NEXT: store i64 [[A:%.*]], ptr addrspace(4) [[A_ADDR_ASCAST]], align 8178// CHECK-NEXT: store i64 [[B:%.*]], ptr addrspace(4) [[B_ADDR_ASCAST]], align 8179// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr addrspace(4) [[A_ADDR_ASCAST]], align 8180// CHECK-NEXT: [[TMP1:%.*]] = load i64, ptr addrspace(4) [[B_ADDR_ASCAST]], align 8181// CHECK-NEXT: [[TMP2:%.*]] = call addrspace(4) i64 @llvm.amdgcn.icmp.i64.i64(i64 [[TMP0]], i64 [[TMP1]], i32 35)182// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[OUT_ADDR_ASCAST]], align 8183// CHECK-NEXT: store i64 [[TMP2]], ptr addrspace(4) [[TMP3]], align 8184// CHECK-NEXT: ret void185//186__global__ void test_uicmp_i64(unsigned long long *out, unsigned long long a, unsigned long long b)187{188 *out = __builtin_amdgcn_uicmpl(a, b, 30+5);189}190 191// Check the 64 bit return value is correctly returned without truncation or assertion.192 193// CHECK-LABEL: @_Z14test_s_memtimePy(194// CHECK-NEXT: entry:195// CHECK-NEXT: [[OUT:%.*]] = alloca ptr addrspace(4), align 8196// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(4), align 8197// CHECK-NEXT: [[OUT_ASCAST:%.*]] = addrspacecast ptr [[OUT]] to ptr addrspace(4)198// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr [[OUT_ADDR]] to ptr addrspace(4)199// CHECK-NEXT: store ptr addrspace(1) [[OUT_COERCE:%.*]], ptr addrspace(4) [[OUT_ASCAST]], align 8200// CHECK-NEXT: [[OUT1:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[OUT_ASCAST]], align 8201// CHECK-NEXT: store ptr addrspace(4) [[OUT1]], ptr addrspace(4) [[OUT_ADDR_ASCAST]], align 8202// CHECK-NEXT: [[TMP0:%.*]] = call addrspace(4) i64 @llvm.amdgcn.s.memtime()203// CHECK-NEXT: [[TMP1:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[OUT_ADDR_ASCAST]], align 8204// CHECK-NEXT: store i64 [[TMP0]], ptr addrspace(4) [[TMP1]], align 8205// CHECK-NEXT: ret void206//207__global__ void test_s_memtime(unsigned long long* out)208{209 *out = __builtin_amdgcn_s_memtime();210}211 212// Check a generic pointer can be passed as a shared pointer and a generic pointer.213__device__ void func(float *x);214 215// CHECK-LABEL: @_Z17test_ds_fmin_funcfPf(216// CHECK-NEXT: entry:217// CHECK-NEXT: [[SHARED:%.*]] = alloca ptr addrspace(4), align 8218// CHECK-NEXT: [[SRC_ADDR:%.*]] = alloca float, align 4219// CHECK-NEXT: [[SHARED_ADDR:%.*]] = alloca ptr addrspace(4), align 8220// CHECK-NEXT: [[X:%.*]] = alloca float, align 4221// CHECK-NEXT: [[SHARED_ASCAST:%.*]] = addrspacecast ptr [[SHARED]] to ptr addrspace(4)222// CHECK-NEXT: [[SRC_ADDR_ASCAST:%.*]] = addrspacecast ptr [[SRC_ADDR]] to ptr addrspace(4)223// CHECK-NEXT: [[SHARED_ADDR_ASCAST:%.*]] = addrspacecast ptr [[SHARED_ADDR]] to ptr addrspace(4)224// CHECK-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr [[X]] to ptr addrspace(4)225// CHECK-NEXT: store ptr addrspace(1) [[SHARED_COERCE:%.*]], ptr addrspace(4) [[SHARED_ASCAST]], align 8226// CHECK-NEXT: [[SHARED1:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[SHARED_ASCAST]], align 8227// CHECK-NEXT: store float [[SRC:%.*]], ptr addrspace(4) [[SRC_ADDR_ASCAST]], align 4228// CHECK-NEXT: store ptr addrspace(4) [[SHARED1]], ptr addrspace(4) [[SHARED_ADDR_ASCAST]], align 8229// CHECK-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[SHARED_ADDR_ASCAST]], align 8230// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr addrspace(4) [[TMP0]] to ptr addrspace(3)231// CHECK-NEXT: [[TMP2:%.*]] = load float, ptr addrspace(4) [[SRC_ADDR_ASCAST]], align 4232// CHECK-NEXT: [[TMP3:%.*]] = atomicrmw fmin ptr addrspace(3) [[TMP1]], float [[TMP2]] monotonic, align 4233// CHECK-NEXT: store volatile float [[TMP3]], ptr addrspace(4) [[X_ASCAST]], align 4234// CHECK-NEXT: [[TMP4:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[SHARED_ADDR_ASCAST]], align 8235// CHECK-NEXT: call spir_func addrspace(4) void @_Z4funcPf(ptr addrspace(4) noundef [[TMP4]]) #[[ATTR6:[0-9]+]]236// CHECK-NEXT: ret void237//238__global__ void test_ds_fmin_func(float src, float *__restrict shared) {239 volatile float x = __builtin_amdgcn_ds_fminf(shared, src, 0, 0, false);240 func(shared);241}242 243// CHECK-LABEL: @_Z14test_is_sharedPf(244// CHECK-NEXT: entry:245// CHECK-NEXT: [[X:%.*]] = alloca ptr addrspace(4), align 8246// CHECK-NEXT: [[X_ADDR:%.*]] = alloca ptr addrspace(4), align 8247// CHECK-NEXT: [[RET:%.*]] = alloca i8, align 1248// CHECK-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr [[X]] to ptr addrspace(4)249// CHECK-NEXT: [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr [[X_ADDR]] to ptr addrspace(4)250// CHECK-NEXT: [[RET_ASCAST:%.*]] = addrspacecast ptr [[RET]] to ptr addrspace(4)251// CHECK-NEXT: store ptr addrspace(1) [[X_COERCE:%.*]], ptr addrspace(4) [[X_ASCAST]], align 8252// CHECK-NEXT: [[X1:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[X_ASCAST]], align 8253// CHECK-NEXT: store ptr addrspace(4) [[X1]], ptr addrspace(4) [[X_ADDR_ASCAST]], align 8254// CHECK-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[X_ADDR_ASCAST]], align 8255// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr addrspace(4) [[TMP0]] to ptr256// CHECK-NEXT: [[TMP2:%.*]] = call addrspace(4) i1 @llvm.amdgcn.is.shared(ptr [[TMP1]])257// CHECK-NEXT: [[STOREDV:%.*]] = zext i1 [[TMP2]] to i8258// CHECK-NEXT: store i8 [[STOREDV]], ptr addrspace(4) [[RET_ASCAST]], align 1259// CHECK-NEXT: ret void260//261__global__ void test_is_shared(float *x){262 bool ret = __builtin_amdgcn_is_shared(x);263}264 265// CHECK-LABEL: @_Z15test_is_privatePi(266// CHECK-NEXT: entry:267// CHECK-NEXT: [[X:%.*]] = alloca ptr addrspace(4), align 8268// CHECK-NEXT: [[X_ADDR:%.*]] = alloca ptr addrspace(4), align 8269// CHECK-NEXT: [[RET:%.*]] = alloca i8, align 1270// CHECK-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr [[X]] to ptr addrspace(4)271// CHECK-NEXT: [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr [[X_ADDR]] to ptr addrspace(4)272// CHECK-NEXT: [[RET_ASCAST:%.*]] = addrspacecast ptr [[RET]] to ptr addrspace(4)273// CHECK-NEXT: store ptr addrspace(1) [[X_COERCE:%.*]], ptr addrspace(4) [[X_ASCAST]], align 8274// CHECK-NEXT: [[X1:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[X_ASCAST]], align 8275// CHECK-NEXT: store ptr addrspace(4) [[X1]], ptr addrspace(4) [[X_ADDR_ASCAST]], align 8276// CHECK-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[X_ADDR_ASCAST]], align 8277// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr addrspace(4) [[TMP0]] to ptr278// CHECK-NEXT: [[TMP2:%.*]] = call addrspace(4) i1 @llvm.amdgcn.is.private(ptr [[TMP1]])279// CHECK-NEXT: [[STOREDV:%.*]] = zext i1 [[TMP2]] to i8280// CHECK-NEXT: store i8 [[STOREDV]], ptr addrspace(4) [[RET_ASCAST]], align 1281// CHECK-NEXT: ret void282//283__global__ void test_is_private(int *x){284 bool ret = __builtin_amdgcn_is_private(x);285}286