https://github.com/arsenm created https://github.com/llvm/llvm-project/pull/220077
Reject an unrecognized synchronization scope string passed to __builtin_amdgcn_fence during semantic analysis instead of relying on the backend to report it during codegen. Co-authored-by: Claude (Claude-Opus-4.8) <[email protected]> >From c553573aa3360e994a71c7600a344d38834cc37a Mon Sep 17 00:00:00 2001 From: Matt Arsenault <[email protected]> Date: Mon, 31 Aug 2026 21:57:27 +0200 Subject: [PATCH] clang/AMDGPU: Diagnose invalid fence sync scope Reject an unrecognized synchronization scope string passed to __builtin_amdgcn_fence during semantic analysis instead of relying on the backend to report it during codegen. Co-authored-by: Claude (Claude-Opus-4.8) <[email protected]> --- clang/include/clang/Basic/DiagnosticSemaKinds.td | 2 ++ clang/lib/Sema/SemaAMDGPU.cpp | 12 ++++++++++++ clang/test/Sema/builtin-amdgcn-fence-failure.cpp | 6 ++---- 3 files changed, 16 insertions(+), 4 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 24ecc88d2fbc0..538c66f9776ea 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -14336,6 +14336,8 @@ def note_amdgpu_named_barrier_reason_inherited : Note< // AMDGCN builtins diagnostics def err_amdgcn_load_lds_size_invalid_value : Error<"invalid size value">; def note_amdgcn_load_lds_size_valid_value : Note<"size must be %select{1, 2, or 4|1, 2, 4, 12 or 16}0">; +def err_amdgcn_fence_invalid_sync_scope + : Error<"unsupported atomic synchronization scope '%0'">; def err_amdgcn_processor_is_arg_not_literal : Error<"the argument to __builtin_amdgcn_processor_is must be a string " "literal">; diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp index 0bda851b0c1cf..fcf13ab54d78a 100644 --- a/clang/lib/Sema/SemaAMDGPU.cpp +++ b/clang/lib/Sema/SemaAMDGPU.cpp @@ -27,6 +27,8 @@ #include "llvm/Support/AMDGPUAddrSpace.h" #include "llvm/Support/AtomicOrdering.h" #include "llvm/TargetParser/AMDGPUTargetParser.h" +#include "llvm/TargetParser/AtomicScope.h" +#include "llvm/TargetParser/Triple.h" #include <cstdint> #include <utility> @@ -152,6 +154,16 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID, return Diag(ScopeExpr->getExprLoc(), diag::err_expr_not_string_literal) << ScopeExpr->getType(); + // Reject any string that is not a valid synchronization scope name. + std::optional<std::string> ScopeName = + ScopeExpr->tryEvaluateString(getASTContext()); + const llvm::Triple &TT = getASTContext().getTargetInfo().getTriple(); + if (ScopeName && !llvm::parseAtomicScopeIRString(TT, *ScopeName)) { + return Diag(ScopeExpr->getExprLoc(), + diag::err_amdgcn_fence_invalid_sync_scope) + << *ScopeName << ScopeExpr->getSourceRange(); + } + return false; } case AMDGPU::BI__builtin_amdgcn_s_setreg: diff --git a/clang/test/Sema/builtin-amdgcn-fence-failure.cpp b/clang/test/Sema/builtin-amdgcn-fence-failure.cpp index 329134fc5abf5..722b9d32e9cfb 100644 --- a/clang/test/Sema/builtin-amdgcn-fence-failure.cpp +++ b/clang/test/Sema/builtin-amdgcn-fence-failure.cpp @@ -1,8 +1,6 @@ // REQUIRES: amdgpu-registered-target -// RUN: not %clang_cc1 %s -o - -S -triple=amdgpu-amd-amdhsa 2>&1 | FileCheck %s +// RUN: %clang_cc1 %s -fsyntax-only -triple=amdgpu-amd-amdhsa -verify void test_amdgcn_fence_failure() { - - // CHECK: error: Unsupported atomic synchronization scope - __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "foobar"); + __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "foobar"); // expected-error {{unsupported atomic synchronization scope 'foobar'}} } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
