[INFO] Initializing environment for https://gitcode.com/pre-commit-clang/mirrors-clang-format.
[INFO] Installing environment for https://gitcode.com/pre-commit-clang/mirrors-clang-format.
[INFO] Once installed this environment will be reused.
[INFO] This may take a few minutes...
clang-format.............................................................Failed
- hook id: clang-format
- files were modified by this hook
Check....................................................................Passed
All changes made by hooks:
diff --git a/blas/spr/arch22/sspr_host.cpp b/blas/spr/arch22/sspr_host.cpp
index 8e6fb9b..61bf2a4 100644
--- a/blas/spr/arch22/sspr_host.cpp
+++ b/blas/spr/arch22/sspr_host.cpp
@@ -26,12 +26,11 @@ void sspr_kernel_do(uint8_t* x, uint8_t* ap, const SsprTilingData& tiling, uint3
 static aclblasStatus_t ValidateSsprParams(
     aclblasFillMode_t uplo, int incx, const float* alpha, const float* x, const float* ap)
 {
-    CHECK_RET(
-        alpha != nullptr, OP_LOGE("aclblasSspr", "alpha must not be nullptr"); return ACLBLAS_STATUS_INVALID_VALUE);
-    CHECK_RET(
-        uplo == ACLBLAS_UPPER || uplo == ACLBLAS_LOWER,
-        OP_LOGE("aclblasSspr", "uplo must be UPPER(121) or LOWER(122), got %d", static_cast<int>(uplo));
-        return ACLBLAS_STATUS_INVALID_VALUE);
+    CHECK_RET(alpha != nullptr, OP_LOGE("aclblasSspr", "alpha must not be nullptr");
+              return ACLBLAS_STATUS_INVALID_VALUE);
+    CHECK_RET(uplo == ACLBLAS_UPPER || uplo == ACLBLAS_LOWER,
+              OP_LOGE("aclblasSspr", "uplo must be UPPER(121) or LOWER(122), got %d", static_cast<int>(uplo));
+              return ACLBLAS_STATUS_INVALID_VALUE);
     CHECK_RET(incx != 0, OP_LOGE("aclblasSspr", "incx must not be zero"); return ACLBLAS_STATUS_INVALID_VALUE);
     CHECK_RET(incx != INT_MIN, OP_LOGE("aclblasSspr", "incx must not be INT_MIN"); return ACLBLAS_STATUS_INVALID_VALUE);
     CHECK_RET(x != nullptr, OP_LOGE("aclblasSspr", "x must not be nullptr"); return ACLBLAS_STATUS_INVALID_VALUE);
@@ -81,9 +80,8 @@ aclblasStatus_t aclblasSspr(
         OP_LOGE("aclblasSspr", "vector core count is 0");
         return ACLBLAS_STATUS_EXECUTION_FAILED;
     }
-    uint32_t useNumBlocks = (incx != 1 && n >= 4096)
-        ? aivCoreNum
-        : std::min(CeilDiv<uint32_t>(n, SIMT_MIN_THREAD_NUM), aivCoreNum);
+    uint32_t useNumBlocks =
+        (incx != 1 && n >= 4096) ? aivCoreNum : std::min(CeilDiv<uint32_t>(n, SIMT_MIN_THREAD_NUM), aivCoreNum);
 
     SsprTilingData tiling = CalSsprTilingData(useNumBlocks, n, uplo, alphaVal, incx);
 
diff --git a/blas/spr/arch22/sspr_kernel.cpp b/blas/spr/arch22/sspr_kernel.cpp
index 88aff04..724faad 100644
--- a/blas/spr/arch22/sspr_kernel.cpp
+++ b/blas/spr/arch22/sspr_kernel.cpp
@@ -63,14 +63,15 @@ constexpr uint32_t kLocalColFloats = 8192;
 // gained at most 2.0% geometric mean but regressed individual sizes by up to 10.2%, so LOWER
 // retains the same coefficients as UPPER. Each term has an O(1) closed form so the boundary
 // search below stays O(log n).
-constexpr uint64_t kCostPerColumn = 2048;    // with kCostPerTransfer: p = 6144 per column
-constexpr uint64_t kCostPerElement = 1;      // per packed element streamed
-constexpr uint64_t kCostPerTransfer = 4096;  // per kChunkFloats-sized transfer
+constexpr uint64_t kCostPerColumn = 2048;   // with kCostPerTransfer: p = 6144 per column
+constexpr uint64_t kCostPerElement = 1;     // per packed element streamed
+constexpr uint64_t kCostPerTransfer = 4096; // per kChunkFloats-sized transfer
 
 // sum_{len=1..m} ceil(len / grain)
 __aicore__ inline uint64_t CeilRunSum(uint64_t m, uint64_t grain)
 {
-    if (m == 0ULL) return 0ULL;
+    if (m == 0ULL)
+        return 0ULL;
     const uint64_t q = m / grain;
     const uint64_t r = m - q * grain;
     return grain * q * (q + 1ULL) / 2ULL + r * (q + 1ULL);
@@ -79,15 +80,13 @@ __aicore__ inline uint64_t CeilRunSum(uint64_t m, uint64_t grain)
 __aicore__ inline uint64_t ColumnWorkPrefix(uint32_t columns, uint32_t n, uint32_t uplo)
 {
     // Column c has length c+1 (UPPER) or n-c (LOWER); the sums below run over [0, columns).
-    const uint64_t elements = uplo == ACLBLAS_UPPER
-        ? static_cast<uint64_t>(columns) * (columns + 1ULL) / 2ULL
-        : static_cast<uint64_t>(columns) * (2ULL * n - columns + 1ULL) / 2ULL;
-    const uint64_t transfers = uplo == ACLBLAS_UPPER
-        ? CeilRunSum(columns, kChunkFloats)
-        : CeilRunSum(n, kChunkFloats) - CeilRunSum(static_cast<uint64_t>(n) - columns, kChunkFloats);
-    return kCostPerColumn * static_cast<uint64_t>(columns)
-         + kCostPerElement * elements
-         + kCostPerTransfer * transfers;
+    const uint64_t elements = uplo == ACLBLAS_UPPER ?
+                                  static_cast<uint64_t>(columns) * (columns + 1ULL) / 2ULL :
+                                  static_cast<uint64_t>(columns) * (2ULL * n - columns + 1ULL) / 2ULL;
+    const uint64_t transfers = uplo == ACLBLAS_UPPER ? CeilRunSum(columns, kChunkFloats) :
+                                                       CeilRunSum(n, kChunkFloats) -
+                                                           CeilRunSum(static_cast<uint64_t>(n) - columns, kChunkFloats);
+    return kCostPerColumn * static_cast<uint64_t>(columns) + kCostPerElement * elements + kCostPerTransfer * transfers;
 }
 
 __aicore__ inline uint64_t PartitionTarget(uint64_t total, uint32_t part, uint32_t parts)
@@ -95,26 +94,26 @@ __aicore__ inline uint64_t PartitionTarget(uint64_t total, uint32_t part, uint32
     return (total / parts) * part + (total % parts) * part / parts;
 }
 
-__aicore__ inline uint32_t WeightedColumnBoundary(
-    uint32_t part, uint32_t parts, uint32_t n, uint32_t uplo)
+__aicore__ inline uint32_t WeightedColumnBoundary(uint32_t part, uint32_t parts, uint32_t n, uint32_t uplo)
 {
-    if (part == 0U) return 0U;
-    if (part >= parts) return n;
+    if (part == 0U)
+        return 0U;
+    if (part >= parts)
+        return n;
     const uint64_t target = PartitionTarget(ColumnWorkPrefix(n, n, uplo), part, parts);
     uint32_t lo = 0U;
     uint32_t hi = n;
     while (lo < hi) {
         const uint32_t mid = lo + (hi - lo) / 2U;
-        if (ColumnWorkPrefix(mid, n, uplo) < target) lo = mid + 1U;
-        else hi = mid;
+        if (ColumnWorkPrefix(mid, n, uplo) < target)
+            lo = mid + 1U;
+        else
+            hi = mid;
     }
     return lo;
 }
 
-__aicore__ inline uint64_t UpperColumnBase(uint32_t col)
-{
-    return static_cast<uint64_t>(col) * (col + 1ULL) / 2ULL;
-}
+__aicore__ inline uint64_t UpperColumnBase(uint32_t col) { return static_cast<uint64_t>(col) * (col + 1ULL) / 2ULL; }
 
 __aicore__ inline uint64_t LowerColumnBase(uint32_t col, uint32_t n)
 {
@@ -124,13 +123,12 @@ __aicore__ inline uint64_t LowerColumnBase(uint32_t col, uint32_t n)
 __aicore__ inline uint64_t XPhysicalIndex(uint32_t logical, uint32_t n, int64_t incx)
 {
     const uint64_t absIncx = static_cast<uint64_t>(incx >= 0 ? incx : -incx);
-    return incx >= 0 ? static_cast<uint64_t>(logical) * absIncx
-                     : static_cast<uint64_t>(n - 1U - logical) * absIncx;
+    return incx >= 0 ? static_cast<uint64_t>(logical) * absIncx : static_cast<uint64_t>(n - 1U - logical) * absIncx;
 }
 
 __aicore__ inline void ScalarColumns(
-    GlobalTensor<float>& xGm, GlobalTensor<float>& apGm, const SsprTilingData& tiling,
-    uint32_t colStart, uint32_t colEnd)
+    GlobalTensor<float>& xGm, GlobalTensor<float>& apGm, const SsprTilingData& tiling, uint32_t colStart,
+    uint32_t colEnd)
 {
     for (uint32_t col = colStart; col < colEnd; ++col) {
         const float axCol = tiling.alpha * xGm.GetValue(XPhysicalIndex(col, tiling.n, tiling.incx));
@@ -138,15 +136,15 @@ __aicore__ inline void ScalarColumns(
             const uint64_t base = UpperColumnBase(col);
             for (uint32_t row = 0; row <= col; ++row) {
                 const uint64_t apIndex = base + row;
-                apGm.SetValue(apIndex, apGm.GetValue(apIndex) +
-                                           axCol * xGm.GetValue(XPhysicalIndex(row, tiling.n, tiling.incx)));
+                apGm.SetValue(
+                    apIndex, apGm.GetValue(apIndex) + axCol * xGm.GetValue(XPhysicalIndex(row, tiling.n, tiling.incx)));
             }
         } else {
             const uint64_t base = LowerColumnBase(col, tiling.n);
             for (uint32_t row = col; row < tiling.n; ++row) {
                 const uint64_t apIndex = base + (row - col);
-                apGm.SetValue(apIndex, apGm.GetValue(apIndex) +
-                                           axCol * xGm.GetValue(XPhysicalIndex(row, tiling.n, tiling.incx)));
+                apGm.SetValue(
+                    apIndex, apGm.GetValue(apIndex) + axCol * xGm.GetValue(XPhysicalIndex(row, tiling.n, tiling.incx)));
             }
         }
     }
@@ -154,8 +152,8 @@ __aicore__ inline void ScalarColumns(
 
 #if __DAV_C220_VEC__
 __aicore__ inline void CanonicalizeStridedRange(
-    LocalTensor<float> cache, LocalTensor<float> span, LocalTensor<uint32_t> offsets,
-    __gm__ float* source, uint32_t n, int64_t incx, uint32_t logicalBase, uint32_t logicalCount)
+    LocalTensor<float> cache, LocalTensor<float> span, LocalTensor<uint32_t> offsets, __gm__ float* source, uint32_t n,
+    int64_t incx, uint32_t logicalBase, uint32_t logicalCount)
 {
     const uint32_t increment = static_cast<uint32_t>(incx > 0 ? incx : -incx);
     auto ubSpan = reinterpret_cast<__ubuf__ float*>(span.GetPhyAddr());
@@ -185,11 +183,9 @@ __aicore__ inline void CanonicalizeStridedRange(
     }
 
     for (uint32_t localStart = 0; localStart < logicalCount; localStart += logicalTile) {
-        const uint32_t count = (logicalCount - localStart < logicalTile)
-                                   ? logicalCount - localStart : logicalTile;
+        const uint32_t count = (logicalCount - localStart < logicalTile) ? logicalCount - localStart : logicalTile;
         const uint32_t globalStart = logicalBase + localStart;
-        const uint32_t physicalStart = incx > 0 ? globalStart * increment
-                                                : (n - globalStart - count) * increment;
+        const uint32_t physicalStart = incx > 0 ? globalStart * increment : (n - globalStart - count) * increment;
         const uint32_t spanCount = (count - 1U) * increment + 1U;
         for (uint32_t i = 0; i < count; ++i) {
             const uint32_t sourceIndex = incx > 0 ? i * increment : (count - 1U - i) * increment;
@@ -197,8 +193,7 @@ __aicore__ inline void CanonicalizeStridedRange(
         }
         set_flag(PIPE_S, PIPE_MTE2, EVENT_ID1);
         wait_flag(PIPE_S, PIPE_MTE2, EVENT_ID1);
-        copy_gm_to_ubuf_align_b32(
-            ubSpan, source + physicalStart, 0, 1, spanCount * sizeof(float), 0, 0, 0, 0);
+        copy_gm_to_ubuf_align_b32(ubSpan, source + physicalStart, 0, 1, spanCount * sizeof(float), 0, 0, 0, 0);
         set_flag(PIPE_MTE2, PIPE_V, EVENT_ID1);
         wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID1);
         Gather(cache[localStart], span, offsets, 0U, count);
@@ -208,8 +203,7 @@ __aicore__ inline void CanonicalizeStridedRange(
     }
 }
 __aicore__ inline void CachedStridedColumns(
-    __gm__ float* gmX, GlobalTensor<float>& apGm, const SsprTilingData& tiling,
-    uint32_t colStart, uint32_t colEnd)
+    __gm__ float* gmX, GlobalTensor<float>& apGm, const SsprTilingData& tiling, uint32_t colStart, uint32_t colEnd)
 {
     constexpr uint32_t kRowTile = kLocalRowFloats;
     TPipe pipe;
@@ -235,7 +229,8 @@ __aicore__ inline void CachedStridedColumns(
     auto ubXCol = reinterpret_cast<__ubuf__ float*>(xCol.GetPhyAddr());
     auto ubOffsets = reinterpret_cast<__ubuf__ uint32_t*>(offsets.GetPhyAddr());
     const uint32_t colCount = colEnd - colStart;
-    if (colCount > kLocalColFloats) return;
+    if (colCount > kLocalColFloats)
+        return;
     CanonicalizeStridedRange(xCol, span, offsets, gmX, tiling.n, tiling.incx, colStart, colCount);
 
     for (uint32_t tileRow = 0; tileRow < tiling.n; tileRow += kRowTile) {
@@ -245,24 +240,22 @@ __aicore__ inline void CachedStridedColumns(
         const DataCopyExtParams fullCopy{1, fullCount * static_cast<uint32_t>(sizeof(float)), 0, 0, 0};
         const DataCopyPadExtParams<float> pad{true, 0, 0, 0.0f};
 
-        const uint32_t rectBegin = tiling.uplo == ACLBLAS_UPPER
-            ? (colStart > tileEnd ? colStart : tileEnd) : colStart;
-        const uint32_t rectEnd = tiling.uplo == ACLBLAS_UPPER
-            ? colEnd : (colEnd < tileRow ? colEnd : tileRow);
+        const uint32_t rectBegin = tiling.uplo == ACLBLAS_UPPER ? (colStart > tileEnd ? colStart : tileEnd) : colStart;
+        const uint32_t rectEnd = tiling.uplo == ACLBLAS_UPPER ? colEnd : (colEnd < tileRow ? colEnd : tileRow);
         if (rectBegin < rectEnd) {
             const uint32_t firstCol = rectBegin;
-            const uint64_t firstOffset = tiling.uplo == ACLBLAS_UPPER
-                ? UpperColumnBase(firstCol) + tileRow
-                : LowerColumnBase(firstCol, tiling.n) + (tileRow - firstCol);
+            const uint64_t firstOffset = tiling.uplo == ACLBLAS_UPPER ?
+                                             UpperColumnBase(firstCol) + tileRow :
+                                             LowerColumnBase(firstCol, tiling.n) + (tileRow - firstCol);
             LocalTensor<float> firstAp = apQueue.AllocTensor<float>();
             DataCopyPad(firstAp, apGm[firstOffset], fullCopy, pad);
             apQueue.EnQue(firstAp);
             for (uint32_t col = rectBegin; col < rectEnd; ++col) {
                 if (col + 1U < rectEnd) {
                     const uint32_t nextCol = col + 1U;
-                    const uint64_t nextOffset = tiling.uplo == ACLBLAS_UPPER
-                        ? UpperColumnBase(nextCol) + tileRow
-                        : LowerColumnBase(nextCol, tiling.n) + (tileRow - nextCol);
+                    const uint64_t nextOffset = tiling.uplo == ACLBLAS_UPPER ?
+                                                    UpperColumnBase(nextCol) + tileRow :
+                                                    LowerColumnBase(nextCol, tiling.n) + (tileRow - nextCol);
                     LocalTensor<float> nextAp = apQueue.AllocTensor<float>();
                     DataCopyPad(nextAp, apGm[nextOffset], fullCopy, pad);
                     apQueue.EnQue(nextAp);
@@ -275,9 +268,9 @@ __aicore__ inline void CachedStridedColumns(
                 outQueue.EnQue(out);
                 apQueue.FreeTensor(apIn);
                 out = outQueue.DeQue<float>();
-                const uint64_t apOffset = tiling.uplo == ACLBLAS_UPPER
-                    ? UpperColumnBase(col) + tileRow
-                    : LowerColumnBase(col, tiling.n) + (tileRow - col);
+                const uint64_t apOffset = tiling.uplo == ACLBLAS_UPPER ?
+                                              UpperColumnBase(col) + tileRow :
+                                              LowerColumnBase(col, tiling.n) + (tileRow - col);
                 DataCopyPad(apGm[apOffset], out, fullCopy);
                 outQueue.FreeTensor(out);
             }
@@ -342,8 +335,8 @@ __aicore__ inline void CachedStridedColumns(
 }
 
 __aicore__ inline void UnifiedContiguousColumns(
-    GlobalTensor<float>& xGm, GlobalTensor<float>& apGm, const SsprTilingData& tiling,
-    uint32_t colStart, uint32_t colEnd)
+    GlobalTensor<float>& xGm, GlobalTensor<float>& apGm, const SsprTilingData& tiling, uint32_t colStart,
+    uint32_t colEnd)
 {
     constexpr uint32_t kRowTile = 16384U;
     constexpr uint32_t kPipelineFloats = 4096U;
@@ -371,9 +364,10 @@ __aicore__ inline void UnifiedContiguousColumns(
             const uint32_t rowEnd = tiling.uplo == ACLBLAS_UPPER ? col + 1U : tiling.n;
             uint32_t row = rowBegin > tileRow ? rowBegin : tileRow;
             const uint32_t segmentEnd = rowEnd < tileEnd ? rowEnd : tileEnd;
-            if (row >= segmentEnd) continue;
-            const uint64_t apBase = tiling.uplo == ACLBLAS_UPPER ? UpperColumnBase(col)
-                                                                 : LowerColumnBase(col, tiling.n);
+            if (row >= segmentEnd)
+                continue;
+            const uint64_t apBase =
+                tiling.uplo == ACLBLAS_UPPER ? UpperColumnBase(col) : LowerColumnBase(col, tiling.n);
             const float axCol = tiling.alpha * xGm.GetValue(col);
 
             // A LOWER diagonal suffix may start at an arbitrary Unified Buffer offset.
@@ -382,7 +376,8 @@ __aicore__ inline void UnifiedContiguousColumns(
             const uint32_t localStart = row - tileRow;
             if ((localStart & (kPipelineFloats - 1U)) != 0U) {
                 uint32_t alignedRow = (row + kPipelineFloats - 1U) & ~(kPipelineFloats - 1U);
-                if (alignedRow > segmentEnd) alignedRow = segmentEnd;
+                if (alignedRow > segmentEnd)
+                    alignedRow = segmentEnd;
                 const uint32_t count = alignedRow - row;
                 const uint64_t apOffset = apBase + row - rowBegin;
                 const DataCopyExtParams copy{1, count * static_cast<uint32_t>(sizeof(float)), 0, 0, 0};
@@ -407,11 +402,9 @@ __aicore__ inline void UnifiedContiguousColumns(
             }
 
             if (row < segmentEnd) {
-                uint32_t currentCount = (segmentEnd - row < kPipelineFloats)
-                    ? segmentEnd - row : kPipelineFloats;
+                uint32_t currentCount = (segmentEnd - row < kPipelineFloats) ? segmentEnd - row : kPipelineFloats;
                 uint64_t currentOffset = apBase + row - rowBegin;
-                DataCopyExtParams currentCopy{
-                    1, currentCount * static_cast<uint32_t>(sizeof(float)), 0, 0, 0};
+                DataCopyExtParams currentCopy{1, currentCount * static_cast<uint32_t>(sizeof(float)), 0, 0, 0};
                 LocalTensor<float> firstAp = apQueue.AllocTensor<float>();
                 DataCopyPad(firstAp, apGm[currentOffset], currentCopy, pad);
                 apQueue.EnQue(firstAp);
@@ -419,11 +412,10 @@ __aicore__ inline void UnifiedContiguousColumns(
                 while (row < segmentEnd) {
                     const uint32_t nextRow = row + currentCount;
                     if (nextRow < segmentEnd) {
-                        const uint32_t nextCount = (segmentEnd - nextRow < kPipelineFloats)
-                            ? segmentEnd - nextRow : kPipelineFloats;
+                        const uint32_t nextCount =
+                            (segmentEnd - nextRow < kPipelineFloats) ? segmentEnd - nextRow : kPipelineFloats;
                         const uint64_t nextOffset = apBase + nextRow - rowBegin;
-                        const DataCopyExtParams nextCopy{
-                            1, nextCount * static_cast<uint32_t>(sizeof(float)), 0, 0, 0};
+                        const DataCopyExtParams nextCopy{1, nextCount * static_cast<uint32_t>(sizeof(float)), 0, 0, 0};
                         LocalTensor<float> nextAp = apQueue.AllocTensor<float>();
                         DataCopyPad(nextAp, apGm[nextOffset], nextCopy, pad);
                         apQueue.EnQue(nextAp);
@@ -441,11 +433,10 @@ __aicore__ inline void UnifiedContiguousColumns(
 
                     row = nextRow;
                     if (row < segmentEnd) {
-                        currentCount = (segmentEnd - row < kPipelineFloats)
-                            ? segmentEnd - row : kPipelineFloats;
+                        currentCount = (segmentEnd - row < kPipelineFloats) ? segmentEnd - row : kPipelineFloats;
                         currentOffset = apBase + row - rowBegin;
-                        currentCopy = DataCopyExtParams{
-                            1, currentCount * static_cast<uint32_t>(sizeof(float)), 0, 0, 0};
+                        currentCopy =
+                            DataCopyExtParams{1, currentCount * static_cast<uint32_t>(sizeof(float)), 0, 0, 0};
                     }
                 }
             }
@@ -454,20 +445,18 @@ __aicore__ inline void UnifiedContiguousColumns(
 }
 
 #endif
-}
+} // namespace
 
-extern "C" __global__ __aicore__ __vector__ void sspr_kernel(
-    GM_ADDR x, GM_ADDR ap, const SsprTilingData tiling)
+extern "C" __global__ __aicore__ __vector__ void sspr_kernel(GM_ADDR x, GM_ADDR ap, const SsprTilingData tiling)
 {
     KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
 
     GlobalTensor<float> xGm;
     GlobalTensor<float> apGm;
     const uint64_t absIncx = static_cast<uint64_t>(tiling.incx >= 0 ? tiling.incx : -tiling.incx);
-    xGm.SetGlobalBuffer(reinterpret_cast<__gm__ float*>(x),
-                        static_cast<uint64_t>(tiling.n - 1U) * absIncx + 1ULL);
-    apGm.SetGlobalBuffer(reinterpret_cast<__gm__ float*>(ap),
-                         static_cast<uint64_t>(tiling.n) * (tiling.n + 1ULL) / 2ULL);
+    xGm.SetGlobalBuffer(reinterpret_cast<__gm__ float*>(x), static_cast<uint64_t>(tiling.n - 1U) * absIncx + 1ULL);
+    apGm.SetGlobalBuffer(
+        reinterpret_cast<__gm__ float*>(ap), static_cast<uint64_t>(tiling.n) * (tiling.n + 1ULL) / 2ULL);
 
     const uint32_t block = GetBlockIdx();
     const uint32_t blocks = GetBlockNum();
@@ -478,7 +467,8 @@ extern "C" __global__ __aicore__ __vector__ void sspr_kernel(
     // strided large-n cases because its element-count model was wrong for short columns.
     colStart = WeightedColumnBoundary(block, blocks, tiling.n, tiling.uplo);
     colEnd = WeightedColumnBoundary(block + 1U, blocks, tiling.n, tiling.uplo);
-    if (colStart >= colEnd) return;
+    if (colStart >= colEnd)
+        return;
 
 #if __DAV_C220_VEC__
     if (tiling.incx == 1) {
@@ -494,8 +484,7 @@ extern "C" __global__ __aicore__ __vector__ void sspr_kernel(
     ScalarColumns(xGm, apGm, tiling, colStart, colEnd);
 }
 
-void sspr_kernel_do(
-    GM_ADDR x, GM_ADDR ap, const SsprTilingData& tiling, uint32_t numBlocks, void* stream)
+void sspr_kernel_do(GM_ADDR x, GM_ADDR ap, const SsprTilingData& tiling, uint32_t numBlocks, void* stream)
 {
     sspr_kernel<<<numBlocks, nullptr, stream>>>(x, ap, tiling);
 }
diff --git a/blas/spr/arch22/sspr_tiling_data.h b/blas/spr/arch22/sspr_tiling_data.h
index e535835..15460c7 100644
--- a/blas/spr/arch22/sspr_tiling_data.h
+++ b/blas/spr/arch22/sspr_tiling_data.h
@@ -23,10 +23,10 @@ inline constexpr uint32_t UB_X_FLOATS = 8192;
 inline constexpr uint32_t UB_THRESHOLD = 128;
 
 struct SsprTilingData {
-    uint32_t numThreads;         // threads per block
-    uint32_t columnsPerBlock;    // columns per block
-    uint32_t n;                  // matrix order
-    uint32_t uplo;               // ACLBLAS_UPPER(121) or ACLBLAS_LOWER(122)
-    float    alpha;              // scalar alpha (dereferenced from pointer)
-    int64_t  incx;               // x vector stride
+    uint32_t numThreads;      // threads per block
+    uint32_t columnsPerBlock; // columns per block
+    uint32_t n;               // matrix order
+    uint32_t uplo;            // ACLBLAS_UPPER(121) or ACLBLAS_LOWER(122)
+    float alpha;              // scalar alpha (dereferenced from pointer)
+    int64_t incx;             // x vector stride
 };
diff --git a/test/spr/arch22/sspr_npu_wrapper.h b/test/spr/arch22/sspr_npu_wrapper.h
index bc396c5..c652dcb 100644
--- a/test/spr/arch22/sspr_npu_wrapper.h
+++ b/test/spr/arch22/sspr_npu_wrapper.h
@@ -23,13 +23,7 @@
 // Otherwise, allocates device memory, copies H2D, invokes kernel,
 // synchronises, copies D2H, and frees.
 inline aclblasStatus_t aclblasSspr_npu(
-    aclblasHandle_t handle,
-    aclblasFillMode_t uplo,
-    int n,
-    const float* alpha,
-    const float* x,
-    int incx,
-    float* ap)
+    aclblasHandle_t handle, aclblasFillMode_t uplo, int n, const float* alpha, const float* x, int incx, float* ap)
 {
     if (handle == nullptr || n <= 0) {
         return aclblasSspr(handle, uplo, n, alpha, x, incx, ap);
@@ -46,26 +40,41 @@ inline aclblasStatus_t aclblasSspr_npu(
 
     if (x != nullptr) {
         aclRet = aclrtMalloc(&dX, xBytes, ACL_MEM_MALLOC_HUGE_FIRST);
-        if (aclRet != ACL_SUCCESS) return ACLBLAS_STATUS_ALLOC_FAILED;
+        if (aclRet != ACL_SUCCESS)
+            return ACLBLAS_STATUS_ALLOC_FAILED;
         aclRet = aclrtMemcpy(dX, xBytes, x, xBytes, ACL_MEMCPY_HOST_TO_DEVICE);
-        if (aclRet != ACL_SUCCESS) { aclrtFree(dX); return ACLBLAS_STATUS_INTERNAL_ERROR; }
+        if (aclRet != ACL_SUCCESS) {
+            aclrtFree(dX);
+            return ACLBLAS_STATUS_INTERNAL_ERROR;
+        }
     }
     if (ap != nullptr) {
         aclRet = aclrtMalloc(&dAP, apBytes, ACL_MEM_MALLOC_HUGE_FIRST);
-        if (aclRet != ACL_SUCCESS) { if (dX) aclrtFree(dX); return ACLBLAS_STATUS_ALLOC_FAILED; }
+        if (aclRet != ACL_SUCCESS) {
+            if (dX)
+                aclrtFree(dX);
+            return ACLBLAS_STATUS_ALLOC_FAILED;
+        }
         aclRet = aclrtMemcpy(dAP, apBytes, ap, apBytes, ACL_MEMCPY_HOST_TO_DEVICE);
-        if (aclRet != ACL_SUCCESS) { if (dX) aclrtFree(dX); aclrtFree(dAP); return ACLBLAS_STATUS_INTERNAL_ERROR; }
+        if (aclRet != ACL_SUCCESS) {
+            if (dX)
+                aclrtFree(dX);
+            aclrtFree(dAP);
+            return ACLBLAS_STATUS_INTERNAL_ERROR;
+        }
     }
 
-    aclblasStatus_t ret = aclblasSspr(handle, uplo, n, alpha,
-        static_cast<const float*>(dX), incx, static_cast<float*>(dAP));
+    aclblasStatus_t ret =
+        aclblasSspr(handle, uplo, n, alpha, static_cast<const float*>(dX), incx, static_cast<float*>(dAP));
 
     aclrtSynchronizeDevice();
     if (ret == ACLBLAS_STATUS_SUCCESS && ap != nullptr) {
         aclrtMemcpy(ap, apBytes, dAP, apBytes, ACL_MEMCPY_DEVICE_TO_HOST);
     }
 
-    if (dX) aclrtFree(dX);
-    if (dAP) aclrtFree(dAP);
+    if (dX)
+        aclrtFree(dX);
+    if (dAP)
+        aclrtFree(dAP);
     return ret;
 }
diff --git a/test/spr/arch22/sspr_test.cpp b/test/spr/arch22/sspr_test.cpp
index ffb498e..f4e962d 100644
--- a/test/spr/arch22/sspr_test.cpp
+++ b/test/spr/arch22/sspr_test.cpp
@@ -18,7 +18,7 @@
 #include "sspr_golden.h"
 #include "sspr_npu_wrapper.h"
 
-class SsprArch22Test : public BlasTest<SsprParam> { };
+class SsprArch22Test : public BlasTest<SsprParam> {};
 
 TEST_F(SsprArch22Test, NullHandle)
 {
@@ -28,8 +28,7 @@ TEST_F(SsprArch22Test, NullHandle)
 }
 
 INSTANTIATE_TEST_SUITE_P(
-    Sspr, SsprArch22Test,
-    ::testing::ValuesIn(GetCasesFromCsv<SsprParam>(ReplaceFileExtension2Csv(__FILE__))),
+    Sspr, SsprArch22Test, ::testing::ValuesIn(GetCasesFromCsv<SsprParam>(ReplaceFileExtension2Csv(__FILE__))),
     PrintCaseInfoString<SsprParam>);
 
 TEST_P(SsprArch22Test, CsvDriven)
@@ -61,16 +60,16 @@ TEST_P(SsprArch22Test, CsvDriven)
     float* apPtr = apHost.empty() ? nullptr : apHost.data();
     const float* alphaPtr = p.alphaNull ? nullptr : &p.alpha;
 
-    aclblasStatus_t ret = aclblasSspr_npu(
-        SsprArch22Test::handle_, p.uplo, p.n, alphaPtr, xPtr, p.incx, apPtr);
+    aclblasStatus_t ret = aclblasSspr_npu(SsprArch22Test::handle_, p.uplo, p.n, alphaPtr, xPtr, p.incx, apPtr);
     EXPECT_EQ(static_cast<int>(ret), static_cast<int>(p.expectResult));
-    if (p.expectResult != ACLBLAS_STATUS_SUCCESS) return;
-    if (p.n == 0) return;
+    if (p.expectResult != ACLBLAS_STATUS_SUCCESS)
+        return;
+    if (p.n == 0)
+        return;
 
     // Compute CPU golden
     std::vector<float> golden = apOrig;
-    aclblasSspr_cpu(
-        SsprArch22Test::handle_, p.uplo, p.n, &p.alpha, xHost.data(), p.incx, golden.data());
+    aclblasSspr_cpu(SsprArch22Test::handle_, p.uplo, p.n, &p.alpha, xHost.data(), p.incx, golden.data());
 
     VerifyConfig cfg;
     if (p.alpha == 0.0f && !p.alphaNull) {