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