From 3ccdc8773d08ce9257600f0c1ad114ab067b9fba Mon Sep 17 00:00:00 2001
From: Blaise Tine <tinebp@yahoo.com>
Date: Sun, 19 May 2024 15:13:21 -0700
Subject: [PATCH] TODO
---
clang/include/clang/Basic/Attr.td | 8 --------
clang/lib/AST/ASTContext.cpp | 2 --
clang/lib/CodeGen/CGDecl.cpp | 16 ----------------
clang/lib/Sema/SemaDeclAttr.cpp | 7 -------
llvm/lib/Target/RISCV/VortexBranchDivergence.cpp | 4 ++--
5 files changed, 2 insertions(+), 35 deletions(-)
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 113f77661d97..46d0a66d59c3 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -2050,14 +2050,6 @@ def Convergent : InheritableAttr {
let SimpleHandler = 1;
}
-// Vortex attributes for marking variables as uniform
-def VortexUniform : InheritableAttr {
- let Spellings = [GNU<"uniform">, Declspec<"__uniform__">];
- let Subjects = SubjectList<[Var]>;
- let LangOpts = [CPlusPlus];
- let Documentation = [Undocumented];
-}
-
def NoInline : DeclOrStmtAttr {
let Spellings = [CustomKeyword<"__noinline__">, GCC<"noinline">,
CXX11<"clang", "noinline">, C23<"clang", "noinline">,
diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp
index 5ec7fe42e09b..1064507f3461 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/clang/lib/AST/ASTContext.cpp
@@ -12316,8 +12316,6 @@ bool ASTContext::DeclMustBeEmitted(const Decl *D) {
if (VD->getDescribedVarTemplate() ||
isa<VarTemplatePartialSpecializationDecl>(VD))
return false;
- if (VD->hasAttr<VortexUniformAttr>())
- return false;
} else if (const auto *FD = dyn_cast<FunctionDecl>(D)) {
// We never need to emit an uninstantiated function template.
if (FD->getTemplatedKind() == FunctionDecl::TK_FunctionTemplate)
diff --git a/clang/lib/CodeGen/CGDecl.cpp b/clang/lib/CodeGen/CGDecl.cpp
index 9ce8c13bd958..c3251bb5ab56 100644
--- a/clang/lib/CodeGen/CGDecl.cpp
+++ b/clang/lib/CodeGen/CGDecl.cpp
@@ -1684,14 +1684,6 @@ CodeGenFunction::EmitAutoVarAlloca(const VarDecl &D) {
if (D.hasAttr<AnnotateAttr>() && HaveInsertPoint())
EmitVarAnnotations(&D, address.emitRawPointer(*this));
- if (D.hasAttr<VortexUniformAttr>() && HaveInsertPoint()) {
- if (auto I = dyn_cast<llvm::Instruction>(address.getPointer())) {
- auto &Context = I->getContext();
- auto MD = llvm::MDNode::get(Context, llvm::MDString::get(Context, "Uniform Variable"));
- I->setMetadata("vortex.uniform", MD);
- }
- }
-
// Make sure we call @llvm.lifetime.end.
if (emission.useLifetimeMarkers())
EHStack.pushCleanup<CallLifetimeEnd>(NormalEHLifetimeMarker,
@@ -2764,14 +2756,6 @@ void CodeGenFunction::EmitParmDecl(const VarDecl &D, ParamValue Arg,
if (D.hasAttr<AnnotateAttr>())
EmitVarAnnotations(&D, DeclPtr.emitRawPointer(*this));
- if (D.hasAttr<VortexUniformAttr>()) {
- if (auto I = dyn_cast<llvm::Instruction>(DeclPtr.getPointer())) {
- auto &Context = I->getContext();
- auto MD = llvm::MDNode::get(Context, llvm::MDString::get(Context, "Uniform Variable"));
- I->setMetadata("vortex.uniform", MD);
- }
- }
-
// We can only check return value nullability if all arguments to the
// function satisfy their nullability preconditions. This makes it necessary
// to emit null checks for args in the function body itself.
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index d41fee8e3498..e2eada24f9fc 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -4563,10 +4563,6 @@ static void handleConstantAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
D->addAttr(::new (S.Context) CUDAConstantAttr(S.Context, AL));
}
-static void handleVortexUniformAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
- handleSimpleAttribute<VortexUniformAttr>(S, D, AL);
-}
-
static void handleSharedAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
const auto *VD = cast<VarDecl>(D);
// extern __shared__ is only allowed on arrays with no length (e.g.
@@ -6473,9 +6469,6 @@ ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, const ParsedAttr &AL,
case ParsedAttr::AT_CUDAConstant:
handleConstantAttr(S, D, AL);
break;
- case ParsedAttr::AT_VortexUniform:
- handleVortexUniformAttr(S, D, AL);
- break;
case ParsedAttr::AT_PassObjectSize:
handlePassObjectSizeAttr(S, D, AL);
break;
diff --git a/llvm/lib/Target/RISCV/VortexBranchDivergence.cpp b/llvm/lib/Target/RISCV/VortexBranchDivergence.cpp
index 95dddd477998..7763dae3c9aa 100644
--- a/llvm/lib/Target/RISCV/VortexBranchDivergence.cpp
+++ b/llvm/lib/Target/RISCV/VortexBranchDivergence.cpp
@@ -984,11 +984,11 @@ void DivergenceTracker::initialize() {
auto var = II->getOperand(0);
if (auto AI = dyn_cast<AllocaInst>(var)) {
var_src = AI;
- LLVM_DEBUG(dbgs() << "*** uniform annotation: " << AI->getName() << ".src(" << var_src << "\n");
+ LLVM_DEBUG(dbgs() << "*** divergent annotation: " << AI->getName() << ".src(" << var_src << "\n");
} else
if (auto CI = dyn_cast<CastInst>(var)) {
var_src = CI->getOperand(0);
- LLVM_DEBUG(dbgs() << "*** uniform annotation: " << CI->getName() << ".src(" << var_src << "\n");
+ LLVM_DEBUG(dbgs() << "*** divergent annotation: " << CI->getName() << ".src(" << var_src << "\n");
}
dv_annotations.insert(var_src);
dv_nodes_.insert(var_src);
--
GitLab