Skip to content

Commit c6517dd

Browse files
[pre-commit.ci] auto fixes from pre-commit.com hooks
for more information, see https://pre-commit.ci
1 parent fb94305 commit c6517dd

1 file changed

Lines changed: 7 additions & 9 deletions

File tree

‎transformer_engine/common/comm_gemm/comm_gemm.cpp‎

Lines changed: 7 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -376,10 +376,8 @@ void cublasmp_gemm(InitMatricesFn init_matrices_fn, NVTECommGemmCtx* ctx, NVTECo
376376
sizeof algo_attr));
377377

378378
const cublasMpMatmulMatrixScale_t scale_mode = CUBLASMP_MATMUL_MATRIX_SCALE_SCALAR_FP32;
379-
auto set_input_scale = [&](const Tensor& t,
380-
cublasMpMatmulDescriptorAttribute_t scale_mode_attr,
381-
cublasMpMatmulDescriptorAttribute_t scale_ptr_attr,
382-
const char* side) {
379+
auto set_input_scale = [&](const Tensor& t, cublasMpMatmulDescriptorAttribute_t scale_mode_attr,
380+
cublasMpMatmulDescriptorAttribute_t scale_ptr_attr, const char* side) {
383381
cublasMpMatmulMatrixScale_t input_scale_mode = scale_mode;
384382
if (is_mxfp_scaling(t.scaling_mode)) {
385383
#if CUBLASMP_VERSION >= 801
@@ -400,8 +398,8 @@ void cublasmp_gemm(InitMatricesFn init_matrices_fn, NVTECommGemmCtx* ctx, NVTECo
400398
NVTE_CHECK(t.scale_inv.dptr, "Scaling must be set for input ", side);
401399
NVTE_CHECK_CUBLASMP(cublasMpMatmulDescriptorSetAttribute(
402400
ctx->matmul_desc.get(), scale_mode_attr, &input_scale_mode, sizeof input_scale_mode));
403-
NVTE_CHECK_CUBLASMP(cublasMpMatmulDescriptorSetAttribute(
404-
ctx->matmul_desc.get(), scale_ptr_attr, &t.scale_inv.dptr, sizeof(void*)));
401+
NVTE_CHECK_CUBLASMP(cublasMpMatmulDescriptorSetAttribute(ctx->matmul_desc.get(), scale_ptr_attr,
402+
&t.scale_inv.dptr, sizeof(void*)));
405403
};
406404
set_input_scale(a_used, CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_A_SCALE_MODE,
407405
CUBLASMP_MATMUL_DESCRIPTOR_ATTRIBUTE_A_SCALE_POINTER, "A");
@@ -553,9 +551,9 @@ void cublasmp_gemm(InitMatricesFn init_matrices_fn, NVTECommGemmCtx* ctx, NVTECo
553551

554552
if (update_nvfp4_alpha) {
555553
NVTE_CHECK_CUDA(cudaStreamSynchronize(main_stream));
556-
alpha *= (nvfp4_amax_a * nvfp4_amax_b) /
557-
(6.0f * 6.0f * static_cast<float>(a_used.nvfp4_e4m3_max) *
558-
static_cast<float>(b_used.nvfp4_e4m3_max));
554+
alpha *=
555+
(nvfp4_amax_a * nvfp4_amax_b) / (6.0f * 6.0f * static_cast<float>(a_used.nvfp4_e4m3_max) *
556+
static_cast<float>(b_used.nvfp4_e4m3_max));
559557
}
560558

561559
NVTE_CHECK_CUBLASMP(

0 commit comments

Comments
 (0)