293 lines · plain
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py2// 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 %s5 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 %s9 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 ptr18// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr19// CHECK-NEXT: [[DISPATCH_PTR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DISPATCH_PTR]] to ptr20// CHECK-NEXT: store ptr addrspace(1) [[OUT_COERCE:%.*]], ptr [[OUT_ASCAST]], align 821// CHECK-NEXT: [[OUT1:%.*]] = load ptr, ptr [[OUT_ASCAST]], align 822// CHECK-NEXT: store ptr [[OUT1]], ptr [[OUT_ADDR_ASCAST]], align 823// CHECK-NEXT: [[TMP0:%.*]] = call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()24// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr addrspace(4) [[TMP0]] to ptr25// CHECK-NEXT: store ptr [[TMP1]], ptr [[DISPATCH_PTR_ASCAST]], align 826// CHECK-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DISPATCH_PTR_ASCAST]], align 827// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 428// CHECK-NEXT: [[TMP4:%.*]] = load ptr, ptr [[OUT_ADDR_ASCAST]], align 829// CHECK-NEXT: store i32 [[TMP3]], ptr [[TMP4]], align 430// CHECK-NEXT: ret void31//32__global__ void use_dispatch_ptr(int* out) {33 const int* dispatch_ptr = (const int*)__builtin_amdgcn_dispatch_ptr();34 *out = *dispatch_ptr;35}36 37// CHECK-LABEL: @_Z13use_queue_ptrPi(38// CHECK-NEXT: entry:39// CHECK-NEXT: [[OUT:%.*]] = alloca ptr, align 8, addrspace(5)40// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)41// CHECK-NEXT: [[QUEUE_PTR:%.*]] = alloca ptr, align 8, addrspace(5)42// CHECK-NEXT: [[OUT_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT]] to ptr43// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr44// CHECK-NEXT: [[QUEUE_PTR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[QUEUE_PTR]] to ptr45// CHECK-NEXT: store ptr addrspace(1) [[OUT_COERCE:%.*]], ptr [[OUT_ASCAST]], align 846// CHECK-NEXT: [[OUT1:%.*]] = load ptr, ptr [[OUT_ASCAST]], align 847// CHECK-NEXT: store ptr [[OUT1]], ptr [[OUT_ADDR_ASCAST]], align 848// CHECK-NEXT: [[TMP0:%.*]] = call ptr addrspace(4) @llvm.amdgcn.queue.ptr()49// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr addrspace(4) [[TMP0]] to ptr50// CHECK-NEXT: store ptr [[TMP1]], ptr [[QUEUE_PTR_ASCAST]], align 851// CHECK-NEXT: [[TMP2:%.*]] = load ptr, ptr [[QUEUE_PTR_ASCAST]], align 852// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 453// CHECK-NEXT: [[TMP4:%.*]] = load ptr, ptr [[OUT_ADDR_ASCAST]], align 854// CHECK-NEXT: store i32 [[TMP3]], ptr [[TMP4]], align 455// CHECK-NEXT: ret void56//57__global__ void use_queue_ptr(int* out) {58 const int* queue_ptr = (const int*)__builtin_amdgcn_queue_ptr();59 *out = *queue_ptr;60}61 62// CHECK-LABEL: @_Z19use_implicitarg_ptrPi(63// CHECK-NEXT: entry:64// CHECK-NEXT: [[OUT:%.*]] = alloca ptr, align 8, addrspace(5)65// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)66// CHECK-NEXT: [[IMPLICITARG_PTR:%.*]] = alloca ptr, align 8, addrspace(5)67// CHECK-NEXT: [[OUT_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT]] to ptr68// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr69// CHECK-NEXT: [[IMPLICITARG_PTR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[IMPLICITARG_PTR]] to ptr70// CHECK-NEXT: store ptr addrspace(1) [[OUT_COERCE:%.*]], ptr [[OUT_ASCAST]], align 871// CHECK-NEXT: [[OUT1:%.*]] = load ptr, ptr [[OUT_ASCAST]], align 872// CHECK-NEXT: store ptr [[OUT1]], ptr [[OUT_ADDR_ASCAST]], align 873// CHECK-NEXT: [[TMP0:%.*]] = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()74// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr addrspace(4) [[TMP0]] to ptr75// CHECK-NEXT: store ptr [[TMP1]], ptr [[IMPLICITARG_PTR_ASCAST]], align 876// CHECK-NEXT: [[TMP2:%.*]] = load ptr, ptr [[IMPLICITARG_PTR_ASCAST]], align 877// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 478// CHECK-NEXT: [[TMP4:%.*]] = load ptr, ptr [[OUT_ADDR_ASCAST]], align 879// CHECK-NEXT: store i32 [[TMP3]], ptr [[TMP4]], align 480// CHECK-NEXT: ret void81//82__global__ void use_implicitarg_ptr(int* out) {83 const int* implicitarg_ptr = (const int*)__builtin_amdgcn_implicitarg_ptr();84 *out = *implicitarg_ptr;85}86 87__global__88 //89 void90// CHECK-LABEL: @_Z12test_ds_fmaxf(91// CHECK-NEXT: entry:92// CHECK-NEXT: [[SRC_ADDR:%.*]] = alloca float, align 4, addrspace(5)93// CHECK-NEXT: [[X:%.*]] = alloca float, align 4, addrspace(5)94// CHECK-NEXT: [[SRC_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SRC_ADDR]] to ptr95// CHECK-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X]] to ptr96// CHECK-NEXT: store float [[SRC:%.*]], ptr [[SRC_ADDR_ASCAST]], align 497// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr [[SRC_ADDR_ASCAST]], align 498// CHECK-NEXT: [[TMP1:%.*]] = atomicrmw fmax ptr addrspace(3) @_ZZ12test_ds_fmaxfE6shared, float [[TMP0]] monotonic, align 499// CHECK-NEXT: store volatile float [[TMP1]], ptr [[X_ASCAST]], align 4100// CHECK-NEXT: ret void101//102 test_ds_fmax(float src) {103 __shared__ float shared;104 volatile float x = __builtin_amdgcn_ds_fmaxf(&shared, src, 0, 0, false);105}106 107// CHECK-LABEL: @_Z12test_ds_faddf(108// CHECK-NEXT: entry:109// CHECK-NEXT: [[SRC_ADDR:%.*]] = alloca float, align 4, addrspace(5)110// CHECK-NEXT: [[X:%.*]] = alloca float, align 4, addrspace(5)111// CHECK-NEXT: [[SRC_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SRC_ADDR]] to ptr112// CHECK-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X]] to ptr113// CHECK-NEXT: store float [[SRC:%.*]], ptr [[SRC_ADDR_ASCAST]], align 4114// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr [[SRC_ADDR_ASCAST]], align 4115// CHECK-NEXT: [[TMP1:%.*]] = atomicrmw fadd ptr addrspace(3) @_ZZ12test_ds_faddfE6shared, float [[TMP0]] monotonic, align 4116// CHECK-NEXT: store volatile float [[TMP1]], ptr [[X_ASCAST]], align 4117// CHECK-NEXT: ret void118//119__global__ void test_ds_fadd(float src) {120 __shared__ float shared;121 volatile float x = __builtin_amdgcn_ds_faddf(&shared, src, 0, 0, false);122}123 124// CHECK-LABEL: @_Z12test_ds_fminfPf(125// CHECK-NEXT: entry:126// CHECK-NEXT: [[SHARED:%.*]] = alloca ptr, align 8, addrspace(5)127// CHECK-NEXT: [[SRC_ADDR:%.*]] = alloca float, align 4, addrspace(5)128// CHECK-NEXT: [[SHARED_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)129// CHECK-NEXT: [[X:%.*]] = alloca float, align 4, addrspace(5)130// CHECK-NEXT: [[SHARED_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SHARED]] to ptr131// CHECK-NEXT: [[SRC_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SRC_ADDR]] to ptr132// CHECK-NEXT: [[SHARED_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SHARED_ADDR]] to ptr133// CHECK-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X]] to ptr134// CHECK-NEXT: store ptr addrspace(1) [[SHARED_COERCE:%.*]], ptr [[SHARED_ASCAST]], align 8135// CHECK-NEXT: [[SHARED1:%.*]] = load ptr, ptr [[SHARED_ASCAST]], align 8136// CHECK-NEXT: store float [[SRC:%.*]], ptr [[SRC_ADDR_ASCAST]], align 4137// CHECK-NEXT: store ptr [[SHARED1]], ptr [[SHARED_ADDR_ASCAST]], align 8138// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[SHARED_ADDR_ASCAST]], align 8139// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(3)140// CHECK-NEXT: [[TMP2:%.*]] = load float, ptr [[SRC_ADDR_ASCAST]], align 4141// CHECK-NEXT: [[TMP3:%.*]] = atomicrmw fmin ptr addrspace(3) [[TMP1]], float [[TMP2]] monotonic, align 4142// CHECK-NEXT: store volatile float [[TMP3]], ptr [[X_ASCAST]], align 4143// CHECK-NEXT: ret void144//145__global__ void test_ds_fmin(float src, float *shared) {146 volatile float x = __builtin_amdgcn_ds_fminf(shared, src, 0, 0, false);147}148 149// CHECK-LABEL: @_Z33test_ret_builtin_nondef_addrspacev(150// CHECK-NEXT: entry:151// CHECK-NEXT: [[X:%.*]] = alloca ptr, align 8, addrspace(5)152// CHECK-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X]] to ptr153// CHECK-NEXT: [[TMP0:%.*]] = call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()154// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr addrspace(4) [[TMP0]] to ptr155// CHECK-NEXT: store ptr [[TMP1]], ptr [[X_ASCAST]], align 8156// CHECK-NEXT: ret void157//158__device__ void test_ret_builtin_nondef_addrspace() {159 void *x = __builtin_amdgcn_dispatch_ptr();160}161 162// CHECK-LABEL: @_Z6endpgmv(163// CHECK-NEXT: entry:164// CHECK-NEXT: call void @llvm.amdgcn.endpgm()165// CHECK-NEXT: ret void166//167__global__ void endpgm() {168 __builtin_amdgcn_endpgm();169}170 171// Check the 64 bit argument is correctly passed to the intrinsic without truncation or assertion.172 173// CHECK-LABEL: @_Z14test_uicmp_i64Pyyy(174// CHECK-NEXT: entry:175// CHECK-NEXT: [[OUT:%.*]] = alloca ptr, align 8, addrspace(5)176// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)177// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8, addrspace(5)178// CHECK-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8, addrspace(5)179// CHECK-NEXT: [[OUT_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT]] to ptr180// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr181// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr182// CHECK-NEXT: [[B_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[B_ADDR]] to ptr183// CHECK-NEXT: store ptr addrspace(1) [[OUT_COERCE:%.*]], ptr [[OUT_ASCAST]], align 8184// CHECK-NEXT: [[OUT1:%.*]] = load ptr, ptr [[OUT_ASCAST]], align 8185// CHECK-NEXT: store ptr [[OUT1]], ptr [[OUT_ADDR_ASCAST]], align 8186// CHECK-NEXT: store i64 [[A:%.*]], ptr [[A_ADDR_ASCAST]], align 8187// CHECK-NEXT: store i64 [[B:%.*]], ptr [[B_ADDR_ASCAST]], align 8188// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[A_ADDR_ASCAST]], align 8189// CHECK-NEXT: [[TMP1:%.*]] = load i64, ptr [[B_ADDR_ASCAST]], align 8190// CHECK-NEXT: [[TMP2:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i64(i64 [[TMP0]], i64 [[TMP1]], i32 35)191// CHECK-NEXT: [[TMP3:%.*]] = load ptr, ptr [[OUT_ADDR_ASCAST]], align 8192// CHECK-NEXT: store i64 [[TMP2]], ptr [[TMP3]], align 8193// CHECK-NEXT: ret void194//195__global__ void test_uicmp_i64(unsigned long long *out, unsigned long long a, unsigned long long b)196{197 *out = __builtin_amdgcn_uicmpl(a, b, 30+5);198}199 200// Check the 64 bit return value is correctly returned without truncation or assertion.201 202// CHECK-LABEL: @_Z14test_s_memtimePy(203// CHECK-NEXT: entry:204// CHECK-NEXT: [[OUT:%.*]] = alloca ptr, align 8, addrspace(5)205// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)206// CHECK-NEXT: [[OUT_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT]] to ptr207// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr208// CHECK-NEXT: store ptr addrspace(1) [[OUT_COERCE:%.*]], ptr [[OUT_ASCAST]], align 8209// CHECK-NEXT: [[OUT1:%.*]] = load ptr, ptr [[OUT_ASCAST]], align 8210// CHECK-NEXT: store ptr [[OUT1]], ptr [[OUT_ADDR_ASCAST]], align 8211// CHECK-NEXT: [[TMP0:%.*]] = call i64 @llvm.amdgcn.s.memtime()212// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[OUT_ADDR_ASCAST]], align 8213// CHECK-NEXT: store i64 [[TMP0]], ptr [[TMP1]], align 8214// CHECK-NEXT: ret void215//216__global__ void test_s_memtime(unsigned long long* out)217{218 *out = __builtin_amdgcn_s_memtime();219}220 221// Check a generic pointer can be passed as a shared pointer and a generic pointer.222__device__ void func(float *x);223 224// CHECK-LABEL: @_Z17test_ds_fmin_funcfPf(225// CHECK-NEXT: entry:226// CHECK-NEXT: [[SHARED:%.*]] = alloca ptr, align 8, addrspace(5)227// CHECK-NEXT: [[SRC_ADDR:%.*]] = alloca float, align 4, addrspace(5)228// CHECK-NEXT: [[SHARED_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)229// CHECK-NEXT: [[X:%.*]] = alloca float, align 4, addrspace(5)230// CHECK-NEXT: [[SHARED_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SHARED]] to ptr231// CHECK-NEXT: [[SRC_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SRC_ADDR]] to ptr232// CHECK-NEXT: [[SHARED_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[SHARED_ADDR]] to ptr233// CHECK-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X]] to ptr234// CHECK-NEXT: store ptr addrspace(1) [[SHARED_COERCE:%.*]], ptr [[SHARED_ASCAST]], align 8235// CHECK-NEXT: [[SHARED1:%.*]] = load ptr, ptr [[SHARED_ASCAST]], align 8236// CHECK-NEXT: store float [[SRC:%.*]], ptr [[SRC_ADDR_ASCAST]], align 4237// CHECK-NEXT: store ptr [[SHARED1]], ptr [[SHARED_ADDR_ASCAST]], align 8238// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[SHARED_ADDR_ASCAST]], align 8239// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(3)240// CHECK-NEXT: [[TMP2:%.*]] = load float, ptr [[SRC_ADDR_ASCAST]], align 4241// CHECK-NEXT: [[TMP3:%.*]] = atomicrmw fmin ptr addrspace(3) [[TMP1]], float [[TMP2]] monotonic, align 4242// CHECK-NEXT: store volatile float [[TMP3]], ptr [[X_ASCAST]], align 4243// CHECK-NEXT: [[TMP4:%.*]] = load ptr, ptr [[SHARED_ADDR_ASCAST]], align 8244// CHECK-NEXT: call void @_Z4funcPf(ptr noundef [[TMP4]]) #[[ATTR7:[0-9]+]]245// CHECK-NEXT: ret void246//247__global__ void test_ds_fmin_func(float src, float *__restrict shared) {248 volatile float x = __builtin_amdgcn_ds_fminf(shared, src, 0, 0, false);249 func(shared);250}251 252// CHECK-LABEL: @_Z14test_is_sharedPf(253// CHECK-NEXT: entry:254// CHECK-NEXT: [[X:%.*]] = alloca ptr, align 8, addrspace(5)255// CHECK-NEXT: [[X_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)256// CHECK-NEXT: [[RET:%.*]] = alloca i8, align 1, addrspace(5)257// CHECK-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X]] to ptr258// CHECK-NEXT: [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X_ADDR]] to ptr259// CHECK-NEXT: [[RET_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RET]] to ptr260// CHECK-NEXT: store ptr addrspace(1) [[X_COERCE:%.*]], ptr [[X_ASCAST]], align 8261// CHECK-NEXT: [[X1:%.*]] = load ptr, ptr [[X_ASCAST]], align 8262// CHECK-NEXT: store ptr [[X1]], ptr [[X_ADDR_ASCAST]], align 8263// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[X_ADDR_ASCAST]], align 8264// CHECK-NEXT: [[TMP1:%.*]] = call i1 @llvm.amdgcn.is.shared(ptr [[TMP0]])265// CHECK-NEXT: [[STOREDV:%.*]] = zext i1 [[TMP1]] to i8266// CHECK-NEXT: store i8 [[STOREDV]], ptr [[RET_ASCAST]], align 1267// CHECK-NEXT: ret void268//269__global__ void test_is_shared(float *x){270 bool ret = __builtin_amdgcn_is_shared(x);271}272 273// CHECK-LABEL: @_Z15test_is_privatePi(274// CHECK-NEXT: entry:275// CHECK-NEXT: [[X:%.*]] = alloca ptr, align 8, addrspace(5)276// CHECK-NEXT: [[X_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)277// CHECK-NEXT: [[RET:%.*]] = alloca i8, align 1, addrspace(5)278// CHECK-NEXT: [[X_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X]] to ptr279// CHECK-NEXT: [[X_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X_ADDR]] to ptr280// CHECK-NEXT: [[RET_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RET]] to ptr281// CHECK-NEXT: store ptr addrspace(1) [[X_COERCE:%.*]], ptr [[X_ASCAST]], align 8282// CHECK-NEXT: [[X1:%.*]] = load ptr, ptr [[X_ASCAST]], align 8283// CHECK-NEXT: store ptr [[X1]], ptr [[X_ADDR_ASCAST]], align 8284// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[X_ADDR_ASCAST]], align 8285// CHECK-NEXT: [[TMP1:%.*]] = call i1 @llvm.amdgcn.is.private(ptr [[TMP0]])286// CHECK-NEXT: [[STOREDV:%.*]] = zext i1 [[TMP1]] to i8287// CHECK-NEXT: store i8 [[STOREDV]], ptr [[RET_ASCAST]], align 1288// CHECK-NEXT: ret void289//290__global__ void test_is_private(int *x){291 bool ret = __builtin_amdgcn_is_private(x);292}293