320 lines · plain
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py2// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu gfx1200 -emit-llvm -o - %s | FileCheck %s3 4// REQUIRES: amdgpu-registered-target5 6typedef unsigned int uint;7 8// CHECK-LABEL: @test_s_sleep_var(9// CHECK-NEXT: entry:10// CHECK-NEXT: [[D_ADDR:%.*]] = alloca i32, align 4, addrspace(5)11// CHECK-NEXT: [[D_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[D_ADDR]] to ptr12// CHECK-NEXT: store i32 [[D:%.*]], ptr [[D_ADDR_ASCAST]], align 413// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[D_ADDR_ASCAST]], align 414// CHECK-NEXT: call void @llvm.amdgcn.s.sleep.var(i32 [[TMP0]])15// CHECK-NEXT: call void @llvm.amdgcn.s.sleep.var(i32 15)16// CHECK-NEXT: ret void17//18void test_s_sleep_var(int d)19{20 __builtin_amdgcn_s_sleep_var(d);21 __builtin_amdgcn_s_sleep_var(15);22}23 24// CHECK-LABEL: @test_permlane16_var(25// CHECK-NEXT: entry:26// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)27// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4, addrspace(5)28// CHECK-NEXT: [[B_ADDR:%.*]] = alloca i32, align 4, addrspace(5)29// CHECK-NEXT: [[C_ADDR:%.*]] = alloca i32, align 4, addrspace(5)30// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr31// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr32// CHECK-NEXT: [[B_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[B_ADDR]] to ptr33// CHECK-NEXT: [[C_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[C_ADDR]] to ptr34// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr [[OUT_ADDR_ASCAST]], align 835// CHECK-NEXT: store i32 [[A:%.*]], ptr [[A_ADDR_ASCAST]], align 436// CHECK-NEXT: store i32 [[B:%.*]], ptr [[B_ADDR_ASCAST]], align 437// CHECK-NEXT: store i32 [[C:%.*]], ptr [[C_ADDR_ASCAST]], align 438// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR_ASCAST]], align 439// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[B_ADDR_ASCAST]], align 440// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[C_ADDR_ASCAST]], align 441// CHECK-NEXT: [[TMP3:%.*]] = call i32 @llvm.amdgcn.permlane16.var(i32 [[TMP0]], i32 [[TMP1]], i32 [[TMP2]], i1 false, i1 false)42// CHECK-NEXT: [[TMP4:%.*]] = load ptr addrspace(1), ptr [[OUT_ADDR_ASCAST]], align 843// CHECK-NEXT: store i32 [[TMP3]], ptr addrspace(1) [[TMP4]], align 444// CHECK-NEXT: ret void45//46void test_permlane16_var(global uint* out, uint a, uint b, uint c) {47 *out = __builtin_amdgcn_permlane16_var(a, b, c, 0, 0);48}49 50// CHECK-LABEL: @test_permlanex16_var(51// CHECK-NEXT: entry:52// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)53// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4, addrspace(5)54// CHECK-NEXT: [[B_ADDR:%.*]] = alloca i32, align 4, addrspace(5)55// CHECK-NEXT: [[C_ADDR:%.*]] = alloca i32, align 4, addrspace(5)56// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr57// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr58// CHECK-NEXT: [[B_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[B_ADDR]] to ptr59// CHECK-NEXT: [[C_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[C_ADDR]] to ptr60// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr [[OUT_ADDR_ASCAST]], align 861// CHECK-NEXT: store i32 [[A:%.*]], ptr [[A_ADDR_ASCAST]], align 462// CHECK-NEXT: store i32 [[B:%.*]], ptr [[B_ADDR_ASCAST]], align 463// CHECK-NEXT: store i32 [[C:%.*]], ptr [[C_ADDR_ASCAST]], align 464// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR_ASCAST]], align 465// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[B_ADDR_ASCAST]], align 466// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[C_ADDR_ASCAST]], align 467// CHECK-NEXT: [[TMP3:%.*]] = call i32 @llvm.amdgcn.permlanex16.var(i32 [[TMP0]], i32 [[TMP1]], i32 [[TMP2]], i1 false, i1 false)68// CHECK-NEXT: [[TMP4:%.*]] = load ptr addrspace(1), ptr [[OUT_ADDR_ASCAST]], align 869// CHECK-NEXT: store i32 [[TMP3]], ptr addrspace(1) [[TMP4]], align 470// CHECK-NEXT: ret void71//72void test_permlanex16_var(global uint* out, uint a, uint b, uint c) {73 *out = __builtin_amdgcn_permlanex16_var(a, b, c, 0, 0);74}75 76// CHECK-LABEL: @test_s_barrier_signal(77// CHECK-NEXT: entry:78// CHECK-NEXT: call void @llvm.amdgcn.s.barrier.signal(i32 -1)79// CHECK-NEXT: call void @llvm.amdgcn.s.barrier.wait(i16 -1)80// CHECK-NEXT: ret void81//82void test_s_barrier_signal()83{84 __builtin_amdgcn_s_barrier_signal(-1);85 __builtin_amdgcn_s_barrier_wait(-1);86}87 88// CHECK-LABEL: @test_s_barrier_signal_var(89// CHECK-NEXT: entry:90// CHECK-NEXT: [[BAR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)91// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4, addrspace(5)92// CHECK-NEXT: [[BAR_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[BAR_ADDR]] to ptr93// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr94// CHECK-NEXT: store ptr [[BAR:%.*]], ptr [[BAR_ADDR_ASCAST]], align 895// CHECK-NEXT: store i32 [[A:%.*]], ptr [[A_ADDR_ASCAST]], align 496// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[BAR_ADDR_ASCAST]], align 897// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(3)98// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[A_ADDR_ASCAST]], align 499// CHECK-NEXT: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) [[TMP1]], i32 [[TMP2]])100// CHECK-NEXT: ret void101//102void test_s_barrier_signal_var(void *bar, int a)103{104 __builtin_amdgcn_s_barrier_signal_var(bar, a);105}106 107// CHECK-LABEL: @test_s_barrier_signal_isfirst(108// CHECK-NEXT: entry:109// CHECK-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)110// CHECK-NEXT: [[B_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)111// CHECK-NEXT: [[C_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)112// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr113// CHECK-NEXT: [[B_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[B_ADDR]] to ptr114// CHECK-NEXT: [[C_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[C_ADDR]] to ptr115// CHECK-NEXT: store ptr [[A:%.*]], ptr [[A_ADDR_ASCAST]], align 8116// CHECK-NEXT: store ptr [[B:%.*]], ptr [[B_ADDR_ASCAST]], align 8117// CHECK-NEXT: store ptr [[C:%.*]], ptr [[C_ADDR_ASCAST]], align 8118// CHECK-NEXT: [[TMP0:%.*]] = call i1 @llvm.amdgcn.s.barrier.signal.isfirst(i32 1)119// CHECK-NEXT: br i1 [[TMP0]], label [[IF_THEN:%.*]], label [[IF_ELSE:%.*]]120// CHECK: if.then:121// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[B_ADDR_ASCAST]], align 8122// CHECK-NEXT: store ptr [[TMP1]], ptr [[A_ADDR_ASCAST]], align 8123// CHECK-NEXT: br label [[IF_END:%.*]]124// CHECK: if.else:125// CHECK-NEXT: [[TMP2:%.*]] = load ptr, ptr [[C_ADDR_ASCAST]], align 8126// CHECK-NEXT: store ptr [[TMP2]], ptr [[A_ADDR_ASCAST]], align 8127// CHECK-NEXT: br label [[IF_END]]128// CHECK: if.end:129// CHECK-NEXT: call void @llvm.amdgcn.s.barrier.wait(i16 1)130// CHECK-NEXT: ret void131//132void test_s_barrier_signal_isfirst(int* a, int* b, int *c)133{134 if(__builtin_amdgcn_s_barrier_signal_isfirst(1))135 a = b;136 else137 a = c;138 139 __builtin_amdgcn_s_barrier_wait(1);140}141 142// CHECK-LABEL: @test_s_barrier_init(143// CHECK-NEXT: entry:144// CHECK-NEXT: [[BAR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)145// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4, addrspace(5)146// CHECK-NEXT: [[BAR_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[BAR_ADDR]] to ptr147// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr148// CHECK-NEXT: store ptr [[BAR:%.*]], ptr [[BAR_ADDR_ASCAST]], align 8149// CHECK-NEXT: store i32 [[A:%.*]], ptr [[A_ADDR_ASCAST]], align 4150// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[BAR_ADDR_ASCAST]], align 8151// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(3)152// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[A_ADDR_ASCAST]], align 4153// CHECK-NEXT: call void @llvm.amdgcn.s.barrier.init(ptr addrspace(3) [[TMP1]], i32 [[TMP2]])154// CHECK-NEXT: ret void155//156void test_s_barrier_init(void *bar, int a)157{158 __builtin_amdgcn_s_barrier_init(bar, a);159}160 161// CHECK-LABEL: @test_s_barrier_join(162// CHECK-NEXT: entry:163// CHECK-NEXT: [[BAR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)164// CHECK-NEXT: [[BAR_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[BAR_ADDR]] to ptr165// CHECK-NEXT: store ptr [[BAR:%.*]], ptr [[BAR_ADDR_ASCAST]], align 8166// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[BAR_ADDR_ASCAST]], align 8167// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(3)168// CHECK-NEXT: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) [[TMP1]])169// CHECK-NEXT: ret void170//171void test_s_barrier_join(void *bar)172{173 __builtin_amdgcn_s_barrier_join(bar);174}175 176// CHECK-LABEL: @test_s_barrier_leave(177// CHECK-NEXT: entry:178// CHECK-NEXT: call void @llvm.amdgcn.s.barrier.leave(i16 1)179// CHECK-NEXT: ret void180//181void test_s_barrier_leave()182{183 __builtin_amdgcn_s_barrier_leave(1);184}185 186// CHECK-LABEL: @test_s_get_barrier_state(187// CHECK-NEXT: entry:188// CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)189// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4, addrspace(5)190// CHECK-NEXT: [[STATE:%.*]] = alloca i32, align 4, addrspace(5)191// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr192// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr193// CHECK-NEXT: store i32 [[A:%.*]], ptr [[A_ADDR_ASCAST]], align 4194// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR_ASCAST]], align 4195// CHECK-NEXT: [[TMP1:%.*]] = call i32 @llvm.amdgcn.s.get.barrier.state(i32 [[TMP0]])196// CHECK-NEXT: store i32 [[TMP1]], ptr addrspace(5) [[STATE]], align 4197// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(5) [[STATE]], align 4198// CHECK-NEXT: ret i32 [[TMP2]]199//200unsigned test_s_get_barrier_state(int a)201{202 unsigned State = __builtin_amdgcn_s_get_barrier_state(a);203 return State;204}205 206// CHECK-LABEL: @test_s_get_named_barrier_state(207// CHECK-NEXT: entry:208// CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)209// CHECK-NEXT: [[BAR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)210// CHECK-NEXT: [[STATE:%.*]] = alloca i32, align 4, addrspace(5)211// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr212// CHECK-NEXT: [[BAR_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[BAR_ADDR]] to ptr213// CHECK-NEXT: store ptr [[BAR:%.*]], ptr [[BAR_ADDR_ASCAST]], align 8214// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[BAR_ADDR_ASCAST]], align 8215// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(3)216// CHECK-NEXT: [[TMP2:%.*]] = call i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(3) [[TMP1]])217// CHECK-NEXT: store i32 [[TMP2]], ptr addrspace(5) [[STATE]], align 4218// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr addrspace(5) [[STATE]], align 4219// CHECK-NEXT: ret i32 [[TMP3]]220//221unsigned test_s_get_named_barrier_state(void *bar)222{223 unsigned State = __builtin_amdgcn_s_get_named_barrier_state(bar);224 return State;225}226 227// CHECK-LABEL: @test_s_ttracedata(228// CHECK-NEXT: entry:229// CHECK-NEXT: call void @llvm.amdgcn.s.ttracedata(i32 1)230// CHECK-NEXT: ret void231//232void test_s_ttracedata()233{234 __builtin_amdgcn_s_ttracedata(1);235}236 237// CHECK-LABEL: @test_s_ttracedata_imm(238// CHECK-NEXT: entry:239// CHECK-NEXT: call void @llvm.amdgcn.s.ttracedata.imm(i16 1)240// CHECK-NEXT: ret void241//242void test_s_ttracedata_imm()243{244 __builtin_amdgcn_s_ttracedata_imm(1);245}246 247// CHECK-LABEL: @test_s_prefetch_data(248// CHECK-NEXT: entry:249// CHECK-NEXT: [[FP_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)250// CHECK-NEXT: [[GP_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)251// CHECK-NEXT: [[CP_ADDR:%.*]] = alloca ptr addrspace(4), align 8, addrspace(5)252// CHECK-NEXT: [[LEN_ADDR:%.*]] = alloca i32, align 4, addrspace(5)253// CHECK-NEXT: [[FP_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[FP_ADDR]] to ptr254// CHECK-NEXT: [[GP_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[GP_ADDR]] to ptr255// CHECK-NEXT: [[CP_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[CP_ADDR]] to ptr256// CHECK-NEXT: [[LEN_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[LEN_ADDR]] to ptr257// CHECK-NEXT: store ptr [[FP:%.*]], ptr [[FP_ADDR_ASCAST]], align 8258// CHECK-NEXT: store ptr addrspace(1) [[GP:%.*]], ptr [[GP_ADDR_ASCAST]], align 8259// CHECK-NEXT: store ptr addrspace(4) [[CP:%.*]], ptr [[CP_ADDR_ASCAST]], align 8260// CHECK-NEXT: store i32 [[LEN:%.*]], ptr [[LEN_ADDR_ASCAST]], align 4261// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[FP_ADDR_ASCAST]], align 8262// CHECK-NEXT: call void @llvm.amdgcn.s.prefetch.data.p0(ptr [[TMP0]], i32 0)263// CHECK-NEXT: [[TMP1:%.*]] = load ptr addrspace(1), ptr [[GP_ADDR_ASCAST]], align 8264// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[LEN_ADDR_ASCAST]], align 4265// CHECK-NEXT: call void @llvm.amdgcn.s.prefetch.data.p1(ptr addrspace(1) [[TMP1]], i32 [[TMP2]])266// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(4), ptr [[CP_ADDR_ASCAST]], align 8267// CHECK-NEXT: call void @llvm.amdgcn.s.prefetch.data.p4(ptr addrspace(4) [[TMP3]], i32 31)268// CHECK-NEXT: ret void269//270void test_s_prefetch_data(int *fp, global float *gp, constant char *cp, unsigned int len)271{272 __builtin_amdgcn_s_prefetch_data(fp, 0);273 __builtin_amdgcn_s_prefetch_data(gp, len);274 __builtin_amdgcn_s_prefetch_data(cp, 31);275}276 277// CHECK-LABEL: @test_s_buffer_prefetch_data(278// CHECK-NEXT: entry:279// CHECK-NEXT: [[RSRC_ADDR:%.*]] = alloca ptr addrspace(8), align 16, addrspace(5)280// CHECK-NEXT: [[LEN_ADDR:%.*]] = alloca i32, align 4, addrspace(5)281// CHECK-NEXT: [[RSRC_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RSRC_ADDR]] to ptr282// CHECK-NEXT: [[LEN_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[LEN_ADDR]] to ptr283// CHECK-NEXT: store ptr addrspace(8) [[RSRC:%.*]], ptr [[RSRC_ADDR_ASCAST]], align 16284// CHECK-NEXT: store i32 [[LEN:%.*]], ptr [[LEN_ADDR_ASCAST]], align 4285// CHECK-NEXT: [[TMP0:%.*]] = load ptr addrspace(8), ptr [[RSRC_ADDR_ASCAST]], align 16286// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[LEN_ADDR_ASCAST]], align 4287// CHECK-NEXT: call void @llvm.amdgcn.s.buffer.prefetch.data(ptr addrspace(8) [[TMP0]], i32 128, i32 [[TMP1]])288// CHECK-NEXT: [[TMP2:%.*]] = load ptr addrspace(8), ptr [[RSRC_ADDR_ASCAST]], align 16289// CHECK-NEXT: call void @llvm.amdgcn.s.buffer.prefetch.data(ptr addrspace(8) [[TMP2]], i32 0, i32 31)290// CHECK-NEXT: ret void291//292void test_s_buffer_prefetch_data(__amdgpu_buffer_rsrc_t rsrc, unsigned int len)293{294 __builtin_amdgcn_s_buffer_prefetch_data(rsrc, 128, len);295 __builtin_amdgcn_s_buffer_prefetch_data(rsrc, 0, 31);296}297 298// CHECK-LABEL: @test_ds_bpermute_fi_b32(299// CHECK-NEXT: entry:300// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)301// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4, addrspace(5)302// CHECK-NEXT: [[B_ADDR:%.*]] = alloca i32, align 4, addrspace(5)303// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr304// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr305// CHECK-NEXT: [[B_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[B_ADDR]] to ptr306// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr [[OUT_ADDR_ASCAST]], align 8307// CHECK-NEXT: store i32 [[A:%.*]], ptr [[A_ADDR_ASCAST]], align 4308// CHECK-NEXT: store i32 [[B:%.*]], ptr [[B_ADDR_ASCAST]], align 4309// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR_ASCAST]], align 4310// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[B_ADDR_ASCAST]], align 4311// CHECK-NEXT: [[TMP2:%.*]] = call i32 @llvm.amdgcn.ds.bpermute.fi.b32(i32 [[TMP0]], i32 [[TMP1]])312// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(1), ptr [[OUT_ADDR_ASCAST]], align 8313// CHECK-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[TMP3]], align 4314// CHECK-NEXT: ret void315//316void test_ds_bpermute_fi_b32(global int* out, int a, int b)317{318 *out = __builtin_amdgcn_ds_bpermute_fi_b32(a, b);319}320