diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 113f77661d97296c2c188d8d7f51edf24037b794..46d0a66d59c3753bb8ea41fa494f9d99aabd9422 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 5ec7fe42e09b0bb7e20d8607e5bbb7b34918997c..1064507f34616a5ba9e33588f7d0b184fd8bcde7 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 9ce8c13bd95848366c58e5d1d4f5e61526af4122..c3251bb5ab5657978203191fa1989b0464cc383a 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 d41fee8e34985e956966dde84994fb5c11a64c71..e2eada24f9fccbfa1e67bc52d617b4bf85c38248 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 95dddd477998370975c57a40ca849a69041809f1..7763dae3c9aa7abb7a1592b8dbb0bed28b21a8b9 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);