278 lines
14 KiB
Common Lisp
278 lines
14 KiB
Common Lisp
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
|
|
// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu gfx1200 -emit-llvm -o - %s | FileCheck %s
|
|
|
|
// REQUIRES: amdgpu-registered-target
|
|
|
|
typedef unsigned int uint;
|
|
|
|
// CHECK-LABEL: @test_s_sleep_var(
|
|
// CHECK-NEXT: entry:
|
|
// CHECK-NEXT: [[D_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
|
|
// CHECK-NEXT: [[D_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[D_ADDR]] to ptr
|
|
// CHECK-NEXT: store i32 [[D:%.*]], ptr [[D_ADDR_ASCAST]], align 4
|
|
// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[D_ADDR_ASCAST]], align 4
|
|
// CHECK-NEXT: call void @llvm.amdgcn.s.sleep.var(i32 [[TMP0]])
|
|
// CHECK-NEXT: call void @llvm.amdgcn.s.sleep.var(i32 15)
|
|
// CHECK-NEXT: ret void
|
|
//
|
|
void test_s_sleep_var(int d)
|
|
{
|
|
__builtin_amdgcn_s_sleep_var(d);
|
|
__builtin_amdgcn_s_sleep_var(15);
|
|
}
|
|
|
|
// CHECK-LABEL: @test_permlane16_var(
|
|
// CHECK-NEXT: entry:
|
|
// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
|
|
// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
|
|
// CHECK-NEXT: [[B_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
|
|
// CHECK-NEXT: [[C_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
|
|
// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr
|
|
// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
|
|
// CHECK-NEXT: [[B_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[B_ADDR]] to ptr
|
|
// CHECK-NEXT: [[C_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[C_ADDR]] to ptr
|
|
// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr [[OUT_ADDR_ASCAST]], align 8
|
|
// CHECK-NEXT: store i32 [[A:%.*]], ptr [[A_ADDR_ASCAST]], align 4
|
|
// CHECK-NEXT: store i32 [[B:%.*]], ptr [[B_ADDR_ASCAST]], align 4
|
|
// CHECK-NEXT: store i32 [[C:%.*]], ptr [[C_ADDR_ASCAST]], align 4
|
|
// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR_ASCAST]], align 4
|
|
// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[B_ADDR_ASCAST]], align 4
|
|
// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[C_ADDR_ASCAST]], align 4
|
|
// CHECK-NEXT: [[TMP3:%.*]] = call i32 @llvm.amdgcn.permlane16.var(i32 [[TMP0]], i32 [[TMP1]], i32 [[TMP2]], i1 false, i1 false)
|
|
// CHECK-NEXT: [[TMP4:%.*]] = load ptr addrspace(1), ptr [[OUT_ADDR_ASCAST]], align 8
|
|
// CHECK-NEXT: store i32 [[TMP3]], ptr addrspace(1) [[TMP4]], align 4
|
|
// CHECK-NEXT: ret void
|
|
//
|
|
void test_permlane16_var(global uint* out, uint a, uint b, uint c) {
|
|
*out = __builtin_amdgcn_permlane16_var(a, b, c, 0, 0);
|
|
}
|
|
|
|
// CHECK-LABEL: @test_permlanex16_var(
|
|
// CHECK-NEXT: entry:
|
|
// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
|
|
// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
|
|
// CHECK-NEXT: [[B_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
|
|
// CHECK-NEXT: [[C_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
|
|
// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr
|
|
// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
|
|
// CHECK-NEXT: [[B_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[B_ADDR]] to ptr
|
|
// CHECK-NEXT: [[C_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[C_ADDR]] to ptr
|
|
// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr [[OUT_ADDR_ASCAST]], align 8
|
|
// CHECK-NEXT: store i32 [[A:%.*]], ptr [[A_ADDR_ASCAST]], align 4
|
|
// CHECK-NEXT: store i32 [[B:%.*]], ptr [[B_ADDR_ASCAST]], align 4
|
|
// CHECK-NEXT: store i32 [[C:%.*]], ptr [[C_ADDR_ASCAST]], align 4
|
|
// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR_ASCAST]], align 4
|
|
// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[B_ADDR_ASCAST]], align 4
|
|
// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[C_ADDR_ASCAST]], align 4
|
|
// CHECK-NEXT: [[TMP3:%.*]] = call i32 @llvm.amdgcn.permlanex16.var(i32 [[TMP0]], i32 [[TMP1]], i32 [[TMP2]], i1 false, i1 false)
|
|
// CHECK-NEXT: [[TMP4:%.*]] = load ptr addrspace(1), ptr [[OUT_ADDR_ASCAST]], align 8
|
|
// CHECK-NEXT: store i32 [[TMP3]], ptr addrspace(1) [[TMP4]], align 4
|
|
// CHECK-NEXT: ret void
|
|
//
|
|
void test_permlanex16_var(global uint* out, uint a, uint b, uint c) {
|
|
*out = __builtin_amdgcn_permlanex16_var(a, b, c, 0, 0);
|
|
}
|
|
|
|
// CHECK-LABEL: @test_s_barrier_signal(
|
|
// CHECK-NEXT: entry:
|
|
// CHECK-NEXT: call void @llvm.amdgcn.s.barrier.signal(i32 -1)
|
|
// CHECK-NEXT: call void @llvm.amdgcn.s.barrier.wait(i16 -1)
|
|
// CHECK-NEXT: ret void
|
|
//
|
|
void test_s_barrier_signal()
|
|
{
|
|
__builtin_amdgcn_s_barrier_signal(-1);
|
|
__builtin_amdgcn_s_barrier_wait(-1);
|
|
}
|
|
|
|
// CHECK-LABEL: @test_s_barrier_signal_var(
|
|
// CHECK-NEXT: entry:
|
|
// CHECK-NEXT: [[BAR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
|
|
// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
|
|
// CHECK-NEXT: [[BAR_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[BAR_ADDR]] to ptr
|
|
// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
|
|
// CHECK-NEXT: store ptr [[BAR:%.*]], ptr [[BAR_ADDR_ASCAST]], align 8
|
|
// CHECK-NEXT: store i32 [[A:%.*]], ptr [[A_ADDR_ASCAST]], align 4
|
|
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[BAR_ADDR_ASCAST]], align 8
|
|
// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(3)
|
|
// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[A_ADDR_ASCAST]], align 4
|
|
// CHECK-NEXT: call void @llvm.amdgcn.s.barrier.signal.var(ptr addrspace(3) [[TMP1]], i32 [[TMP2]])
|
|
// CHECK-NEXT: ret void
|
|
//
|
|
void test_s_barrier_signal_var(void *bar, int a)
|
|
{
|
|
__builtin_amdgcn_s_barrier_signal_var(bar, a);
|
|
}
|
|
|
|
// CHECK-LABEL: @test_s_barrier_signal_isfirst(
|
|
// CHECK-NEXT: entry:
|
|
// CHECK-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
|
|
// CHECK-NEXT: [[B_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
|
|
// CHECK-NEXT: [[C_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
|
|
// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
|
|
// CHECK-NEXT: [[B_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[B_ADDR]] to ptr
|
|
// CHECK-NEXT: [[C_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[C_ADDR]] to ptr
|
|
// CHECK-NEXT: store ptr [[A:%.*]], ptr [[A_ADDR_ASCAST]], align 8
|
|
// CHECK-NEXT: store ptr [[B:%.*]], ptr [[B_ADDR_ASCAST]], align 8
|
|
// CHECK-NEXT: store ptr [[C:%.*]], ptr [[C_ADDR_ASCAST]], align 8
|
|
// CHECK-NEXT: [[TMP0:%.*]] = call i1 @llvm.amdgcn.s.barrier.signal.isfirst(i32 1)
|
|
// CHECK-NEXT: br i1 [[TMP0]], label [[IF_THEN:%.*]], label [[IF_ELSE:%.*]]
|
|
// CHECK: if.then:
|
|
// CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[B_ADDR_ASCAST]], align 8
|
|
// CHECK-NEXT: store ptr [[TMP1]], ptr [[A_ADDR_ASCAST]], align 8
|
|
// CHECK-NEXT: br label [[IF_END:%.*]]
|
|
// CHECK: if.else:
|
|
// CHECK-NEXT: [[TMP2:%.*]] = load ptr, ptr [[C_ADDR_ASCAST]], align 8
|
|
// CHECK-NEXT: store ptr [[TMP2]], ptr [[A_ADDR_ASCAST]], align 8
|
|
// CHECK-NEXT: br label [[IF_END]]
|
|
// CHECK: if.end:
|
|
// CHECK-NEXT: call void @llvm.amdgcn.s.barrier.wait(i16 1)
|
|
// CHECK-NEXT: ret void
|
|
//
|
|
void test_s_barrier_signal_isfirst(int* a, int* b, int *c)
|
|
{
|
|
if(__builtin_amdgcn_s_barrier_signal_isfirst(1))
|
|
a = b;
|
|
else
|
|
a = c;
|
|
|
|
__builtin_amdgcn_s_barrier_wait(1);
|
|
}
|
|
|
|
// CHECK-LABEL: @test_s_get_barrier_state(
|
|
// CHECK-NEXT: entry:
|
|
// CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
|
|
// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
|
|
// CHECK-NEXT: [[STATE:%.*]] = alloca i32, align 4, addrspace(5)
|
|
// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
|
|
// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
|
|
// CHECK-NEXT: [[STATE_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[STATE]] to ptr
|
|
// CHECK-NEXT: store i32 [[A:%.*]], ptr [[A_ADDR_ASCAST]], align 4
|
|
// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR_ASCAST]], align 4
|
|
// CHECK-NEXT: [[TMP1:%.*]] = call i32 @llvm.amdgcn.s.get.barrier.state(i32 [[TMP0]])
|
|
// CHECK-NEXT: store i32 [[TMP1]], ptr [[STATE_ASCAST]], align 4
|
|
// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[STATE_ASCAST]], align 4
|
|
// CHECK-NEXT: ret i32 [[TMP2]]
|
|
//
|
|
unsigned test_s_get_barrier_state(int a)
|
|
{
|
|
unsigned State = __builtin_amdgcn_s_get_barrier_state(a);
|
|
return State;
|
|
}
|
|
|
|
// CHECK-LABEL: @test_s_get_named_barrier_state(
|
|
// CHECK-NEXT: entry:
|
|
// CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)
|
|
// CHECK-NEXT: [[BAR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
|
|
// CHECK-NEXT: [[STATE:%.*]] = alloca i32, align 4, addrspace(5)
|
|
// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr
|
|
// CHECK-NEXT: [[BAR_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[BAR_ADDR]] to ptr
|
|
// CHECK-NEXT: [[STATE_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[STATE]] to ptr
|
|
// CHECK-NEXT: store ptr [[BAR:%.*]], ptr [[BAR_ADDR_ASCAST]], align 8
|
|
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[BAR_ADDR_ASCAST]], align 8
|
|
// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(3)
|
|
// CHECK-NEXT: [[TMP2:%.*]] = call i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(3) [[TMP1]])
|
|
// CHECK-NEXT: store i32 [[TMP2]], ptr [[STATE_ASCAST]], align 4
|
|
// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr [[STATE_ASCAST]], align 4
|
|
// CHECK-NEXT: ret i32 [[TMP3]]
|
|
//
|
|
unsigned test_s_get_named_barrier_state(void *bar)
|
|
{
|
|
unsigned State = __builtin_amdgcn_s_get_named_barrier_state(bar);
|
|
return State;
|
|
}
|
|
|
|
// CHECK-LABEL: @test_s_ttracedata(
|
|
// CHECK-NEXT: entry:
|
|
// CHECK-NEXT: call void @llvm.amdgcn.s.ttracedata(i32 1)
|
|
// CHECK-NEXT: ret void
|
|
//
|
|
void test_s_ttracedata()
|
|
{
|
|
__builtin_amdgcn_s_ttracedata(1);
|
|
}
|
|
|
|
// CHECK-LABEL: @test_s_ttracedata_imm(
|
|
// CHECK-NEXT: entry:
|
|
// CHECK-NEXT: call void @llvm.amdgcn.s.ttracedata.imm(i16 1)
|
|
// CHECK-NEXT: ret void
|
|
//
|
|
void test_s_ttracedata_imm()
|
|
{
|
|
__builtin_amdgcn_s_ttracedata_imm(1);
|
|
}
|
|
|
|
// CHECK-LABEL: @test_s_prefetch_data(
|
|
// CHECK-NEXT: entry:
|
|
// CHECK-NEXT: [[FP_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
|
|
// CHECK-NEXT: [[GP_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
|
|
// CHECK-NEXT: [[CP_ADDR:%.*]] = alloca ptr addrspace(4), align 8, addrspace(5)
|
|
// CHECK-NEXT: [[LEN_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
|
|
// CHECK-NEXT: [[FP_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[FP_ADDR]] to ptr
|
|
// CHECK-NEXT: [[GP_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[GP_ADDR]] to ptr
|
|
// CHECK-NEXT: [[CP_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[CP_ADDR]] to ptr
|
|
// CHECK-NEXT: [[LEN_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[LEN_ADDR]] to ptr
|
|
// CHECK-NEXT: store ptr [[FP:%.*]], ptr [[FP_ADDR_ASCAST]], align 8
|
|
// CHECK-NEXT: store ptr addrspace(1) [[GP:%.*]], ptr [[GP_ADDR_ASCAST]], align 8
|
|
// CHECK-NEXT: store ptr addrspace(4) [[CP:%.*]], ptr [[CP_ADDR_ASCAST]], align 8
|
|
// CHECK-NEXT: store i32 [[LEN:%.*]], ptr [[LEN_ADDR_ASCAST]], align 4
|
|
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[FP_ADDR_ASCAST]], align 8
|
|
// CHECK-NEXT: call void @llvm.amdgcn.s.prefetch.data.p0(ptr [[TMP0]], i32 0)
|
|
// CHECK-NEXT: [[TMP1:%.*]] = load ptr addrspace(1), ptr [[GP_ADDR_ASCAST]], align 8
|
|
// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[LEN_ADDR_ASCAST]], align 4
|
|
// CHECK-NEXT: call void @llvm.amdgcn.s.prefetch.data.p1(ptr addrspace(1) [[TMP1]], i32 [[TMP2]])
|
|
// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(4), ptr [[CP_ADDR_ASCAST]], align 8
|
|
// CHECK-NEXT: call void @llvm.amdgcn.s.prefetch.data.p4(ptr addrspace(4) [[TMP3]], i32 31)
|
|
// CHECK-NEXT: ret void
|
|
//
|
|
void test_s_prefetch_data(int *fp, global float *gp, constant char *cp, unsigned int len)
|
|
{
|
|
__builtin_amdgcn_s_prefetch_data(fp, 0);
|
|
__builtin_amdgcn_s_prefetch_data(gp, len);
|
|
__builtin_amdgcn_s_prefetch_data(cp, 31);
|
|
}
|
|
|
|
// CHECK-LABEL: @test_s_buffer_prefetch_data(
|
|
// CHECK-NEXT: entry:
|
|
// CHECK-NEXT: [[RSRC_ADDR:%.*]] = alloca ptr addrspace(8), align 16, addrspace(5)
|
|
// CHECK-NEXT: [[LEN_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
|
|
// CHECK-NEXT: [[RSRC_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RSRC_ADDR]] to ptr
|
|
// CHECK-NEXT: [[LEN_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[LEN_ADDR]] to ptr
|
|
// CHECK-NEXT: store ptr addrspace(8) [[RSRC:%.*]], ptr [[RSRC_ADDR_ASCAST]], align 16
|
|
// CHECK-NEXT: store i32 [[LEN:%.*]], ptr [[LEN_ADDR_ASCAST]], align 4
|
|
// CHECK-NEXT: [[TMP0:%.*]] = load ptr addrspace(8), ptr [[RSRC_ADDR_ASCAST]], align 16
|
|
// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[LEN_ADDR_ASCAST]], align 4
|
|
// CHECK-NEXT: call void @llvm.amdgcn.s.buffer.prefetch.data(ptr addrspace(8) [[TMP0]], i32 128, i32 [[TMP1]])
|
|
// CHECK-NEXT: [[TMP2:%.*]] = load ptr addrspace(8), ptr [[RSRC_ADDR_ASCAST]], align 16
|
|
// CHECK-NEXT: call void @llvm.amdgcn.s.buffer.prefetch.data(ptr addrspace(8) [[TMP2]], i32 0, i32 31)
|
|
// CHECK-NEXT: ret void
|
|
//
|
|
void test_s_buffer_prefetch_data(__amdgpu_buffer_rsrc_t rsrc, unsigned int len)
|
|
{
|
|
__builtin_amdgcn_s_buffer_prefetch_data(rsrc, 128, len);
|
|
__builtin_amdgcn_s_buffer_prefetch_data(rsrc, 0, 31);
|
|
}
|
|
|
|
// CHECK-LABEL: @test_ds_bpermute_fi_b32(
|
|
// CHECK-NEXT: entry:
|
|
// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
|
|
// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
|
|
// CHECK-NEXT: [[B_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
|
|
// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr
|
|
// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
|
|
// CHECK-NEXT: [[B_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[B_ADDR]] to ptr
|
|
// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr [[OUT_ADDR_ASCAST]], align 8
|
|
// CHECK-NEXT: store i32 [[A:%.*]], ptr [[A_ADDR_ASCAST]], align 4
|
|
// CHECK-NEXT: store i32 [[B:%.*]], ptr [[B_ADDR_ASCAST]], align 4
|
|
// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR_ASCAST]], align 4
|
|
// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[B_ADDR_ASCAST]], align 4
|
|
// CHECK-NEXT: [[TMP2:%.*]] = call i32 @llvm.amdgcn.ds.bpermute.fi.b32(i32 [[TMP0]], i32 [[TMP1]])
|
|
// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(1), ptr [[OUT_ADDR_ASCAST]], align 8
|
|
// CHECK-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[TMP3]], align 4
|
|
// CHECK-NEXT: ret void
|
|
//
|
|
void test_ds_bpermute_fi_b32(global int* out, int a, int b)
|
|
{
|
|
*out = __builtin_amdgcn_ds_bpermute_fi_b32(a, b);
|
|
}
|