[clang] [Clang][HIP] Deprecate __hip_atomic_* builtins (PR #189897)
Sameer Sahasrabuddhe via cfe-commits
cfe-commits at lists.llvm.org
Thu Apr 2 23:54:55 PDT 2026
https://github.com/ssahasra updated https://github.com/llvm/llvm-project/pull/189897
>From dfeb02bd65a1aca9edb3ca9a06ad2096b9e45c4f Mon Sep 17 00:00:00 2001
From: Sameer Sahasrabuddhe <sameer.sahasrabuddhe at amd.com>
Date: Wed, 1 Apr 2026 07:14:07 +0530
Subject: [PATCH 1/2] [Clang][HIP] Deprecate __hip_atomic_* builtins
Emit -Wdeprecated-builtins warnings for all __hip_atomic_* builtins,
pointing users to their __scoped_atomic_* equivalents.
Provide a fixit when the scope is a compile-time constant and there is a direct
mapping from the HIP builtin to a Clang builtin. The compare_exchange builtins
differ in how they accept the desired value, so only a warning (without a fixit)
is emitted for those.
Assisted-By: Claude Opus 4.6
---
clang/lib/Sema/SemaChecking.cpp | 88 +++++++++++++++++++
clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu | 11 ++-
clang/test/CodeGenCUDA/atomic-ops.cu | 2 +-
.../CodeGenHIP/atomic-deprecated-fixit.hip | 38 ++++++++
clang/test/SemaCUDA/atomic-ops.cu | 2 +-
.../test/SemaCUDA/spirv-amdgcn-atomic-ops.cu | 2 +-
clang/test/SemaHIP/atomic-deprecated.hip | 58 ++++++++++++
7 files changed, 195 insertions(+), 6 deletions(-)
create mode 100644 clang/test/CodeGenHIP/atomic-deprecated-fixit.hip
create mode 100644 clang/test/SemaHIP/atomic-deprecated.hip
diff --git a/clang/lib/Sema/SemaChecking.cpp b/clang/lib/Sema/SemaChecking.cpp
index de8b965144971..5223565dd3e11 100644
--- a/clang/lib/Sema/SemaChecking.cpp
+++ b/clang/lib/Sema/SemaChecking.cpp
@@ -4565,6 +4565,91 @@ ExprResult Sema::AtomicOpsOverloaded(ExprResult TheCallResult,
Op);
}
+/// Deprecate __hip_atomic_* builtins in favour of __scoped_atomic_*
+/// equivalents. Provide a fixit when the scope is a compile-time constant and
+/// there is a direct mapping from the HIP builtin to a Clang builtin. The
+/// compare_exchange builtins differ in how they accept the desired value, so
+/// only a warning (without a fixit) is emitted for those.
+static void DiagnoseDeprecatedHIPAtomic(Sema &S, SourceRange ExprRange,
+ MultiExprArg Args,
+ AtomicExpr::AtomicOp Op) {
+ StringRef OldName;
+ StringRef NewName;
+ bool CanFixIt;
+
+ switch (Op) {
+#define HIP_ATOMIC_FIXABLE(hip, scoped) \
+ case AtomicExpr::AO__hip_atomic_##hip: \
+ OldName = "__hip_atomic_" #hip; \
+ NewName = "__scoped_atomic_" #scoped; \
+ CanFixIt = true; \
+ break;
+ HIP_ATOMIC_FIXABLE(load, load_n)
+ HIP_ATOMIC_FIXABLE(store, store_n)
+ HIP_ATOMIC_FIXABLE(exchange, exchange_n)
+ HIP_ATOMIC_FIXABLE(fetch_add, fetch_add)
+ HIP_ATOMIC_FIXABLE(fetch_sub, fetch_sub)
+ HIP_ATOMIC_FIXABLE(fetch_and, fetch_and)
+ HIP_ATOMIC_FIXABLE(fetch_or, fetch_or)
+ HIP_ATOMIC_FIXABLE(fetch_xor, fetch_xor)
+ HIP_ATOMIC_FIXABLE(fetch_min, fetch_min)
+ HIP_ATOMIC_FIXABLE(fetch_max, fetch_max)
+#undef HIP_ATOMIC_FIXABLE
+ case AtomicExpr::AO__hip_atomic_compare_exchange_weak:
+ OldName = "__hip_atomic_compare_exchange_weak";
+ NewName = "__scoped_atomic_compare_exchange";
+ CanFixIt = false;
+ break;
+ case AtomicExpr::AO__hip_atomic_compare_exchange_strong:
+ OldName = "__hip_atomic_compare_exchange_strong";
+ NewName = "__scoped_atomic_compare_exchange";
+ CanFixIt = false;
+ break;
+ default:
+ llvm_unreachable("unhandled HIP atomic op");
+ }
+
+ auto DB = S.Diag(ExprRange.getBegin(), diag::warn_deprecated_builtin)
+ << OldName << NewName;
+ if (!CanFixIt)
+ return;
+
+ DB << FixItHint::CreateReplacement(ExprRange, NewName);
+
+ Expr *Scope = Args[Args.size() - 1];
+ std::optional<llvm::APSInt> ScopeVal =
+ Scope->getIntegerConstantExpr(S.Context);
+ if (!ScopeVal)
+ return;
+
+ StringRef ScopeName;
+ switch (ScopeVal->getZExtValue()) {
+ case AtomicScopeHIPModel::SingleThread:
+ ScopeName = "__MEMORY_SCOPE_SINGLE";
+ break;
+ case AtomicScopeHIPModel::Wavefront:
+ ScopeName = "__MEMORY_SCOPE_WVFRNT";
+ break;
+ case AtomicScopeHIPModel::Workgroup:
+ ScopeName = "__MEMORY_SCOPE_WRKGRP";
+ break;
+ case AtomicScopeHIPModel::Agent:
+ ScopeName = "__MEMORY_SCOPE_DEVICE";
+ break;
+ case AtomicScopeHIPModel::System:
+ ScopeName = "__MEMORY_SCOPE_SYSTEM";
+ break;
+ case AtomicScopeHIPModel::Cluster:
+ ScopeName = "__MEMORY_SCOPE_CLUSTR";
+ break;
+ default:
+ return;
+ }
+
+ DB << FixItHint::CreateReplacement(
+ CharSourceRange::getTokenRange(Scope->getSourceRange()), ScopeName);
+}
+
ExprResult Sema::BuildAtomicExpr(SourceRange CallRange, SourceRange ExprRange,
SourceLocation RParenLoc, MultiExprArg Args,
AtomicExpr::AtomicOp Op,
@@ -5161,6 +5246,9 @@ ExprResult Sema::BuildAtomicExpr(SourceRange CallRange, SourceRange ExprRange,
SubExprs.push_back(Scope);
}
+ if (IsHIP)
+ DiagnoseDeprecatedHIPAtomic(*this, ExprRange, Args, Op);
+
AtomicExpr *AE = new (Context)
AtomicExpr(ExprRange.getBegin(), SubExprs, ResultType, Op, RParenLoc);
diff --git a/clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu b/clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu
index 22c40e6d38ea2..bf28c62abc29c 100644
--- a/clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu
+++ b/clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu
@@ -1,18 +1,23 @@
// RUN: %clang_cc1 -x hip %s -emit-llvm -o - -triple=amdgcn-amd-amdhsa \
// RUN: -fcuda-is-device -target-cpu gfx906 -fnative-half-type \
-// RUN: -fnative-half-arguments-and-returns | FileCheck -check-prefixes=FUN,CHECK,SAFEIR %s
+// RUN: -fnative-half-arguments-and-returns -Wno-deprecated-builtins \
+// RUN: | FileCheck -check-prefixes=FUN,CHECK,SAFEIR %s
// RUN: %clang_cc1 -x hip %s -emit-llvm -o - -triple=amdgcn-amd-amdhsa \
// RUN: -fcuda-is-device -target-cpu gfx906 -fnative-half-type \
-// RUN: -fnative-half-arguments-and-returns -munsafe-fp-atomics | FileCheck -check-prefixes=FUN,CHECK,UNSAFEIR %s
+// RUN: -fnative-half-arguments-and-returns -munsafe-fp-atomics \
+// RUN: -Wno-deprecated-builtins \
+// RUN: | FileCheck -check-prefixes=FUN,CHECK,UNSAFEIR %s
// RUN: %clang_cc1 -x hip %s -O3 -S -o - -triple=amdgcn-amd-amdhsa \
// RUN: -fcuda-is-device -target-cpu gfx1100 -fnative-half-type \
-// RUN: -fnative-half-arguments-and-returns | FileCheck -check-prefixes=FUN,SAFE %s
+// RUN: -fnative-half-arguments-and-returns -Wno-deprecated-builtins \
+// RUN: | FileCheck -check-prefixes=FUN,SAFE %s
// RUN: %clang_cc1 -x hip %s -O3 -S -o - -triple=amdgcn-amd-amdhsa \
// RUN: -fcuda-is-device -target-cpu gfx942 -fnative-half-type \
// RUN: -fnative-half-arguments-and-returns -munsafe-fp-atomics \
+// RUN: -Wno-deprecated-builtins \
// RUN: | FileCheck -check-prefixes=FUN,UNSAFE %s
// REQUIRES: amdgpu-registered-target
diff --git a/clang/test/CodeGenCUDA/atomic-ops.cu b/clang/test/CodeGenCUDA/atomic-ops.cu
index a41e6a6fb2dc7..0954d43805894 100644
--- a/clang/test/CodeGenCUDA/atomic-ops.cu
+++ b/clang/test/CodeGenCUDA/atomic-ops.cu
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -x hip -std=c++11 -triple amdgcn -fcuda-is-device -emit-llvm %s -o - | FileCheck -enable-var-scope %s
+// RUN: %clang_cc1 -x hip -std=c++11 -triple amdgcn -fcuda-is-device -emit-llvm -Wno-deprecated-builtins %s -o - | FileCheck -enable-var-scope %s
#include "Inputs/cuda.h"
// CHECK-LABEL: @_Z24atomic32_op_singlethreadPiii
diff --git a/clang/test/CodeGenHIP/atomic-deprecated-fixit.hip b/clang/test/CodeGenHIP/atomic-deprecated-fixit.hip
new file mode 100644
index 0000000000000..8291612c0e39d
--- /dev/null
+++ b/clang/test/CodeGenHIP/atomic-deprecated-fixit.hip
@@ -0,0 +1,38 @@
+// Verify that applying the deprecation fix-its for __hip_atomic_* builtins
+// produces semantically identical LLVM IR. Both the original source and the
+// fixed source are compiled and checked against the same CHECK lines.
+//
+// RUN: %clang_cc1 -x hip -std=c++11 -triple amdgcn -fcuda-is-device \
+// RUN: -emit-llvm -Wno-deprecated-builtins %s -o - | FileCheck %s
+// RUN: cp %s %t
+// RUN: %clang_cc1 -x hip -std=c++11 -triple amdgcn -fcuda-is-device \
+// RUN: -Wdeprecated-builtins -fixit %t
+// RUN: %clang_cc1 -x hip -std=c++11 -triple amdgcn -fcuda-is-device \
+// RUN: -emit-llvm %t -o - | FileCheck %s
+
+#define __device__ __attribute__((device))
+
+// CHECK-LABEL: @_Z14test_all_fixedPiii
+// CHECK: load atomic i32, ptr {{%.*}} syncscope("singlethread") monotonic
+// CHECK: store atomic i32 %{{.*}}, ptr %{{.*}} syncscope("wavefront") monotonic
+// CHECK: atomicrmw xchg ptr {{%.*}}, i32 {{%.*}} syncscope("workgroup") monotonic
+// CHECK: atomicrmw add ptr {{%.*}}, i32 {{%.*}} syncscope("cluster") monotonic
+// CHECK: atomicrmw sub ptr {{%.*}}, i32 {{%.*}} syncscope("agent") monotonic
+// CHECK: atomicrmw and ptr {{%.*}}, i32 {{%.*}} monotonic
+// CHECK: atomicrmw or ptr {{%.*}}, i32 {{%.*}} syncscope("singlethread") monotonic
+// CHECK: atomicrmw xor ptr {{%.*}}, i32 {{%.*}} syncscope("wavefront") monotonic
+// CHECK: atomicrmw min ptr {{%.*}}, i32 {{%.*}} syncscope("workgroup") monotonic
+// CHECK: atomicrmw max ptr {{%.*}}, i32 {{%.*}} syncscope("agent") monotonic
+__device__ int test_all_fixed(int *p, int v, int d) {
+ v = __hip_atomic_load(p, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SINGLETHREAD);
+ __hip_atomic_store(p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_WAVEFRONT);
+ v = __hip_atomic_exchange(p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_WORKGROUP);
+ v = __hip_atomic_fetch_add(p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_CLUSTER);
+ v = __hip_atomic_fetch_sub(p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+ v = __hip_atomic_fetch_and(p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
+ v = __hip_atomic_fetch_or(p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SINGLETHREAD);
+ v = __hip_atomic_fetch_xor(p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_WAVEFRONT);
+ v = __hip_atomic_fetch_min(p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_WORKGROUP);
+ v = __hip_atomic_fetch_max(p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
+ return v;
+}
diff --git a/clang/test/SemaCUDA/atomic-ops.cu b/clang/test/SemaCUDA/atomic-ops.cu
index 40e110c4b9b77..9d78dd3ebcd48 100644
--- a/clang/test/SemaCUDA/atomic-ops.cu
+++ b/clang/test/SemaCUDA/atomic-ops.cu
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -x hip -std=c++11 -triple amdgcn -fcuda-is-device -verify -fsyntax-only %s
+// RUN: %clang_cc1 -x hip -std=c++11 -triple amdgcn -fcuda-is-device -verify -fsyntax-only -Wno-deprecated-builtins %s
#include "Inputs/cuda.h"
diff --git a/clang/test/SemaCUDA/spirv-amdgcn-atomic-ops.cu b/clang/test/SemaCUDA/spirv-amdgcn-atomic-ops.cu
index 503e786877819..adddb140857ff 100644
--- a/clang/test/SemaCUDA/spirv-amdgcn-atomic-ops.cu
+++ b/clang/test/SemaCUDA/spirv-amdgcn-atomic-ops.cu
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -x hip -std=c++11 -triple spirv64-amd-amdhsa -fcuda-is-device -verify -fsyntax-only %s
+// RUN: %clang_cc1 -x hip -std=c++11 -triple spirv64-amd-amdhsa -fcuda-is-device -verify -fsyntax-only -Wno-deprecated-builtins %s
#include "Inputs/cuda.h"
diff --git a/clang/test/SemaHIP/atomic-deprecated.hip b/clang/test/SemaHIP/atomic-deprecated.hip
new file mode 100644
index 0000000000000..4cff9f99bf91b
--- /dev/null
+++ b/clang/test/SemaHIP/atomic-deprecated.hip
@@ -0,0 +1,58 @@
+// RUN: %clang_cc1 -x hip -triple amdgcn -fcuda-is-device -verify -fsyntax-only -Wdeprecated-builtins %s
+// RUN: %clang_cc1 -x hip -triple amdgcn -fcuda-is-device -fdiagnostics-parseable-fixits -Wdeprecated-builtins %s 2>&1 | FileCheck %s
+
+#define __device__ __attribute__((device))
+
+__device__ void test(int *p, int v, int s) {
+ __hip_atomic_load(p, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM); // expected-warning {{builtin __hip_atomic_load is deprecated; use __scoped_atomic_load_n instead}}
+ __hip_atomic_store(p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM); // expected-warning {{builtin __hip_atomic_store is deprecated; use __scoped_atomic_store_n instead}}
+ __hip_atomic_load(p, __ATOMIC_RELAXED, s); // expected-warning {{builtin __hip_atomic_load is deprecated; use __scoped_atomic_load_n instead}}
+ __hip_atomic_store(p, v, __ATOMIC_RELAXED, s); // expected-warning {{builtin __hip_atomic_store is deprecated; use __scoped_atomic_store_n instead}}
+}
+
+// Statically known scope: both the builtin name and the scope argument get fix-its.
+// CHECK: fix-it:{{.*}}:{[[@LINE-7]]:3-[[@LINE-7]]:20}:"__scoped_atomic_load_n"
+// CHECK: fix-it:{{.*}}:{[[@LINE-8]]:42-[[@LINE-8]]:67}:"__MEMORY_SCOPE_SYSTEM"
+// CHECK: fix-it:{{.*}}:{[[@LINE-8]]:3-[[@LINE-8]]:21}:"__scoped_atomic_store_n"
+// CHECK: fix-it:{{.*}}:{[[@LINE-9]]:46-[[@LINE-9]]:71}:"__MEMORY_SCOPE_SYSTEM"
+// Dynamic scope: only the builtin name gets a fix-it.
+// CHECK: fix-it:{{.*}}:{[[@LINE-10]]:3-[[@LINE-10]]:20}:"__scoped_atomic_load_n"
+// CHECK: fix-it:{{.*}}:{[[@LINE-10]]:3-[[@LINE-10]]:21}:"__scoped_atomic_store_n"
+
+// Fetch ops: argument layout matches __scoped_atomic_fetch_*. Scope varies
+// across all six __HIP_MEMORY_SCOPE_* values.
+__device__ void test_fetch(int *p, int v) {
+ __hip_atomic_fetch_add(p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SINGLETHREAD); // expected-warning {{builtin __hip_atomic_fetch_add is deprecated; use __scoped_atomic_fetch_add instead}}
+ __hip_atomic_fetch_sub(p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_WAVEFRONT); // expected-warning {{builtin __hip_atomic_fetch_sub is deprecated; use __scoped_atomic_fetch_sub instead}}
+ __hip_atomic_fetch_and(p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_WORKGROUP); // expected-warning {{builtin __hip_atomic_fetch_and is deprecated; use __scoped_atomic_fetch_and instead}}
+ __hip_atomic_fetch_or(p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); // expected-warning {{builtin __hip_atomic_fetch_or is deprecated; use __scoped_atomic_fetch_or instead}}
+ __hip_atomic_fetch_xor(p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM); // expected-warning {{builtin __hip_atomic_fetch_xor is deprecated; use __scoped_atomic_fetch_xor instead}}
+ __hip_atomic_fetch_min(p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_CLUSTER); // expected-warning {{builtin __hip_atomic_fetch_min is deprecated; use __scoped_atomic_fetch_min instead}}
+ __hip_atomic_fetch_max(p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SINGLETHREAD); // expected-warning {{builtin __hip_atomic_fetch_max is deprecated; use __scoped_atomic_fetch_max instead}}
+}
+// CHECK: fix-it:{{.*}}:{[[@LINE-8]]:3-[[@LINE-8]]:25}:"__scoped_atomic_fetch_add"
+// CHECK: fix-it:{{.*}}:{[[@LINE-9]]:50-[[@LINE-9]]:81}:"__MEMORY_SCOPE_SINGLE"
+// CHECK: fix-it:{{.*}}:{[[@LINE-9]]:3-[[@LINE-9]]:25}:"__scoped_atomic_fetch_sub"
+// CHECK: fix-it:{{.*}}:{[[@LINE-10]]:50-[[@LINE-10]]:78}:"__MEMORY_SCOPE_WVFRNT"
+// CHECK: fix-it:{{.*}}:{[[@LINE-10]]:3-[[@LINE-10]]:25}:"__scoped_atomic_fetch_and"
+// CHECK: fix-it:{{.*}}:{[[@LINE-11]]:50-[[@LINE-11]]:78}:"__MEMORY_SCOPE_WRKGRP"
+// CHECK: fix-it:{{.*}}:{[[@LINE-11]]:3-[[@LINE-11]]:24}:"__scoped_atomic_fetch_or"
+// CHECK: fix-it:{{.*}}:{[[@LINE-12]]:49-[[@LINE-12]]:73}:"__MEMORY_SCOPE_DEVICE"
+// CHECK: fix-it:{{.*}}:{[[@LINE-12]]:3-[[@LINE-12]]:25}:"__scoped_atomic_fetch_xor"
+// CHECK: fix-it:{{.*}}:{[[@LINE-13]]:50-[[@LINE-13]]:75}:"__MEMORY_SCOPE_SYSTEM"
+// CHECK: fix-it:{{.*}}:{[[@LINE-13]]:3-[[@LINE-13]]:25}:"__scoped_atomic_fetch_min"
+// CHECK: fix-it:{{.*}}:{[[@LINE-14]]:50-[[@LINE-14]]:76}:"__MEMORY_SCOPE_CLUSTR"
+// CHECK: fix-it:{{.*}}:{[[@LINE-14]]:3-[[@LINE-14]]:25}:"__scoped_atomic_fetch_max"
+// CHECK: fix-it:{{.*}}:{[[@LINE-15]]:50-[[@LINE-15]]:81}:"__MEMORY_SCOPE_SINGLE"
+
+// Exchange op: argument layout matches __scoped_atomic_exchange_n, so both
+// name and scope get fix-its. compare_exchange_{weak,strong} have a different
+// argument layout (scoped form takes an additional 'weak' boolean and passes
+// 'desired' by pointer), so only a deprecation warning is emitted.
+__device__ void test_exchange(int *p, int *ep, int v) {
+ __hip_atomic_exchange(p, v, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_WAVEFRONT); // expected-warning {{builtin __hip_atomic_exchange is deprecated; use __scoped_atomic_exchange_n instead}}
+ __hip_atomic_compare_exchange_weak(p, ep, v, __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_WORKGROUP); // expected-warning {{builtin __hip_atomic_compare_exchange_weak is deprecated; use __scoped_atomic_compare_exchange instead}}
+ __hip_atomic_compare_exchange_strong(p, ep, v, __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT); // expected-warning {{builtin __hip_atomic_compare_exchange_strong is deprecated; use __scoped_atomic_compare_exchange instead}}
+}
+// CHECK: fix-it:{{.*}}:{[[@LINE-4]]:3-[[@LINE-4]]:24}:"__scoped_atomic_exchange_n"
+// CHECK: fix-it:{{.*}}:{[[@LINE-5]]:49-[[@LINE-5]]:77}:"__MEMORY_SCOPE_WVFRNT"
>From bc40a119682621a4c743d4d025e06ed69211fc26 Mon Sep 17 00:00:00 2001
From: Sameer Sahasrabuddhe <sameer.sahasrabuddhe at amd.com>
Date: Fri, 3 Apr 2026 11:31:30 +0530
Subject: [PATCH 2/2] warning default ignored, new group
-Whip-deprecated-builtins
---
clang/include/clang/Basic/DiagnosticGroups.td | 1 +
clang/include/clang/Basic/DiagnosticSemaKinds.td | 3 +++
clang/lib/Sema/SemaChecking.cpp | 2 +-
clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu | 11 +++--------
clang/test/CodeGenCUDA/atomic-ops.cu | 2 +-
clang/test/CodeGenHIP/atomic-deprecated-fixit.hip | 4 ++--
clang/test/SemaCUDA/atomic-ops.cu | 2 +-
clang/test/SemaCUDA/spirv-amdgcn-atomic-ops.cu | 2 +-
clang/test/SemaHIP/atomic-deprecated.hip | 4 ++--
9 files changed, 15 insertions(+), 16 deletions(-)
diff --git a/clang/include/clang/Basic/DiagnosticGroups.td b/clang/include/clang/Basic/DiagnosticGroups.td
index 6e82203d858f5..7614e40ac8f46 100644
--- a/clang/include/clang/Basic/DiagnosticGroups.td
+++ b/clang/include/clang/Basic/DiagnosticGroups.td
@@ -270,6 +270,7 @@ def UnguardedAvailability : DiagGroup<"unguarded-availability",
def : DiagGroup<"partial-availability", [UnguardedAvailability]>;
def DeprecatedDynamicExceptionSpec
: DiagGroup<"deprecated-dynamic-exception-spec">;
+def HipDeprecatedBuiltins : DiagGroup<"hip-deprecated-builtins">;
def DeprecatedBuiltins : DiagGroup<"deprecated-builtins">;
def DeprecatedImplementations :DiagGroup<"deprecated-implementations">;
def DeprecatedIncrementBool : DiagGroup<"deprecated-increment-bool">;
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 6b9fa4a257397..5f0ecbac2abe0 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -6330,6 +6330,9 @@ def warn_unavailable_def : Warning<
def warn_deprecated_builtin : Warning<
"builtin %0 is deprecated; use %1 instead">,
InGroup<DeprecatedBuiltins>;
+def warn_hip_deprecated_builtin : Warning<
+ "builtin %0 is deprecated; use %1 instead">,
+ InGroup<HipDeprecatedBuiltins>, DefaultIgnore;
def warn_deprecated_builtin_no_suggestion : Warning<"builtin %0 is deprecated">,
InGroup<DeprecatedBuiltins>;
def err_unavailable : Error<"%0 is unavailable">;
diff --git a/clang/lib/Sema/SemaChecking.cpp b/clang/lib/Sema/SemaChecking.cpp
index 5223565dd3e11..21858c4077bf8 100644
--- a/clang/lib/Sema/SemaChecking.cpp
+++ b/clang/lib/Sema/SemaChecking.cpp
@@ -4609,7 +4609,7 @@ static void DiagnoseDeprecatedHIPAtomic(Sema &S, SourceRange ExprRange,
llvm_unreachable("unhandled HIP atomic op");
}
- auto DB = S.Diag(ExprRange.getBegin(), diag::warn_deprecated_builtin)
+ auto DB = S.Diag(ExprRange.getBegin(), diag::warn_hip_deprecated_builtin)
<< OldName << NewName;
if (!CanFixIt)
return;
diff --git a/clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu b/clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu
index bf28c62abc29c..22c40e6d38ea2 100644
--- a/clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu
+++ b/clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu
@@ -1,23 +1,18 @@
// RUN: %clang_cc1 -x hip %s -emit-llvm -o - -triple=amdgcn-amd-amdhsa \
// RUN: -fcuda-is-device -target-cpu gfx906 -fnative-half-type \
-// RUN: -fnative-half-arguments-and-returns -Wno-deprecated-builtins \
-// RUN: | FileCheck -check-prefixes=FUN,CHECK,SAFEIR %s
+// RUN: -fnative-half-arguments-and-returns | FileCheck -check-prefixes=FUN,CHECK,SAFEIR %s
// RUN: %clang_cc1 -x hip %s -emit-llvm -o - -triple=amdgcn-amd-amdhsa \
// RUN: -fcuda-is-device -target-cpu gfx906 -fnative-half-type \
-// RUN: -fnative-half-arguments-and-returns -munsafe-fp-atomics \
-// RUN: -Wno-deprecated-builtins \
-// RUN: | FileCheck -check-prefixes=FUN,CHECK,UNSAFEIR %s
+// RUN: -fnative-half-arguments-and-returns -munsafe-fp-atomics | FileCheck -check-prefixes=FUN,CHECK,UNSAFEIR %s
// RUN: %clang_cc1 -x hip %s -O3 -S -o - -triple=amdgcn-amd-amdhsa \
// RUN: -fcuda-is-device -target-cpu gfx1100 -fnative-half-type \
-// RUN: -fnative-half-arguments-and-returns -Wno-deprecated-builtins \
-// RUN: | FileCheck -check-prefixes=FUN,SAFE %s
+// RUN: -fnative-half-arguments-and-returns | FileCheck -check-prefixes=FUN,SAFE %s
// RUN: %clang_cc1 -x hip %s -O3 -S -o - -triple=amdgcn-amd-amdhsa \
// RUN: -fcuda-is-device -target-cpu gfx942 -fnative-half-type \
// RUN: -fnative-half-arguments-and-returns -munsafe-fp-atomics \
-// RUN: -Wno-deprecated-builtins \
// RUN: | FileCheck -check-prefixes=FUN,UNSAFE %s
// REQUIRES: amdgpu-registered-target
diff --git a/clang/test/CodeGenCUDA/atomic-ops.cu b/clang/test/CodeGenCUDA/atomic-ops.cu
index 0954d43805894..a41e6a6fb2dc7 100644
--- a/clang/test/CodeGenCUDA/atomic-ops.cu
+++ b/clang/test/CodeGenCUDA/atomic-ops.cu
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -x hip -std=c++11 -triple amdgcn -fcuda-is-device -emit-llvm -Wno-deprecated-builtins %s -o - | FileCheck -enable-var-scope %s
+// RUN: %clang_cc1 -x hip -std=c++11 -triple amdgcn -fcuda-is-device -emit-llvm %s -o - | FileCheck -enable-var-scope %s
#include "Inputs/cuda.h"
// CHECK-LABEL: @_Z24atomic32_op_singlethreadPiii
diff --git a/clang/test/CodeGenHIP/atomic-deprecated-fixit.hip b/clang/test/CodeGenHIP/atomic-deprecated-fixit.hip
index 8291612c0e39d..c26c405ea3f3e 100644
--- a/clang/test/CodeGenHIP/atomic-deprecated-fixit.hip
+++ b/clang/test/CodeGenHIP/atomic-deprecated-fixit.hip
@@ -3,10 +3,10 @@
// fixed source are compiled and checked against the same CHECK lines.
//
// RUN: %clang_cc1 -x hip -std=c++11 -triple amdgcn -fcuda-is-device \
-// RUN: -emit-llvm -Wno-deprecated-builtins %s -o - | FileCheck %s
+// RUN: -emit-llvm %s -o - | FileCheck %s
// RUN: cp %s %t
// RUN: %clang_cc1 -x hip -std=c++11 -triple amdgcn -fcuda-is-device \
-// RUN: -Wdeprecated-builtins -fixit %t
+// RUN: -fixit %t
// RUN: %clang_cc1 -x hip -std=c++11 -triple amdgcn -fcuda-is-device \
// RUN: -emit-llvm %t -o - | FileCheck %s
diff --git a/clang/test/SemaCUDA/atomic-ops.cu b/clang/test/SemaCUDA/atomic-ops.cu
index 9d78dd3ebcd48..40e110c4b9b77 100644
--- a/clang/test/SemaCUDA/atomic-ops.cu
+++ b/clang/test/SemaCUDA/atomic-ops.cu
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -x hip -std=c++11 -triple amdgcn -fcuda-is-device -verify -fsyntax-only -Wno-deprecated-builtins %s
+// RUN: %clang_cc1 -x hip -std=c++11 -triple amdgcn -fcuda-is-device -verify -fsyntax-only %s
#include "Inputs/cuda.h"
diff --git a/clang/test/SemaCUDA/spirv-amdgcn-atomic-ops.cu b/clang/test/SemaCUDA/spirv-amdgcn-atomic-ops.cu
index adddb140857ff..503e786877819 100644
--- a/clang/test/SemaCUDA/spirv-amdgcn-atomic-ops.cu
+++ b/clang/test/SemaCUDA/spirv-amdgcn-atomic-ops.cu
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -x hip -std=c++11 -triple spirv64-amd-amdhsa -fcuda-is-device -verify -fsyntax-only -Wno-deprecated-builtins %s
+// RUN: %clang_cc1 -x hip -std=c++11 -triple spirv64-amd-amdhsa -fcuda-is-device -verify -fsyntax-only %s
#include "Inputs/cuda.h"
diff --git a/clang/test/SemaHIP/atomic-deprecated.hip b/clang/test/SemaHIP/atomic-deprecated.hip
index 4cff9f99bf91b..fcdaa1d8ac710 100644
--- a/clang/test/SemaHIP/atomic-deprecated.hip
+++ b/clang/test/SemaHIP/atomic-deprecated.hip
@@ -1,5 +1,5 @@
-// RUN: %clang_cc1 -x hip -triple amdgcn -fcuda-is-device -verify -fsyntax-only -Wdeprecated-builtins %s
-// RUN: %clang_cc1 -x hip -triple amdgcn -fcuda-is-device -fdiagnostics-parseable-fixits -Wdeprecated-builtins %s 2>&1 | FileCheck %s
+// RUN: %clang_cc1 -x hip -triple amdgcn -fcuda-is-device -verify -fsyntax-only -Whip-deprecated-builtins %s
+// RUN: %clang_cc1 -x hip -triple amdgcn -fcuda-is-device -fdiagnostics-parseable-fixits -Whip-deprecated-builtins %s 2>&1 | FileCheck %s
#define __device__ __attribute__((device))
More information about the cfe-commits
mailing list