diff options
| author | Austin Kerbow <Austin.Kerbow@amd.com> | 2022-03-04 10:59:26 +0300 |
|---|---|---|
| committer | Austin Kerbow <Austin.Kerbow@amd.com> | 2022-03-13 09:15:42 +0300 |
| commit | 62bcfcb5a588e5e844f8e4e42a2e4d15c907a746 (patch) | |
| tree | f80654be03e234bcebc69a32ad8152401b2afce0 /clang | |
| parent | e825f49b80df3479108171d6aac3208965744d95 (diff) | |
[AMDGPU] Add llvm.amdgcn.s.setprio intrinsic
Reviewed By: rampitec, arsenm
Differential Revision: https://reviews.llvm.org/D120976
Diffstat (limited to 'clang')
| -rw-r--r-- | clang/include/clang/Basic/BuiltinsAMDGPU.def | 1 | ||||
| -rw-r--r-- | clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 9 | ||||
| -rw-r--r-- | clang/test/SemaOpenCL/builtins-amdgcn-error.cl | 6 |
3 files changed, 16 insertions, 0 deletions
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index 3b7ff75a9410..d2e60f85b9fe 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -116,6 +116,7 @@ BUILTIN(__builtin_amdgcn_cubema, "ffff", "nc") BUILTIN(__builtin_amdgcn_s_sleep, "vIi", "n") BUILTIN(__builtin_amdgcn_s_incperflevel, "vIi", "n") BUILTIN(__builtin_amdgcn_s_decperflevel, "vIi", "n") +BUILTIN(__builtin_amdgcn_s_setprio, "vIs", "n") BUILTIN(__builtin_amdgcn_uicmp, "WUiUiUiIi", "nc") BUILTIN(__builtin_amdgcn_uicmpl, "WUiWUiWUiIi", "nc") BUILTIN(__builtin_amdgcn_sicmp, "WUiiiIi", "nc") diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl index b02e6308c343..ee6d2e0589c5 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -423,6 +423,15 @@ void test_s_decperflevel() __builtin_amdgcn_s_decperflevel(15); } +// CHECK-LABEL: @test_s_setprio +// CHECK: call void @llvm.amdgcn.s.setprio(i16 0) +// CHECK: call void @llvm.amdgcn.s.setprio(i16 3) +void test_s_setprio() +{ + __builtin_amdgcn_s_setprio(0); + __builtin_amdgcn_s_setprio(3); +} + // CHECK-LABEL: @test_cubeid( // CHECK: call float @llvm.amdgcn.cubeid(float %a, float %b, float %c) void test_cubeid(global float* out, float a, float b, float c) { diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error.cl index 20a520812e3f..af4fdf32c272 100644 --- a/clang/test/SemaOpenCL/builtins-amdgcn-error.cl +++ b/clang/test/SemaOpenCL/builtins-amdgcn-error.cl @@ -54,6 +54,12 @@ void test_s_decperflevel(int x) __builtin_amdgcn_s_decperflevel(x); // expected-error {{argument to '__builtin_amdgcn_s_decperflevel' must be a constant integer}} } +void test_s_setprio(int x) +{ + __builtin_amdgcn_s_setprio(x); // expected-error {{argument to '__builtin_amdgcn_s_setprio' must be a constant integer}} + __builtin_amdgcn_s_setprio(65536); // expected-warning {{implicit conversion from 'int' to 'short' changes value from 65536 to 0}} +} + void test_sicmp_i32(global ulong* out, int a, int b, uint c) { *out = __builtin_amdgcn_sicmp(a, b, c); // expected-error {{argument to '__builtin_amdgcn_sicmp' must be a constant integer}} |