maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

AMDGPU: Remove gcc builtin names from workitem intrinsics 

We'll need to emit these manually in clang to add range metadata

Reviewers: arsenm

Differential Revision: http://reviews.llvm.org/D20691

llvm-svn: 273318
This commit is contained in:
Jan Vesely 2016-06-21 20:46:22 +00:00
parent fea814d531
commit 26c6bb103f
1 changed files with 30 additions and 26 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -11,28 +11,37 @@
//
//===----------------------------------------------------------------------===//
class AMDGPUReadPreloadRegisterIntrinsic<string name>
: Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>,
GCCBuiltin<name>;
class AMDGPUReadPreloadRegisterIntrinsic
: Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>;
class AMDGPUReadPreloadRegisterIntrinsicNamed<string name>
: Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>, GCCBuiltin<name>;
let TargetPrefix = "r600" in {
multiclass AMDGPUReadPreloadRegisterIntrinsic_xyz<string prefix> {
def _x : AMDGPUReadPreloadRegisterIntrinsic<!strconcat(prefix, "_x")>;
def _y : AMDGPUReadPreloadRegisterIntrinsic<!strconcat(prefix, "_y")>;
def _z : AMDGPUReadPreloadRegisterIntrinsic<!strconcat(prefix, "_z")>;
multiclass AMDGPUReadPreloadRegisterIntrinsic_xyz {
def _x : AMDGPUReadPreloadRegisterIntrinsic;
def _y : AMDGPUReadPreloadRegisterIntrinsic;
def _z : AMDGPUReadPreloadRegisterIntrinsic;
}
defm int_r600_read_global_size : AMDGPUReadPreloadRegisterIntrinsic_xyz <
"__builtin_r600_read_global_size">;
defm int_r600_read_local_size : AMDGPUReadPreloadRegisterIntrinsic_xyz <
"__builtin_r600_read_local_size">;
defm int_r600_read_ngroups : AMDGPUReadPreloadRegisterIntrinsic_xyz <
"__builtin_r600_read_ngroups">;
defm int_r600_read_tgid : AMDGPUReadPreloadRegisterIntrinsic_xyz <
"__builtin_r600_read_tgid">;
defm int_r600_read_tidig : AMDGPUReadPreloadRegisterIntrinsic_xyz <
"__builtin_r600_read_tidig">;
multiclass AMDGPUReadPreloadRegisterIntrinsic_xyz_named<string prefix> {
def _x : AMDGPUReadPreloadRegisterIntrinsicNamed<!strconcat(prefix, "_x")>;
def _y : AMDGPUReadPreloadRegisterIntrinsicNamed<!strconcat(prefix, "_y")>;
def _z : AMDGPUReadPreloadRegisterIntrinsicNamed<!strconcat(prefix, "_z")>;
}
defm int_r600_read_global_size : AMDGPUReadPreloadRegisterIntrinsic_xyz_named
<"__builtin_r600_read_global_size">;
defm int_r600_read_ngroups : AMDGPUReadPreloadRegisterIntrinsic_xyz_named
<"__builtin_r600_read_ngroups">;
defm int_r600_read_tgid : AMDGPUReadPreloadRegisterIntrinsic_xyz_named
<"__builtin_r600_read_tgid">;
defm int_r600_read_local_size : AMDGPUReadPreloadRegisterIntrinsic_xyz;
defm int_r600_read_tidig : AMDGPUReadPreloadRegisterIntrinsic_xyz;
def int_r600_read_workdim : AMDGPUReadPreloadRegisterIntrinsic;
def int_r600_rat_store_typed :
// 1st parameter: Data
@ -45,9 +54,6 @@ def int_r600_rsq : Intrinsic<
[llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem]
>;
def int_r600_read_workdim : AMDGPUReadPreloadRegisterIntrinsic <
"__builtin_r600_read_workdim"
>;
} // End TargetPrefix = "r600"
@ -60,10 +66,9 @@ def int_AMDGPU_ldexp : Intrinsic<
let TargetPrefix = "amdgcn" in {
defm int_amdgcn_workitem_id : AMDGPUReadPreloadRegisterIntrinsic_xyz <
"__builtin_amdgcn_workitem_id">;
defm int_amdgcn_workgroup_id : AMDGPUReadPreloadRegisterIntrinsic_xyz <
"__builtin_amdgcn_workgroup_id">;
defm int_amdgcn_workitem_id : AMDGPUReadPreloadRegisterIntrinsic_xyz;
defm int_amdgcn_workgroup_id : AMDGPUReadPreloadRegisterIntrinsic_xyz_named
<"__builtin_amdgcn_workgroup_id">;
def int_amdgcn_s_barrier : GCCBuiltin<"__builtin_amdgcn_s_barrier">,
Intrinsic<[], [], [IntrConvergent]>;
@ -289,8 +294,7 @@ def int_amdgcn_buffer_atomic_cmpswap : Intrinsic<
llvm_i1_ty], // slc(imm)
[]>;
def int_amdgcn_read_workdim : AMDGPUReadPreloadRegisterIntrinsic <
"__builtin_amdgcn_read_workdim">;
def int_amdgcn_read_workdim : AMDGPUReadPreloadRegisterIntrinsic;
def int_amdgcn_buffer_wbinvl1_sc :