RPCS3/llvm-mirror
1
0
Fork 0
mirror of https://github.com/RPCS3/llvm-mirror.git synced 2024-11-22 10:42:39 +01:00
Code Issues Projects Releases Wiki Activity

AMDGPU: Add new amdgcn intrinsics for cube instructions

Browse Source 
More cleanup to try to get all intrinsics using the correct
amdgcn prefix that are as close to the instruction as possible.

llvm-svn: 258786
This commit is contained in:
Matt Arsenault 2016-01-26 04:29:56 +00:00
parent 667cd15c1c
commit 581518df24
9 changed files with 136 additions and 7 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

20
include/llvm/IR/IntrinsicsAMDGPU.td
View File

@ -110,6 +110,26 @@ def int_amdgcn_class : Intrinsic<
[llvm_i1_ty], [llvm_anyfloat_ty, llvm_i32_ty], [IntrNoMem]
>;
def int_amdgcn_cubeid : GCCBuiltin<"__builtin_amdgcn_cubeid">,
Intrinsic<[llvm_float_ty],
[llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem]
>;
def int_amdgcn_cubema : GCCBuiltin<"__builtin_amdgcn_cubema">,
Intrinsic<[llvm_float_ty],
[llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem]
>;
def int_amdgcn_cubesc : GCCBuiltin<"__builtin_amdgcn_cubesc">,
Intrinsic<[llvm_float_ty],
[llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem]
>;
def int_amdgcn_cubetc : GCCBuiltin<"__builtin_amdgcn_cubetc">,
Intrinsic<[llvm_float_ty],
[llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem]
>;
def int_amdgcn_read_workdim : AMDGPUReadPreloadRegisterIntrinsic <
"__builtin_amdgcn_read_workdim">;

6
lib/Target/AMDGPU/AMDGPUIntrinsics.td
View File

@ -33,12 +33,16 @@ let TargetPrefix = "AMDGPU", isTarget = 1 in {
def int_AMDGPU_cvt_f32_ubyte1 : Intrinsic<[llvm_float_ty], [llvm_i32_ty], [IntrNoMem]>;
def int_AMDGPU_cvt_f32_ubyte2 : Intrinsic<[llvm_float_ty], [llvm_i32_ty], [IntrNoMem]>;
def int_AMDGPU_cvt_f32_ubyte3 : Intrinsic<[llvm_float_ty], [llvm_i32_ty], [IntrNoMem]>;
def int_AMDGPU_flbit_i32 : Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem]>;
// Deprecated in favor of separate int_amdgcn_cube* intrinsics.
def int_AMDGPU_cube : Intrinsic<[llvm_v4f32_ty], [llvm_v4f32_ty], [IntrNoMem]>;
// Deprecated in favor of expanded bit operations
def int_AMDGPU_bfi : Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
def int_AMDGPU_bfe_i32 : Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
def int_AMDGPU_bfe_u32 : Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
def int_AMDGPU_bfm : Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
def int_AMDGPU_flbit_i32 : Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem]>;
def int_AMDGPU_rsq_clamped : Intrinsic<
[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]

8
lib/Target/AMDGPU/SIInstructions.td
View File

@ -1636,16 +1636,16 @@ defm V_MAD_U32_U24 : VOP3Inst <vop3<0x143, 0x1c3>, "v_mad_u32_u24",
} // End isCommutable = 1
defm V_CUBEID_F32 : VOP3Inst <vop3<0x144, 0x1c4>, "v_cubeid_f32",
VOP_F32_F32_F32_F32
VOP_F32_F32_F32_F32, int_amdgcn_cubeid
>;
defm V_CUBESC_F32 : VOP3Inst <vop3<0x145, 0x1c5>, "v_cubesc_f32",
VOP_F32_F32_F32_F32
VOP_F32_F32_F32_F32, int_amdgcn_cubesc
>;
defm V_CUBETC_F32 : VOP3Inst <vop3<0x146, 0x1c6>, "v_cubetc_f32",
VOP_F32_F32_F32_F32
VOP_F32_F32_F32_F32, int_amdgcn_cubetc
>;
defm V_CUBEMA_F32 : VOP3Inst <vop3<0x147, 0x1c7>, "v_cubema_f32",
VOP_F32_F32_F32_F32
VOP_F32_F32_F32_F32, int_amdgcn_cubema
>;
defm V_BFE_U32 : VOP3Inst <vop3<0x148, 0x1c8>, "v_bfe_u32",

46
test/CodeGen/AMDGPU/cube.ll Normal file
View File

@ -0,0 +1,46 @@
; RUN: llc -march=amdgcn -mcpu=tahiti -verify-machineinstrs < %s | FileCheck -check-prefix=GCN -check-prefix=SI %s
; RUN: llc -march=amdgcn -mcpu=fiji -verify-machineinstrs < %s | FileCheck -check-prefix=GCN -check-prefix=VI %s
declare float @llvm.amdgcn.cubeid(float, float, float) #0
declare float @llvm.amdgcn.cubesc(float, float, float) #0
declare float @llvm.amdgcn.cubetc(float, float, float) #0
declare float @llvm.amdgcn.cubema(float, float, float) #0
declare <4 x float> @llvm.AMDGPU.cube(<4 x float>) #0
; GCN-LABEL: {{^}}cube:
; GCN-DAG: v_cubeid_f32 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
; GCN-DAG: v_cubesc_f32 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
; GCN-DAG: v_cubetc_f32 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
; GCN-DAG: v_cubema_f32 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
; GCN: buffer_store_dwordx4
define void @cube(<4 x float> addrspace(1)* %out, float %a, float %b, float %c) #1 {
%cubeid = call float @llvm.amdgcn.cubeid(float %a, float %b, float %c)
%cubesc = call float @llvm.amdgcn.cubesc(float %a, float %b, float %c)
%cubetc = call float @llvm.amdgcn.cubetc(float %a, float %b, float %c)
%cubema = call float @llvm.amdgcn.cubema(float %a, float %b, float %c)
%vec0 = insertelement <4 x float> undef, float %cubeid, i32 0
%vec1 = insertelement <4 x float> %vec0, float %cubesc, i32 1
%vec2 = insertelement <4 x float> %vec1, float %cubetc, i32 2
%vec3 = insertelement <4 x float> %vec2, float %cubema, i32 3
store <4 x float> %vec3, <4 x float> addrspace(1)* %out
ret void
}
; GCN-LABEL: {{^}}legacy_cube:
; GCN-DAG: v_cubeid_f32 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
; GCN-DAG: v_cubesc_f32 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
; GCN-DAG: v_cubetc_f32 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
; GCN-DAG: v_cubema_f32 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
; GCN: buffer_store_dwordx4
define void @legacy_cube(<4 x float> addrspace(1)* %out, <4 x float> %abcx) #1 {
%cube = call <4 x float> @llvm.AMDGPU.cube(<4 x float> %abcx)
store <4 x float> %cube, <4 x float> addrspace(1)* %out
ret void
}
attributes #0 = { nounwind readnone }
attributes #1 = { nounwind }

3
test/CodeGen/AMDGPU/llvm.AMDGPU.cube.ll
View File

@ -1,7 +1,6 @@
; RUN: llc < %s -march=r600 -mcpu=redwood | FileCheck %s
; CHECK: {{^}}cube:
; CHECK-LABEL: {{^}}cube:
; CHECK: CUBE T{{[0-9]}}.X
; CHECK: CUBE T{{[0-9]}}.Y
; CHECK: CUBE T{{[0-9]}}.Z

15
test/CodeGen/AMDGPU/llvm.amdgcn.cubeid.ll Normal file
View File

@ -0,0 +1,15 @@
; RUN: llc -march=amdgcn -mcpu=tahiti -verify-machineinstrs < %s | FileCheck -check-prefix=GCN -check-prefix=SI %s
; RUN: llc -march=amdgcn -mcpu=fiji -verify-machineinstrs < %s | FileCheck -check-prefix=GCN -check-prefix=VI %s
declare float @llvm.amdgcn.cubeid(float, float, float) #0
; GCN-LABEL: {{^}}test_cubeid:
; GCN: v_cubeid_f32 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
define void @test_cubeid(float addrspace(1)* %out, float %a, float %b, float %c) #1 {
%result = call float @llvm.amdgcn.cubeid(float %a, float %b, float %c)
store float %result, float addrspace(1)* %out
ret void
}
attributes #0 = { nounwind readnone }
attributes #1 = { nounwind }

15
test/CodeGen/AMDGPU/llvm.amdgcn.cubema.ll Normal file
View File

@ -0,0 +1,15 @@
; RUN: llc -march=amdgcn -mcpu=tahiti -verify-machineinstrs < %s | FileCheck -check-prefix=GCN -check-prefix=SI %s
; RUN: llc -march=amdgcn -mcpu=fiji -verify-machineinstrs < %s | FileCheck -check-prefix=GCN -check-prefix=VI %s
declare float @llvm.amdgcn.cubema(float, float, float) #0
; GCN-LABEL: {{^}}test_cubema:
; GCN: v_cubema_f32 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
define void @test_cubema(float addrspace(1)* %out, float %a, float %b, float %c) #1 {
%result = call float @llvm.amdgcn.cubema(float %a, float %b, float %c)
store float %result, float addrspace(1)* %out
ret void
}
attributes #0 = { nounwind readnone }
attributes #1 = { nounwind }

15
test/CodeGen/AMDGPU/llvm.amdgcn.cubesc.ll Normal file
View File

@ -0,0 +1,15 @@
; RUN: llc -march=amdgcn -mcpu=tahiti -verify-machineinstrs < %s | FileCheck -check-prefix=GCN -check-prefix=SI %s
; RUN: llc -march=amdgcn -mcpu=fiji -verify-machineinstrs < %s | FileCheck -check-prefix=GCN -check-prefix=VI %s
declare float @llvm.amdgcn.cubesc(float, float, float) #0
; GCN-LABEL: {{^}}test_cubesc:
; GCN: v_cubesc_f32 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
define void @test_cubesc(float addrspace(1)* %out, float %a, float %b, float %c) #1 {
%result = call float @llvm.amdgcn.cubesc(float %a, float %b, float %c)
store float %result, float addrspace(1)* %out
ret void
}
attributes #0 = { nounwind readnone }
attributes #1 = { nounwind }

15
test/CodeGen/AMDGPU/llvm.amdgcn.cubetc.ll Normal file
View File

@ -0,0 +1,15 @@
; RUN: llc -march=amdgcn -mcpu=tahiti -verify-machineinstrs < %s | FileCheck -check-prefix=GCN -check-prefix=SI %s
; RUN: llc -march=amdgcn -mcpu=fiji -verify-machineinstrs < %s | FileCheck -check-prefix=GCN -check-prefix=VI %s
declare float @llvm.amdgcn.cubetc(float, float, float) #0
; GCN-LABEL: {{^}}test_cubetc:
; GCN: v_cubetc_f32 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
define void @test_cubetc(float addrspace(1)* %out, float %a, float %b, float %c) #1 {
%result = call float @llvm.amdgcn.cubetc(float %a, float %b, float %c)
store float %result, float addrspace(1)* %out
ret void
}
attributes #0 = { nounwind readnone }
attributes #1 = { nounwind }