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

AMDGPU: Add read_exec_lo/hi builtins 

llvm-svn: 315238
This commit is contained in:
Matt Arsenault 2017-10-09 20:06:37 +00:00
parent c1d5955684
commit f12e3b848a
3 changed files with 25 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
clang/include/clang/Basic/BuiltinsAMDGPU.def
View File

@ -121,6 +121,8 @@ TARGET_BUILTIN(__builtin_amdgcn_fmed3h, "hhhh", "nc", "gfx9-insts")
// Special builtins.
//===----------------------------------------------------------------------===//
BUILTIN(__builtin_amdgcn_read_exec, "LUi", "nc")
BUILTIN(__builtin_amdgcn_read_exec_lo, "Ui", "nc")
BUILTIN(__builtin_amdgcn_read_exec_hi, "Ui", "nc")
//===----------------------------------------------------------------------===//
// R600-NI only builtins.

9
clang/lib/CodeGen/CGBuiltin.cpp
View File

@ -9103,6 +9103,15 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
CI->setConvergent();
return CI;
}
case AMDGPU::BI__builtin_amdgcn_read_exec_lo:
case AMDGPU::BI__builtin_amdgcn_read_exec_hi: {
StringRef RegName = BuiltinID == AMDGPU::BI__builtin_amdgcn_read_exec_lo ?
"exec_lo" : "exec_hi";
CallInst *CI = cast<CallInst>(
EmitSpecialRegisterBuiltin(*this, E, Int32Ty, Int32Ty, true, RegName));
CI->setConvergent();
return CI;
}
// amdgcn workitem
case AMDGPU::BI__builtin_amdgcn_workitem_id_x:

14
clang/test/CodeGenOpenCL/builtins-amdgcn.cl
View File

@ -421,6 +421,18 @@ void test_read_exec(global ulong* out) {
// CHECK: declare i64 @llvm.read_register.i64(metadata) #[[NOUNWIND_READONLY:[0-9]+]]
// CHECK-LABEL: @test_read_exec_lo(
// CHECK: call i32 @llvm.read_register.i32(metadata ![[EXEC_LO:[0-9]+]]) #[[READ_EXEC_ATTRS]]
void test_read_exec_lo(global uint* out) {
*out = __builtin_amdgcn_read_exec_lo();
}
// CHECK-LABEL: @test_read_exec_hi(
// CHECK: call i32 @llvm.read_register.i32(metadata ![[EXEC_HI:[0-9]+]]) #[[READ_EXEC_ATTRS]]
void test_read_exec_hi(global uint* out) {
*out = __builtin_amdgcn_read_exec_hi();
}
// CHECK-LABEL: @test_dispatch_ptr
// CHECK: call i8 addrspace(2)* @llvm.amdgcn.dispatch.ptr()
void test_dispatch_ptr(__attribute__((address_space(2))) unsigned char ** out)
@ -499,3 +511,5 @@ void test_s_getpc(global ulong* out)
// CHECK-DAG: attributes #[[NOUNWIND_READONLY:[0-9]+]] = { nounwind readonly }
// CHECK-DAG: attributes #[[READ_EXEC_ATTRS]] = { convergent }
// CHECK-DAG: ![[EXEC]] = !{!"exec"}
// CHECK-DAG: ![[EXEC_LO]] = !{!"exec_lo"}
// CHECK-DAG: ![[EXEC_HI]] = !{!"exec_hi"}