Remove CA references under compiler_pipeline and vecz

Colin Davidson · Colin Davidson · commit 758d573e40b8 · 2025-08-19T13:44:50.000+01:00
There are multiple references which are not publicly visible. Remove
the ones under compiler_pipeline and vecz as part of this process.
diff --git a/modules/compiler/compiler_pipeline/include/compiler/utils/device_info.h b/modules/compiler/compiler_pipeline/include/compiler/utils/device_info.h
@@ -35,7 +35,7 @@ namespace utils {
 ///
 /// NOTE: Must be kept in sync with mux_floating_point_capabilities_e in
 /// mux/include/mux/mux.h! This should probably be placed in an intermediary
-/// mux/compiler library and shared as part of CA-4236.
+/// mux/compiler library and shared.
 enum device_floating_point_capabilities_e {
   /// @brief Denormals supported.
   device_floating_point_capabilities_denorm = 0x1,
diff --git a/modules/compiler/compiler_pipeline/source/builtin_info.cpp b/modules/compiler/compiler_pipeline/source/builtin_info.cpp
@@ -314,7 +314,7 @@ BuiltinUniformity BuiltinInfo::isBuiltinUniform(const Builtin &B,
     case eMuxBuiltinGetLocalLinearId:
     case eMuxBuiltinGetGlobalLinearId:
       // TODO: This is fine for vectorizing in the x-axis, but currently we do
-      // not support vectorizing along y or z (see CA-2843).
+      // not support vectorizing along y or z.
       return SimdDimIdx ? eBuiltinUniformityNever
                         : eBuiltinUniformityInstanceID;
   }
diff --git a/modules/compiler/compiler_pipeline/source/cl_builtin_info.cpp b/modules/compiler/compiler_pipeline/source/cl_builtin_info.cpp
@@ -1533,7 +1533,7 @@ Value *CLBuiltinInfo::emitBuiltinInline(Function *F, IRBuilder<> &B,
         NameMangler Mangler(&F->getContext());
         const auto name = Mangler.demangleName(F->getName());
         if (name == "vload_half") {
-          // TODO CA-4691 handle "vload_halfn"
+          // TODO handle "vload_halfn"
           return emitBuiltinInlineVLoadHalf(F, B, Args);
         }
       } break;
@@ -1551,7 +1551,7 @@ Value *CLBuiltinInfo::emitBuiltinInline(Function *F, IRBuilder<> &B,
         NameMangler Mangler(&F->getContext());
         Lexer L(Mangler.demangleName(F->getName()));
         if (L.Consume("vstore_half")) {
-          // TODO CA-4691 handle "vstore_halfn"
+          // TODO handle "vstore_halfn"
           return emitBuiltinInlineVStoreHalf(F, L.TextLeft(), B, Args);
         }
       } break;
@@ -1969,7 +1969,7 @@ Value *CLBuiltinInfo::emitBuiltinInlineAsLLVMBinaryIntrinsic(
   const Triple TT(B.GetInsertBlock()->getModule()->getTargetTriple());
   if (TT.getArch() == Triple::arm || TT.getArch() == Triple::aarch64) {
     // fmin and fmax fail CTS on arm targets.
-    // This is a HACK and should be removed when CA-3595 is resolved.
+    // This is a HACK and should be removed when it is resolved.
     return nullptr;
   }
 
diff --git a/modules/compiler/compiler_pipeline/source/mux_builtin_info.cpp b/modules/compiler/compiler_pipeline/source/mux_builtin_info.cpp
@@ -531,7 +531,6 @@ Function *BIMuxInfoConcept::defineMemBarrier(Function &F, unsigned,
   // our set of default set of targets can't make use of anything but a
   // single-threaded fence. We're also ignoring the kind of memory being
   // controlled by the barrier.
-  // See CA-2997 and CA-3042 for related discussions.
   auto &M = *F.getParent();
   setDefaultBuiltinAttributes(F);
   F.setLinkage(GlobalValue::InternalLinkage);
diff --git a/modules/compiler/compiler_pipeline/source/replace_local_module_scope_variables_pass.cpp b/modules/compiler/compiler_pipeline/source/replace_local_module_scope_variables_pass.cpp
@@ -696,7 +696,6 @@ PreservedAnalyses compiler::utils::ReplaceLocalModuleScopeVariablesPass::run(
 
       // We can't guarantee a subprogram for all functions.
       // FIXME: Should we be able to? Do we need to clone subprograms somehow?
-      // See CA-4241.
       if (!DISubprogram) {
         continue;
       }
diff --git a/modules/compiler/compiler_pipeline/source/replace_mem_intrinsics_pass.cpp b/modules/compiler/compiler_pipeline/source/replace_mem_intrinsics_pass.cpp
@@ -40,7 +40,7 @@ PreservedAnalyses ReplaceMemIntrinsicsPass::run(Function &F,
             break;
           case Intrinsic::memmove:
             // The expandMemMoveAsLoop fails if the address spaces differ so we
-            // will only do it if they are the same - see CA-4682
+            // will only do it if they are the same
             if (CI->getArgOperand(0)->getType()->getPointerAddressSpace() ==
                 CI->getArgOperand(1)->getType()->getPointerAddressSpace()) {
               CallsToProcess.push_back(CI);
diff --git a/modules/compiler/compiler_pipeline/source/unique_opaque_structs_pass.cpp b/modules/compiler/compiler_pipeline/source/unique_opaque_structs_pass.cpp
@@ -74,7 +74,7 @@ static bool shouldClone(compiler::utils::StructTypeRemapper &StructTypeRemapper,
     }
   }
 
-  // TODO: Check globals (see CA-3833).
+  // TODO: Check globals.
 
   // If an instruction makes use of a type but
   // isn't of that type e.g. a cast it will necessarily get caught by
diff --git a/modules/compiler/compiler_pipeline/source/work_item_loops_pass.cpp b/modules/compiler/compiler_pipeline/source/work_item_loops_pass.cpp
@@ -1921,7 +1921,7 @@ PreservedAnalyses compiler::utils::WorkItemLoopsPass::run(
       // FIXME: Also mark them as internal in case they contain symbols we
       // haven't resolved as part of the work-item loop wrapping process. We
       // rely on GlobalOptPass to remove such functions; this is the same root
-      // issue as CA-4126.
+      // issue as some mux builtins require DCE for correctness.
       F.setLinkage(GlobalValue::InternalLinkage);
     }
   }
diff --git a/modules/compiler/vecz/source/control_flow_boscc.cpp b/modules/compiler/vecz/source/control_flow_boscc.cpp
@@ -293,7 +293,6 @@ bool ControlFlowConversionState::BOSCCGadget::createUniformRegions(
     // take several elements into account:
     // - The length of the duplicated code
     // - branch probability
-    // - TODO: CA-1221
     // size_t cost =
     //    std::accumulate(Region->predicatedBlocks.begin(),
     //    Region->predicatedBlocks.end(), 0,
@@ -352,11 +351,11 @@ bool ControlFlowConversionState::BOSCCGadget::createUniformRegions(
   // doesn't seem to matter, as long as we can fully identify the predicated
   // subset of the SESE region, so we are really working with Multiple-Entry,
   // Single-Exit regions here. This was the cause of the BOSCC Back Door bug
-  // that was encountered previously (CA-2711), where the entry block of a
+  // that was encountered previously, where the entry block of a
   // supposed SESE region did not actually dominate everything in the region,
   // which in this case was caused by an additional non-divergent code path
   // (the "back door" entry point), but it is equally possible for two
-  // divergence-causing branches to enter a predicated region (CA-3194).
+  // divergence-causing branches to enter a predicated region.
   //
   // a)    A*      b)    A       c)    A       d)    A      .
   //      / \           / \           / \           / \     .
@@ -383,7 +382,7 @@ bool ControlFlowConversionState::BOSCCGadget::createUniformRegions(
   // immediate post-dominator of B, the first-encountered divergence causing
   // block. Therefore the two overlapping regions have different exit blocks.
   //
-  // Another situation can arise (CA-3851) where the SESE region can contain
+  // Another situation can arise where the SESE region can contain
   // two completely unconnected predicated subregions. Although the DCBI is
   // SESE compact, a SESE region can still contain other, nested SESE regions.
   // Since an entry point into the predicated subregion is not necessarily the
@@ -648,7 +647,7 @@ bool ControlFlowConversionState::BOSCCGadget::connectBOSCCRegions() {
   VECZ_FAIL_IF(!computeBlockOrdering());
 
   // NOTE doing the Liveness Analysis here is potentially dangerous, since we
-  // have yet to fully restore SSA form (CA-3703).
+  // have yet to fully restore SSA form.
   liveness = &AM.getResult<LivenessAnalysis>(F);
   RC->recalculate(F);
   VECZ_FAIL_IF(!blendFinalize());
diff --git a/modules/compiler/vecz/source/pass.cpp b/modules/compiler/vecz/source/pass.cpp
@@ -112,7 +112,8 @@ PreservedAnalyses RunVeczPass::run(Module &M, ModuleAnalysisManager &MAM) {
   } else {
     if (auto Err = Mach.getPB().parsePassPipeline(PM, VeczPassPipeline)) {
       // NOTE this is a command line user error print, not a debug print.
-      // We may want to hoist this out of Vecz once CA-4134 is resolved.
+      // We may want to hoist this out of Vecz once replacing RunVeczPass with
+      // a passbuilder is resolved.
       errs() << "vecz pipeline: " << toString(std::move(Err)) << "\n";
       return PreservedAnalyses::all();
     }
diff --git a/modules/compiler/vecz/source/transform/basic_mem2reg_pass.cpp b/modules/compiler/vecz/source/transform/basic_mem2reg_pass.cpp
@@ -119,7 +119,7 @@ bool BasicMem2RegPass::canPromoteAlloca(AllocaInst *Alloca) const {
       //   %v = load i16, ptr %a
       // We can only promote the alloca if we can bitcast between the two
       // underlying types as well.
-      // We could probably zero-extend or trunc if we had to? See CA-4382.
+      // We could probably zero-extend or trunc if we had to?
       const unsigned DstPointeeBits = U->getType()->getPrimitiveSizeInBits();
       if (!DstPointeeBits || SrcPointeeBits != DstPointeeBits) {
         return false;
@@ -213,8 +213,7 @@ bool BasicMem2RegPass::promoteAlloca(AllocaInst *Alloca) const {
       //   %a = alloca i32
       //   store i16, ptr %a
       //   %v = load i32, ptr %a
-      // Note: we could do other things if the type sizes didn't match. See
-      // CA-4382.
+      // Note: we could do other things if the type sizes didn't match.
       if (Load->getType()->getPrimitiveSizeInBits() !=
           NewValue->getType()->getPrimitiveSizeInBits()) {
         return false;
diff --git a/modules/compiler/vecz/source/transform/control_flow_conversion_pass.cpp b/modules/compiler/vecz/source/transform/control_flow_conversion_pass.cpp
@@ -1333,7 +1333,7 @@ bool ControlFlowConversionState::Impl::applyMaskToCall(CallInst *CI,
   if (!callee) {
     callee = dyn_cast<Function>(CI->getCalledOperand()->stripPointerCasts());
   }
-  VECZ_FAIL_IF(!callee);  // TODO: CA-1505: Support indirect function calls.
+  VECZ_FAIL_IF(!callee);  // TODO: Support indirect function calls.
   // Check to see if this is a function that we know we won't be able to
   // handle in any other way.
   VECZ_FAIL_IF(callee->cannotDuplicate());
diff --git a/modules/compiler/vecz/source/transform/packetizer.cpp b/modules/compiler/vecz/source/transform/packetizer.cpp
@@ -536,7 +536,7 @@ bool Packetizer::Impl::packetize() {
               B, VectorType::get(B.getInt32Ty(), SimdWidth), "index.vec");
         }
 
-        // CA-3943 this implementation looks unlikely to be correct, but for
+        // This implementation looks unlikely to be correct, but for
         // now we just maintain the original behaviour, until we have a better
         // idea of what is going on or whether any of this is still needed.
         // This case will never be encountered during kernel vectorization.
@@ -1239,7 +1239,7 @@ Value *Packetizer::Impl::packetizeGroupReduction(Instruction *I) {
 
   // Reduce the packet values in-place.
   // TODO: can we add 'reassoc' to the floating-point reductions to absolve
-  // them of ordering? See CA-3969.
+  // them of ordering?
   op.getPacketValues(packetWidth, opPackets);
 
   assert((!VL || packetWidth) &&
@@ -2088,8 +2088,7 @@ ValuePacket Packetizer::Impl::packetizeCall(CallInst *CI) {
     // Some llvm intrinsic functions like abs have argument that are constants
     // and define as llvm_i1_ty. This means that thoses operand can't
     // be packetized. To solve that temporary, we use this vector so every
-    // cases can set independently what operand must be skipped
-    // CA-3696
+    // cases can set independently what operand must be skipped.
     SmallVector<bool, maxOperands> operandsToSkip(maxOperands, false);
     switch (IntrID) {
       case Intrinsic::abs:
diff --git a/modules/compiler/vecz/source/transform/ternary_transform_pass.cpp b/modules/compiler/vecz/source/transform/ternary_transform_pass.cpp
@@ -107,7 +107,7 @@ bool shouldTransform(SelectInst *Select, const StrideAnalysisResult &SAR) {
     // scalar loads and stores. Performing this transform on vectors was
     // historically banned due to internal limitations, but these days we
     // *should* be able to. It's just that we don't know whether it's
-    // beneficial: see CA-4337.
+    // beneficial.
     for (User *U : GEP->users()) {
       if (auto *const LI = dyn_cast<LoadInst>(U)) {
         if (LI->getType()->isVectorTy()) {
diff --git a/modules/compiler/vecz/source/vector_target_info.cpp b/modules/compiler/vecz/source/vector_target_info.cpp
@@ -1300,7 +1300,7 @@ unsigned TargetInfo::getVectorWidthForType(const llvm::TargetTransformInfo &TTI,
   }
 
   // The floor of 8 prevents poor double precision performance.
-  // Not sure why (CA-3461 related?)
+  // Not sure why.
   return std::max(MaxVecRegBitWidth / BitWidth, 8u);
 }
 
diff --git a/modules/compiler/vecz/source/vector_target_info_arm.cpp b/modules/compiler/vecz/source/vector_target_info_arm.cpp
@@ -171,7 +171,7 @@ bool TargetInfoArm::optimizeInterleavedGroup(IRBuilder<> &B,
   VECZ_FAIL_IF(HasMask);
   VECZ_FAIL_IF(stride < 0);
 
-  // TODO CA-3100 fetch information on SubTargetInfo
+  // TODO fetch information on SubTargetInfo
   // load instructions seems to be easily split in the backend whereas stores
   // generate a backend error because of invalid data type on vector operands.
   // Vector operands are enabled in the backend only when SubTargetInfo ensures
@@ -337,7 +337,7 @@ bool TargetInfoAArch64::optimizeInterleavedGroup(
   VECZ_FAIL_IF(HasMask);
   VECZ_FAIL_IF(stride < 0);
 
-  // TODO CA-3100 fetch information on SubTargetInfo
+  // TODO fetch information on SubTargetInfo
   // load instructions seems to be easily split in the backend whereas stores
   // generate a backend error because of invalid data type on vector operands.
   // Vector operands are enabled in the backend only when SubTargetInfo ensures
diff --git a/modules/compiler/vecz/source/vector_target_info_riscv.cpp b/modules/compiler/vecz/source/vector_target_info_riscv.cpp
@@ -101,7 +101,7 @@ bool TargetInfoRISCV::isVPVectorLegal(const llvm::Function &F,
 
 // Should be target-dependent. Take RISCV legal types for now.
 // FIXME: LLVM 14 adds better support for legalization of vp intrinsics, but
-// not RISCV ones like vrgather_vv. See CA-4071.
+// not RISCV ones like vrgather_vv.
 bool TargetInfoRISCV::isVectorTypeLegal(Type *Ty) const {
   assert(Ty->isVectorTy() && "Expecting a vector type.");
   (void)Ty;
diff --git a/modules/compiler/vecz/source/vectorization_context.cpp b/modules/compiler/vecz/source/vectorization_context.cpp
@@ -233,7 +233,7 @@ Function *VectorizationContext::getOrCreateMaskedFunction(CallInst *CI) {
   if (!F) {
     F = dyn_cast<Function>(CI->getCalledOperand()->stripPointerCasts());
   }
-  VECZ_FAIL_IF(!F);  // TODO: CA-1505: Support indirect function calls.
+  VECZ_FAIL_IF(!F);  // TODO: Support indirect function calls.
   LLVMContext &ctx = F->getContext();
 
   // We will handle printf statements, but handling every possible vararg
diff --git a/modules/compiler/vecz/test/lit/llvm/basic_vecz_mem2reg.ll b/modules/compiler/vecz/test/lit/llvm/basic_vecz_mem2reg.ll
@@ -59,14 +59,14 @@ declare i64 @__mux_get_global_id(i32)
 ; CHECK:  %4 = bitcast i32 %3 to <2 x i16>
 
 ; Note: we can't optimize this as the allocated type size and loaded type sizes
-; don't match. Maybe we could trunc %3 from i32 to i16? See CA-4382.
+; don't match. Maybe we could trunc %3 from i32 to i16?
 
 ; CHECK: define spir_kernel void @__vecz_v4_load_type_size_mismatch_no_bitcast(ptr addrspace(1) %p)
 ; CHECK:  %data = alloca i32, align 4
 ; CHECK:  %4 = load i16, ptr %data, align 2
 
 ; Note: we can't optimize this as the allocated type size and loaded type sizes
-; don't match. Maybe we could trunc %3 from i32 to i16? See CA-4382.
+; don't match. Maybe we could trunc %3 from i32 to i16?
 
 ; CHECK: define spir_kernel void @__vecz_v4_store_type_size_mismatch_no_bitcast(ptr addrspace(1) %p)
 ; CHECK:  %data = alloca i32, align 4
diff --git a/modules/compiler/vecz/test/lit/llvm/builtin_pointer_return.ll b/modules/compiler/vecz/test/lit/llvm/builtin_pointer_return.ll
@@ -26,8 +26,7 @@ declare spir_func <2 x float> @_Z5fractDv2_fPS_(<2 x float>, <2 x float>*)
 declare spir_func <4 x float> @_Z5fractDv4_fPS_(<4 x float>, <4 x float>*)
 declare spir_func <8 x float> @_Z5fractDv8_fPS_(<8 x float>, <8 x float>*)
 
-; FIXME: Both of these are instantiating when we have vector equivalents: see
-; CA-4046.
+; FIXME: Both of these are instantiating when we have vector equivalents.
 
 define spir_kernel void @fract_v1(float* %xptr, float* %outptr, float* %ioutptr) {
   %iouta = alloca float
diff --git a/modules/compiler/vecz/test/lit/llvm/partial_linearization_exit_masks.ll b/modules/compiler/vecz/test/lit/llvm/partial_linearization_exit_masks.ll
@@ -18,7 +18,7 @@
 
 ; This test ensures that VECZ does not crash during control flow conversion due
 ; to a missing exit mask. As such, we need only verify that the return code from
-; veczc is 0, and FileCheck is not required. See CA-3117 for details.
+; veczc is 0, and FileCheck is not required.
 
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 target triple = "spir64-unknown-unknown"
diff --git a/modules/compiler/vecz/test/lit/llvm/ternary_transform.ll b/modules/compiler/vecz/test/lit/llvm/ternary_transform.ll
@@ -107,7 +107,7 @@ declare i64 @__mux_get_global_id(i32)
 ; CHECK: %c2 = select i1 %cond, ptr %c0, ptr %c1
 ; CHECK: store i64 %b, ptr %c2, align 4
 
-; Note: we don't perform this transform on vector accesses - see CA-4337.
+; Note: we don't perform this transform on vector accesses.
 ; CHECK: define spir_kernel void @__vecz_v4_test_vector_scalar_cond(i64 %a, <2 x i32> %b, ptr %c)
 ; CHECK:   %gid = call i64 @__mux_get_global_id(i32 0)
 ; CHECK:   %cond = icmp eq i64 %a, %gid

Original file line number	Diff line number	Diff line change
`@@ -314,7 +314,7 @@ BuiltinUniformity BuiltinInfo::isBuiltinUniform(const Builtin &B,`
`314`	`314`	`case eMuxBuiltinGetLocalLinearId:`
`315`	`315`	`case eMuxBuiltinGetGlobalLinearId:`
`316`	`316`	`// TODO: This is fine for vectorizing in the x-axis, but currently we do`
`317`		`- // not support vectorizing along y or z (see CA-2843).`
	`317`	`+ // not support vectorizing along y or z.`
`318`	`318`	`return SimdDimIdx ? eBuiltinUniformityNever`
`319`	`319`	`: eBuiltinUniformityInstanceID;`
`320`	`320`	`}`
Original file line number	Diff line number	Diff line change
`@@ -696,7 +696,6 @@ PreservedAnalyses compiler::utils::ReplaceLocalModuleScopeVariablesPass::run(`
`696`	`696`
`697`	`697`	`// We can't guarantee a subprogram for all functions.`
`698`	`698`	`// FIXME: Should we be able to? Do we need to clone subprograms somehow?`
`699`		`- // See CA-4241.`
`700`	`699`	`if (!DISubprogram) {`
`701`	`700`	`continue;`
`702`	`701`	`}`
Original file line number	Diff line number	Diff line change
`@@ -74,7 +74,7 @@ static bool shouldClone(compiler::utils::StructTypeRemapper &StructTypeRemapper,`
`74`	`74`	`}`
`75`	`75`	`}`
`76`	`76`
`77`		`- // TODO: Check globals (see CA-3833).`
	`77`	`+ // TODO: Check globals.`
`78`	`78`
`79`	`79`	`// If an instruction makes use of a type but`
`80`	`80`	`// isn't of that type e.g. a cast it will necessarily get caught by`
Original file line number	Diff line number	Diff line change
`@@ -1921,7 +1921,7 @@ PreservedAnalyses compiler::utils::WorkItemLoopsPass::run(`
`1921`	`1921`	`// FIXME: Also mark them as internal in case they contain symbols we`
`1922`	`1922`	`// haven't resolved as part of the work-item loop wrapping process. We`
`1923`	`1923`	`// rely on GlobalOptPass to remove such functions; this is the same root`
`1924`		`- // issue as CA-4126.`
	`1924`	`+ // issue as some mux builtins require DCE for correctness.`
`1925`	`1925`	`F.setLinkage(GlobalValue::InternalLinkage);`
`1926`	`1926`	`}`
`1927`	`1927`	`}`