clang-format | Format fbsource with clang-format 21.

Reviewed By: ChristianK275

Differential Revision: D85317706

fbshipit-source-id: b399c5c4b75252999442b7d7d2778e7a241b0025

Author: nlutsenko

bench/BenchUtils.h

@@ -274,8 +274,9 @@ void performance_test(
     aligned_vector<float> B(Bint.begin(), Bint.end());
     std::vector<std::unique_ptr<PackedGemmMatrixB<btype>>> Bp;
     for (int i = 0; i < num_instances; ++i) {
-      Bp.emplace_back(std::unique_ptr<PackedGemmMatrixB<btype>>(
-          new PackedGemmMatrixB<btype>(btran, k, n, alpha, B.data())));
+      Bp.emplace_back(
+          std::unique_ptr<PackedGemmMatrixB<btype>>(
+              new PackedGemmMatrixB<btype>(btran, k, n, alpha, B.data())));
     }
     auto kAligned = ((k * sizeof(float) + 64) & ~63) / sizeof(float);
     auto nAligned = ((n * sizeof(float) + 64) & ~63) / sizeof(float);

fbgemm_gpu/codegen/training/index_select/batch_index_select_dim0_cpu_host.cpp

@@ -403,8 +403,9 @@ class BatchIndexSelectDim0TensorCPUOp
         input_columns,
         permute_output_dim_0_1);
     ctx->saved_data["permute_output_dim_0_1"] = permute_output_dim_0_1;
-    ctx->save_for_backward(std::vector<Tensor>{
-        indices, input_num_indices, input_rows, input_columns, res[1]});
+    ctx->save_for_backward(
+        std::vector<Tensor>{
+            indices, input_num_indices, input_rows, input_columns, res[1]});
     res.resize(1);
     return res;
   }

fbgemm_gpu/codegen/training/index_select/batch_index_select_dim0_host.cpp

@@ -264,8 +264,9 @@ class BatchIndexSelectDim0GPUOp
 
     ctx->saved_data["permute_output_dim_0_1"] = permute_output_dim_0_1;
 
-    ctx->save_for_backward(std::vector<Tensor>{
-        inputs, indices, res[1], res[2], res[3], res[4], res[5], res[6]});
+    ctx->save_for_backward(
+        std::vector<Tensor>{
+            inputs, indices, res[1], res[2], res[3], res[4], res[5], res[6]});
 
     res.resize(1);
     return res;
@@ -584,8 +585,9 @@ class BatchIndexSelectDim0TensorGPUOp
 
     ctx->saved_data["permute_output_dim_0_1"] = permute_output_dim_0_1;
 
-    ctx->save_for_backward(std::vector<Tensor>{
-        inputs, indices, res[1], res[2], res[3], res[4], res[5], res[6]});
+    ctx->save_for_backward(
+        std::vector<Tensor>{
+            inputs, indices, res[1], res[2], res[3], res[4], res[5], res[6]});
 
     // res.resize(1);
     return res;

fbgemm_gpu/experimental/gen_ai/src/kv_cache/kv_cache_defs.cpp

@@ -18,16 +18,21 @@ TORCH_LIBRARY_FRAGMENT(fbgemm, m) {
   m.def("rope_qkv_varseq_prefill(Tensor XQ, Tensor(a!)? XK, Tensor? XV, Tensor(b!) cache_K, Tensor(c!) cache_V,  Tensor varseq_batch, Tensor varseq_seqpos, float theta, int? num_groups=1, Tensor? block_tables=None, int page_size=" STRING(
         DEFAULT_PAGE_SIZE) ", Tensor? actual_batch_size=None, Tensor? varseq_cache_seqpos=None, int cache_logical_dtype_int=0, bool rope_scaling=False, int old_context_len=8192"
         ", float scaling_factor=16, float lo_freq_factor=1, float hi_freq_factor=32,  Tensor? qparam_k=None, Tensor? qparam_v=None, bool write_k_back=False, bool k_norm=False,bool update_kv=True, Tensor?amax_qkv=None, Tensor?kv_quant_scale_precomputed=None) -> Tensor");
-  m.def("rope_qkv_decoding(Tensor XQ, Tensor? XK, Tensor? XV, Tensor(a!) cache_K, Tensor(b!) cache_V,  Tensor seqpos, float theta, int? num_groups=1, Tensor? block_tables=None, int page_size=" STRING(
-      DEFAULT_PAGE_SIZE) ", Tensor? actual_batch_size=None, Tensor? batch=None, Tensor? cache_seqpos=None,  int cache_logical_dtype_int=0, bool rope_scaling=False, int old_context_len=8192, float scaling_factor=16, float lo_freq_factor=1, float hi_freq_factor=32, Tensor? qparam_k=None, Tensor? qparam_v=None, bool k_norm=False, bool update_kv=True, Tensor?amax_qkv=None) -> Tensor");
-  m.def("nope_qkv_varseq_prefill(Tensor XQ, Tensor? XK, Tensor? XV, Tensor(a!) cache_K, Tensor(b!) cache_V,  Tensor varseq_batch, Tensor varseq_seqpos, Tensor? block_tables=None, int page_size=" STRING(
-      DEFAULT_PAGE_SIZE) ", Tensor? actual_batch_size=None, Tensor? varseq_cache_seqpos=None, int cache_logical_dtype_int=0, int? num_groups=1, Tensor? qparam_k=None, Tensor? qparam_v=None, bool k_norm=False, bool update_kv=True, Tensor?amax_qkv=None, Tensor?kv_quant_scale_precomputed=None) -> Tensor");
-  m.def("nope_qkv_decoding(Tensor XQ, Tensor? XK, Tensor? XV, Tensor(a!) cache_K, Tensor(b!) cache_V,  Tensor seqpos, Tensor? block_tables=None, int page_size=" STRING(
-      DEFAULT_PAGE_SIZE) ", Tensor? actual_batch_size=None, Tensor? batch=None, Tensor? cache_seqpos=None, int cache_logical_dtype_int=0, int? num_groups=1, Tensor? qparam_k=None, Tensor? qparam_v=None, bool k_norm=False, bool update_kv=True, Tensor?amax_qkv=None) -> Tensor");
-  m.def("xpos_qkv_varseq_prefill(Tensor XQ, Tensor XK, Tensor XV, Tensor(a!) cache_K, Tensor(b!) cache_V, Tensor varseq_batch, Tensor varseq_seqpos, float theta, float gamma, float scale_base, float exponent_offset, int? num_groups=1, Tensor? block_tables=None, int page_size=" STRING(
-      DEFAULT_PAGE_SIZE) ", Tensor? actual_batch_size=None, Tensor? varseq_cache_seqpos=None, int cache_logical_dtype_int=0, bool rope_scaling=False, int old_context_len=8192, float scaling_factor=16, float lo_freq_factor=1, float hi_freq_factor=32,  Tensor? qparam_k=None, Tensor? qparam_v=None) -> Tensor");
-  m.def("xpos_qkv_decoding(Tensor XQ, Tensor XK, Tensor XV, Tensor(a!) cache_K, Tensor(b!) cache_V,  Tensor seqpos, float theta, float gamma, float scale_base, float exponent_offset, int? num_groups=1, Tensor? block_tables=None, int page_size=" STRING(
-      DEFAULT_PAGE_SIZE) ", Tensor? actual_batch_size=None, Tensor? batch=None, Tensor? cache_seqpos=None, int cache_logical_dtype_int=0, bool rope_scaling=False, int old_context_len=8192, float scaling_factor=16, float lo_freq_factor=1, float hi_freq_factor=32,  Tensor? qparam_k=None, Tensor? qparam_v=None) -> Tensor");
+  m.def(
+      "rope_qkv_decoding(Tensor XQ, Tensor? XK, Tensor? XV, Tensor(a!) cache_K, Tensor(b!) cache_V,  Tensor seqpos, float theta, int? num_groups=1, Tensor? block_tables=None, int page_size=" STRING(
+          DEFAULT_PAGE_SIZE) ", Tensor? actual_batch_size=None, Tensor? batch=None, Tensor? cache_seqpos=None,  int cache_logical_dtype_int=0, bool rope_scaling=False, int old_context_len=8192, float scaling_factor=16, float lo_freq_factor=1, float hi_freq_factor=32, Tensor? qparam_k=None, Tensor? qparam_v=None, bool k_norm=False, bool update_kv=True, Tensor?amax_qkv=None) -> Tensor");
+  m.def(
+      "nope_qkv_varseq_prefill(Tensor XQ, Tensor? XK, Tensor? XV, Tensor(a!) cache_K, Tensor(b!) cache_V,  Tensor varseq_batch, Tensor varseq_seqpos, Tensor? block_tables=None, int page_size=" STRING(
+          DEFAULT_PAGE_SIZE) ", Tensor? actual_batch_size=None, Tensor? varseq_cache_seqpos=None, int cache_logical_dtype_int=0, int? num_groups=1, Tensor? qparam_k=None, Tensor? qparam_v=None, bool k_norm=False, bool update_kv=True, Tensor?amax_qkv=None, Tensor?kv_quant_scale_precomputed=None) -> Tensor");
+  m.def(
+      "nope_qkv_decoding(Tensor XQ, Tensor? XK, Tensor? XV, Tensor(a!) cache_K, Tensor(b!) cache_V,  Tensor seqpos, Tensor? block_tables=None, int page_size=" STRING(
+          DEFAULT_PAGE_SIZE) ", Tensor? actual_batch_size=None, Tensor? batch=None, Tensor? cache_seqpos=None, int cache_logical_dtype_int=0, int? num_groups=1, Tensor? qparam_k=None, Tensor? qparam_v=None, bool k_norm=False, bool update_kv=True, Tensor?amax_qkv=None) -> Tensor");
+  m.def(
+      "xpos_qkv_varseq_prefill(Tensor XQ, Tensor XK, Tensor XV, Tensor(a!) cache_K, Tensor(b!) cache_V, Tensor varseq_batch, Tensor varseq_seqpos, float theta, float gamma, float scale_base, float exponent_offset, int? num_groups=1, Tensor? block_tables=None, int page_size=" STRING(
+          DEFAULT_PAGE_SIZE) ", Tensor? actual_batch_size=None, Tensor? varseq_cache_seqpos=None, int cache_logical_dtype_int=0, bool rope_scaling=False, int old_context_len=8192, float scaling_factor=16, float lo_freq_factor=1, float hi_freq_factor=32,  Tensor? qparam_k=None, Tensor? qparam_v=None) -> Tensor");
+  m.def(
+      "xpos_qkv_decoding(Tensor XQ, Tensor XK, Tensor XV, Tensor(a!) cache_K, Tensor(b!) cache_V,  Tensor seqpos, float theta, float gamma, float scale_base, float exponent_offset, int? num_groups=1, Tensor? block_tables=None, int page_size=" STRING(
+          DEFAULT_PAGE_SIZE) ", Tensor? actual_batch_size=None, Tensor? batch=None, Tensor? cache_seqpos=None, int cache_logical_dtype_int=0, bool rope_scaling=False, int old_context_len=8192, float scaling_factor=16, float lo_freq_factor=1, float hi_freq_factor=32,  Tensor? qparam_k=None, Tensor? qparam_v=None) -> Tensor");
   m.def(
       "dequantize_int4_cache(Tensor cache_K, Tensor cache_V, Tensor kv_seqlen, int? num_groups=1, Tensor? qparam_k=None, Tensor? qparam_v=None) -> (Tensor, Tensor)");
   m.def(

fbgemm_gpu/experimental/gen_ai/src/quantize/ck_extensions/bf16_grouped/kernels/bf16_grouped_kernel_manifest.h

@@ -13,8 +13,7 @@
 
 #include <ATen/ATen.h>
 
-#define KERNEL_NAME_MAP_ENTRY(name) \
-  { #name, name }
+#define KERNEL_NAME_MAP_ENTRY(name) {#name, name}
 
 template <typename InputType, typename OutputType>
 using GroupedKernel =

fbgemm_gpu/experimental/gen_ai/src/quantize/quantize.cu

@@ -444,8 +444,9 @@ __global__ void scaleMatrix1(
     const int64_t lda) {
   for (int64_t i = threadIdx.x + blockIdx.x * blockDim.x; i < numel;
        i += (size_t)blockDim.x * gridDim.x) {
-    output[i] = T_OUT(scale<QUANTIZE>(
-        static_cast<float>(input[i]), static_cast<float>(input_scale[0])));
+    output[i] = T_OUT(
+        scale<QUANTIZE>(
+            static_cast<float>(input[i]), static_cast<float>(input_scale[0])));
   }
 }
 
@@ -534,9 +535,10 @@ __global__ void scaleMatrixRowwise1(
     const int64_t lda) {
   for (int64_t i = threadIdx.x + blockIdx.x * blockDim.x; i < numel;
        i += (size_t)blockDim.x * gridDim.x) {
-    output[i] = T_OUT(scale<QUANTIZE>(
-        static_cast<float>(input[i]),
-        static_cast<float>(input_scale[i / lda])));
+    output[i] = T_OUT(
+        scale<QUANTIZE>(
+            static_cast<float>(input[i]),
+            static_cast<float>(input_scale[i / lda])));
   }
 }
 
@@ -549,9 +551,10 @@ __global__ void scaleMatrixColwise(
     const int64_t lda) {
   for (int64_t i = threadIdx.x + blockIdx.x * blockDim.x; i < numel;
        i += (size_t)blockDim.x * gridDim.x) {
-    output[i] = T_OUT(scale<QUANTIZE>(
-        static_cast<float>(input[i]),
-        static_cast<float>(input_scale[i % lda])));
+    output[i] = T_OUT(
+        scale<QUANTIZE>(
+            static_cast<float>(input[i]),
+            static_cast<float>(input_scale[i % lda])));
   }
 }
 
@@ -1764,8 +1767,9 @@ __device__ __host__ __inline__ void compute_scale_with_global(
     const double two_level_scale =
         static_cast<double>(scale_format_max) * (elem_format_max / global_amax);
 
-    const double local_unscale_q = quantize_amax_e4m3(static_cast<float>(
-                                       local_unscale * two_level_scale)) /
+    const double local_unscale_q =
+        quantize_amax_e4m3(
+            static_cast<float>(local_unscale * two_level_scale)) /
         two_level_scale;
     double eps = FLT_MIN;
     scale = 1. / (local_unscale_q + eps);

fbgemm_gpu/experimental/hstu/src/hstu_ampere/hstu_fwd.h

@@ -564,13 +564,10 @@ inline __device__ void hstu_compute_attn_1rowblock(
   Tensor tOcO = gmem_thr_copy_O.partition_D(cO);
   Tensor tOpO = make_tensor<bool>(make_shape(size<2>(tOgO)));
   // Clear_OOB_K must be false since we don't want to write zeros to gmem
-  flash::
-      copy</*Is_even_MN=*/false, /*Clear_OOB_MN=*/false, /*Clear_OOB_K=*/false>(
-          gmem_tiled_copy_O,
-          tOrO,
-          tOgO,
-          tOcO,
-          actual_seqlen_q - m_block * kBlockM);
+  flash::copy</*Is_even_MN=*/false,
+              /*Clear_OOB_MN=*/false,
+              /*Clear_OOB_K=*/false>(
+      gmem_tiled_copy_O, tOrO, tOgO, tOcO, actual_seqlen_q - m_block * kBlockM);
 }
 
 template <typename Kernel_traits, typename Params>

fbgemm_gpu/experimental/hstu/src/hstu_hopper/epilogue_bwd_sm90_tma.hpp

@@ -155,11 +155,11 @@ struct CollectiveEpilogueBwd {
         tdKrdK_out); // ((Atom,AtomNum), MMA_M, MMA_N)
     Tensor taccdVrdV = smem_thr_copy_dKV.retile_S(
         tdVrdV_out); // ((Atom,AtomNum), MMA_M, MMA_N)
-    Tensor taccdKsdK =
-        smem_thr_copy_dKV.partition_D(cute::conditional_return<!dKV_swapAB>(
+    Tensor taccdKsdK = smem_thr_copy_dKV.partition_D(
+        cute::conditional_return<!dKV_swapAB>(
             sdK, sdKt)); // ((Atom,AtomNum),PIPE_M,PIPE_N)
-    Tensor taccdVsdV =
-        smem_thr_copy_dKV.partition_D(cute::conditional_return<!dKV_swapAB>(
+    Tensor taccdVsdV = smem_thr_copy_dKV.partition_D(
+        cute::conditional_return<!dKV_swapAB>(
             sdV, sdVt)); // ((Atom,AtomNum),PIPE_M,PIPE_N)
 
     // Make sure all WGs have finished reading K and V

fbgemm_gpu/experimental/hstu/src/hstu_hopper/hstu_bwd_kernel.h

@@ -408,12 +408,10 @@ __global__ void __launch_bounds__(
       using TileShape_MNK = typename Ktraits::TileShape_MNK;
       Tensor tdKrdK = partition_fragment_C(
           tiled_mma_dKV,
-          select < !dKV_swapAB ? 1 : 2,
-          !dKV_swapAB ? 2 : 1 > (TileShape_MNK{}));
+          select<!dKV_swapAB ? 1 : 2, !dKV_swapAB ? 2 : 1>(TileShape_MNK{}));
       Tensor tdVrdV = partition_fragment_C(
           tiled_mma_dKV,
-          select < !dKV_swapAB ? 1 : 2,
-          !dKV_swapAB ? 2 : 1 > (TileShape_MNK{}));
+          select<!dKV_swapAB ? 1 : 2, !dKV_swapAB ? 2 : 1>(TileShape_MNK{}));
       collective_mainloop.mma(
           mainloop_params,
           pipeline_q,

fbgemm_gpu/experimental/hstu/src/hstu_hopper/hstu_bwd_postprocess_kernel.h

@@ -254,8 +254,7 @@ class FlashAttnBwdPostprocessConvertdQ {
     TiledMma tiled_mma_dQ;
     Tensor taccdQrdQaccum = partition_fragment_C(
         tiled_mma_dQ,
-        select < !dQ_swapAB ? 0 : 1,
-        !dQ_swapAB ? 1 : 0 > (TileShape_MK{}));
+        select<!dQ_swapAB ? 0 : 1, !dQ_swapAB ? 1 : 0>(TileShape_MK{}));
     CUTE_STATIC_ASSERT_V(size(taccdQrdQaccum) == size(tdQsdQaccum));
     Tensor tdQrdQaccum = s2r_thr_copy_dQaccum.retile_D(taccdQrdQaccum);
     cute::copy(s2r_tiled_copy_dQaccum, tdQsdQaccum, tdQrdQaccum);