[pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

Author: pre-commit-ci[bot]

transformer_engine/common/cast/nvfp4/specialized/group_quantize_transpose_nvfp4_tuned_1D.cuh

@@ -459,7 +459,8 @@ __device__ __forceinline__ size_t get_grouped_scale_base_offset(
     const int64_t *const __restrict__ last_dims_ptr, const bool rowwise) {
   size_t scale_base = 0;
   for (size_t t = 0; t < tensor_id; ++t) {
-    const size_t rows = get_tensor_rows_num(t, shape_rep, first_logical_dim, first_dims_ptr, num_tensors);
+    const size_t rows =
+        get_tensor_rows_num(t, shape_rep, first_logical_dim, first_dims_ptr, num_tensors);
     const size_t cols = get_tensor_cols_num(t, shape_rep, last_logical_dim, last_dims_ptr);
 
     const size_t scale_rows = rowwise ? rows : cols;
@@ -500,7 +501,8 @@ __device__ __forceinline__ JobDescriptor decode_job(
   job.tensor_id = get_current_tensor_id(shape_rep, num_tensors, job.block_global_offset,
                                         static_cast<size_t>(ctaid_Y), first_logical_dim,
                                         last_logical_dim, offsets_ptr);
-  job.rows = get_tensor_rows_num(job.tensor_id, shape_rep, first_logical_dim, first_dims_ptr, num_tensors);
+  job.rows =
+      get_tensor_rows_num(job.tensor_id, shape_rep, first_logical_dim, first_dims_ptr, num_tensors);
   job.cols = get_tensor_cols_num(job.tensor_id, shape_rep, last_logical_dim, last_dims_ptr);
   return job;
 }
@@ -720,12 +722,14 @@ __global__ void __launch_bounds__(THREADS_NUM) group_quantize_transpose_nvfp4_tu
 
   constexpr int shmem_buff_size = buff_size_aligned_in / BUFFS_NUM;
 
-  const float S_enc_rowwise = (amax_rowwise_ptr == nullptr)
-                              ? 1.0f
-                              : core::compute_global_encode_scaling_factor_FP4(*amax_rowwise_ptr);
-  const float S_enc_colwise = (amax_colwise_ptr == nullptr)
-                              ? S_enc_rowwise
-                              : core::compute_global_encode_scaling_factor_FP4(*amax_colwise_ptr);
+  const float S_enc_rowwise =
+      (amax_rowwise_ptr == nullptr)
+          ? 1.0f
+          : core::compute_global_encode_scaling_factor_FP4(*amax_rowwise_ptr);
+  const float S_enc_colwise =
+      (amax_colwise_ptr == nullptr)
+          ? S_enc_rowwise
+          : core::compute_global_encode_scaling_factor_FP4(*amax_colwise_ptr);
 
   __shared__ size_t rowwise_scale_base[MAX_SUPPORTED_TENSOR_DESCRIPTORS + 1];
   __shared__ size_t colwise_scale_base[MAX_SUPPORTED_TENSOR_DESCRIPTORS + 1];
@@ -739,10 +743,11 @@ __global__ void __launch_bounds__(THREADS_NUM) group_quantize_transpose_nvfp4_tu
       size_t rowwise_scale_base_acc = 0;
       size_t colwise_scale_base_acc = 0;
 
-      for (size_t t = 0; t < num_tensors; ++t) { 
-        const size_t rows = get_tensor_rows_num(t, shape_rep, first_logical_dim, first_dims_ptr, num_tensors);
+      for (size_t t = 0; t < num_tensors; ++t) {
+        const size_t rows =
+            get_tensor_rows_num(t, shape_rep, first_logical_dim, first_dims_ptr, num_tensors);
         const size_t cols = get_tensor_cols_num(t, shape_rep, last_logical_dim, last_dims_ptr);
-    
+
         rowwise_scale_base_acc += rows * get_nvfp4_scale_stride(cols);
         colwise_scale_base_acc += cols * get_nvfp4_scale_stride(rows);
         rowwise_scale_base[t + 1] = rowwise_scale_base_acc;
@@ -856,9 +861,8 @@ __global__ void __launch_bounds__(THREADS_NUM) group_quantize_transpose_nvfp4_tu
     const size_t scale_stride_t = get_nvfp4_scale_stride(rows);
 
     nvfp4_scale_t *const scales_rowwise = scales_ptr + rowwise_scale_base[current_job.tensor_id];
-    nvfp4_scale_t *const scales_colwise = RETURN_TRANSPOSE
-                                          ? (scales_t_ptr + colwise_scale_base[current_job.tensor_id])
-                                          : nullptr;
+    nvfp4_scale_t *const scales_colwise =
+        RETURN_TRANSPOSE ? (scales_t_ptr + colwise_scale_base[current_job.tensor_id]) : nullptr;
 
     const CUtensorMap &tensor_map_input = g_tensor_maps_input[current_job.tensor_id];
     const CUtensorMap &tensor_map_output = g_tensor_maps_output[current_job.tensor_id];
@@ -1205,9 +1209,8 @@ inline void group_quantize_transpose(const GroupedTensor *input, const Tensor *n
             NVTE_CHECK_CUDA(cudaFuncSetAttribute(
                 kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, dshmem_size));
             kernel<<<grid, block_size, dshmem_size, stream>>>(
-                shape_rep,
-                num_tensors, first_logical_dim, last_logical_dim, offsets_ptr, first_dims_ptr,
-                last_dims_ptr, scales_ptr, scales_t_ptr, noop_ptr, amax_rowwise_ptr,
+                shape_rep, num_tensors, first_logical_dim, last_logical_dim, offsets_ptr,
+                first_dims_ptr, last_dims_ptr, scales_ptr, scales_t_ptr, noop_ptr, amax_rowwise_ptr,
                 amax_colwise_ptr, work_blocks_X, work_blocks_Y, rng_state);
             NVTE_CHECK_CUDA(cudaGetLastError());
           });););