https://github.com/shiltian updated https://github.com/llvm/llvm-project/pull/176838
>From f83fe1d94945f2301a6e1b899eb5d935577ffd56 Mon Sep 17 00:00:00 2001 From: Shilei Tian <[email protected]> Date: Mon, 19 Jan 2026 18:04:35 -0500 Subject: [PATCH 1/3] [Clang][AMDGPU] Add a Sema check for the imm argument of ` __builtin_amdgcn_s_setreg` Our backend cannot select the corresponding intrinsic if the imm argument is not a `int16_t` or `uint16_t`, which is not really helpful. --- clang/lib/Sema/SemaAMDGPU.cpp | 3 +++ clang/test/SemaOpenCL/builtins-amdgcn-error.cl | 1 + 2 files changed, 4 insertions(+) diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp index b6eebf35296ef..4261e1849133f 100644 --- a/clang/lib/Sema/SemaAMDGPU.cpp +++ b/clang/lib/Sema/SemaAMDGPU.cpp @@ -86,6 +86,9 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID, OrderIndex = 0; ScopeIndex = 1; break; + case AMDGPU::BI__builtin_amdgcn_s_setreg: + return SemaRef.BuiltinConstantArgRange(TheCall, /*ArgNum=*/0, /*Low=*/0, + /*High=*/UINT16_MAX); case AMDGPU::BI__builtin_amdgcn_mov_dpp: return checkMovDPPFunctionCall(TheCall, 5, 1); case AMDGPU::BI__builtin_amdgcn_mov_dpp8: diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error.cl index 7a550f026bc1b..13c40b9301237 100644 --- a/clang/test/SemaOpenCL/builtins-amdgcn-error.cl +++ b/clang/test/SemaOpenCL/builtins-amdgcn-error.cl @@ -166,6 +166,7 @@ void test_fence() { void test_s_setreg(int x, int y) { __builtin_amdgcn_s_setreg(x, 0); // expected-error {{argument to '__builtin_amdgcn_s_setreg' must be a constant integer}} __builtin_amdgcn_s_setreg(x, y); // expected-error {{argument to '__builtin_amdgcn_s_setreg' must be a constant integer}} + __builtin_amdgcn_s_setreg(193768, y); // expected-error {{argument value 193768 is outside the valid range [0, 65535]}} } void test_atomic_inc32() { >From 909ce51a1013ac1299b1f86f35ac2d22fdb0470c Mon Sep 17 00:00:00 2001 From: Shilei Tian <[email protected]> Date: Mon, 19 Jan 2026 18:27:59 -0500 Subject: [PATCH 2/3] fix comment --- clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl index 80b585513f71a..04140ed3f10b0 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -1229,9 +1229,9 @@ kernel void test_mqsad_u32_u8(global uint4* out, ulong src0, uint src1, uint4 sr } // CHECK-LABEL: test_s_setreg( -// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.setreg(i32 8193, i32 %val) +// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.setreg(i32 65535, i32 %val) kernel void test_s_setreg(uint val) { - __builtin_amdgcn_s_setreg(8193, val); + __builtin_amdgcn_s_setreg(65535, val); } // CHECK-LABEL test_atomic_inc_dec( >From 769e13807c4401e87a593a0e9ed5a36a31ea3c68 Mon Sep 17 00:00:00 2001 From: Shilei Tian <[email protected]> Date: Mon, 19 Jan 2026 18:46:30 -0500 Subject: [PATCH 3/3] add one more negative test --- clang/test/SemaOpenCL/builtins-amdgcn-error.cl | 1 + 1 file changed, 1 insertion(+) diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error.cl index 13c40b9301237..eb1a86bdcdeb0 100644 --- a/clang/test/SemaOpenCL/builtins-amdgcn-error.cl +++ b/clang/test/SemaOpenCL/builtins-amdgcn-error.cl @@ -167,6 +167,7 @@ void test_s_setreg(int x, int y) { __builtin_amdgcn_s_setreg(x, 0); // expected-error {{argument to '__builtin_amdgcn_s_setreg' must be a constant integer}} __builtin_amdgcn_s_setreg(x, y); // expected-error {{argument to '__builtin_amdgcn_s_setreg' must be a constant integer}} __builtin_amdgcn_s_setreg(193768, y); // expected-error {{argument value 193768 is outside the valid range [0, 65535]}} + __builtin_amdgcn_s_setreg(65536, y); // expected-error {{argument value 65536 is outside the valid range [0, 65535]}} } void test_atomic_inc32() { _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
