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() { _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
