diff options
| author | Matt Arsenault <Matthew.Arsenault@amd.com> | 2017-10-09 20:06:37 +0000 |
|---|---|---|
| committer | Matt Arsenault <Matthew.Arsenault@amd.com> | 2017-10-09 20:06:37 +0000 |
| commit | 2feacd003b61f030266f29fdaa5c29b525f46ae1 (patch) | |
| tree | 1fd07cac0999fbfc3ac9b040781b66fa63282ad0 | |
| parent | 21e1fad25b70e1a6cf1d5f61a2d5c39765a92900 (diff) | |
| download | clang-2feacd003b61f030266f29fdaa5c29b525f46ae1.tar.gz | |
AMDGPU: Add read_exec_lo/hi builtins
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@315238 91177308-0d34-0410-b5e6-96231b3b80d8
| -rw-r--r-- | include/clang/Basic/BuiltinsAMDGPU.def | 2 | ||||
| -rw-r--r-- | lib/CodeGen/CGBuiltin.cpp | 9 | ||||
| -rw-r--r-- | test/CodeGenOpenCL/builtins-amdgcn.cl | 14 |
3 files changed, 25 insertions, 0 deletions
diff --git a/include/clang/Basic/BuiltinsAMDGPU.def b/include/clang/Basic/BuiltinsAMDGPU.def index 3e5b0d13ae..ec6a0fb917 100644 --- a/include/clang/Basic/BuiltinsAMDGPU.def +++ b/include/clang/Basic/BuiltinsAMDGPU.def @@ -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. diff --git a/lib/CodeGen/CGBuiltin.cpp b/lib/CodeGen/CGBuiltin.cpp index 7ae558f9ae..a4e6452a7d 100644 --- a/lib/CodeGen/CGBuiltin.cpp +++ b/lib/CodeGen/CGBuiltin.cpp @@ -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: diff --git a/test/CodeGenOpenCL/builtins-amdgcn.cl b/test/CodeGenOpenCL/builtins-amdgcn.cl index 65b0666ad1..9f036547bf 100644 --- a/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -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"} |