[clang] 4efbe98 - [Clang][AMDGPU] Add a Sema check for the imm argument of ` __builtin_amdgcn_s_setreg` (#176838)

[via cfe-commits]

Author: Shilei Tian
Date: 2026-01-20T11:48:52-05:00
New Revision: 4efbe986596afee43848e8f5e0d33c00ef6462a8

URL: https://github.com/llvm/llvm-project/commit/4efbe986596afee43848e8f5e0d33c00ef6462a8
DIFF: https://github.com/llvm/llvm-project/commit/4efbe986596afee43848e8f5e0d33c00ef6462a8.diff

LOG: [Clang][AMDGPU] Add a Sema check for the imm argument of ` __builtin_amdgcn_s_setreg` (#176838)

Our backend cannot select the corresponding intrinsic if the imm
argument is not a `int16_t` or `uint16_t`, which is not really helpful.

Added: 
    

Modified: 
    clang/lib/Sema/SemaAMDGPU.cpp
    clang/test/CodeGenOpenCL/builtins-amdgcn.cl
    clang/test/SemaOpenCL/builtins-amdgcn-error.cl

Removed: 
    


################################################################################
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/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(

diff  --git a/clang/test/SemaOpenCL/builtins-amdgcn-error.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error.cl
index 7a550f026bc1b..eb1a86bdcdeb0 100644
--- a/clang/test/SemaOpenCL/builtins-amdgcn-error.cl
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-error.cl
@@ -166,6 +166,8 @@ 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]}}
+  __builtin_amdgcn_s_setreg(65536, y); // expected-error {{argument value 65536 is outside the valid range [0, 65535]}}
 }
 
 void test_atomic_inc32() {