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

AMDGPU: Add builtin for getreg intrinsic 

llvm-svn: 292636
This commit is contained in:
Matt Arsenault 2017-01-20 19:24:22 +00:00
parent 6342baf2b1
commit 24b5ae4497
3 changed files with 17 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -35,6 +35,7 @@ BUILTIN(__builtin_amdgcn_workitem_id_z, "Ui", "nc")
//===----------------------------------------------------------------------===//
// Instruction builtins.
//===----------------------------------------------------------------------===//
BUILTIN(__builtin_amdgcn_s_getreg, "UiIi", "n")
BUILTIN(__builtin_amdgcn_s_barrier, "v", "n")
BUILTIN(__builtin_amdgcn_wave_barrier, "v", "n")
BUILTIN(__builtin_amdgcn_div_scale, "dddbb*", "n")

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

@ -371,6 +371,17 @@ void test_get_group_id(int d, global int *out)
}
}
// CHECK-LABEL: @test_s_getreg(
// CHECK: tail call i32 @llvm.amdgcn.s.getreg(i32 0)
// CHECK: tail call i32 @llvm.amdgcn.s.getreg(i32 1)
// CHECK: tail call i32 @llvm.amdgcn.s.getreg(i32 65535)
void test_s_getreg(volatile global uint *out)
{
*out = __builtin_amdgcn_s_getreg(0);
*out = __builtin_amdgcn_s_getreg(1);
*out = __builtin_amdgcn_s_getreg(65535);
}
// CHECK-LABEL: @test_get_local_id(
// CHECK: tail call i32 @llvm.amdgcn.workitem.id.x(), !range [[WI_RANGE:![0-9]*]]
// CHECK: tail call i32 @llvm.amdgcn.workitem.id.y(), !range [[WI_RANGE]]

5
clang/test/SemaOpenCL/builtins-amdgcn-error.cl
View File

@ -62,3 +62,8 @@ void test_ds_swizzle(global int* out, int a, int b)
{
*out = __builtin_amdgcn_ds_swizzle(a, b); // expected-error {{argument to '__builtin_amdgcn_ds_swizzle' must be a constant integer}}
}
void test_s_getreg(global int* out, int a)
{
*out = __builtin_amdgcn_s_getreg(a); // expected-error {{argument to '__builtin_amdgcn_s_getreg' must be a constant integer}}
}