[QST] (C7504) & (C7509) Potential Performance Loss: about 'setmaxnreg' and wgmma.mma_async

[chenhongyu2048]

What is your question?

I wrote a kernel based on cute and cutlass, which used cutlass::arch::warpgroup_reg_alloc and wgmma.mma_async.
I wrote the kernel in a kernel.cuh file and included the kernel.cuh in src1.cu. Since my project contains multiple cu files, such as src2.cu, src3.cu. I need to use rdc=True to compile my project.
However, the following warning appeared during the compilation process.

ptxas info    : (C7504) Potential Performance Loss: 'setmaxnreg' ignored to maintain compatibility across compilation units.
ptxas info    : (C7509) Potential Performance Loss: wgmma.mma_async instructions are serialized due to the presence of Extern calls in the function 

I can be sure that I have used the -DNDEBUG flag. I'm using cuda12.8 on H20.

[chenhongyu2048]

The code is as below.
If I comment out this part of the code, the above warning will not appear. Or if I write this part of the code directly in a src.cu file and only compile the src.cu file, no warning will be generated, too.

#pragma once

#include <iostream>
#include <fstream>
#include <sstream>
#include <vector>
#include <cfloat>
#include <cuda_runtime.h>

#include "cutlass/cutlass.h"

#include "cute/tensor.hpp"
#include "cutlass/tensor_ref.h"
#include "cutlass/epilogue/collective/default_epilogue.hpp"
#include "cutlass/epilogue/thread/linear_combination.h"
#include "cutlass/gemm/dispatch_policy.hpp"
#include "cutlass/gemm/group_array_problem_shape.hpp"
#include "cutlass/gemm/collective/collective_builder.hpp"
#include "cutlass/epilogue/collective/collective_builder.hpp"
#include "cutlass/gemm/device/gemm_universal_adapter.h"
#include "cutlass/gemm/kernel/gemm_universal.hpp"

#include "cutlass/util/command_line.h"
#include "cutlass/util/distribution.h"
#include "cutlass/util/host_tensor.h"
#include "cutlass/util/packed_stride.hpp"
#include "cutlass/util/tensor_view_io.h"
#include "cutlass/util/reference/device/gemm.h"
#include "cutlass/util/reference/device/tensor_compare.h"
#include "cutlass/util/reference/device/tensor_fill.h"

#include "cuda_utils.cuh"

using namespace cute;
using ProblemShape = cutlass::gemm::GroupProblemShape<Shape<int, int, int>>;
using ElementAccumulator = float;                                 // Element type for internal accumulation
using ArchTag = cutlass::arch::Sm90;                              // Tag indicating the minimum SM that supports the intended feature
using OperatorClass = cutlass::arch::OpClassTensorOp;             // Operator class tag
using StageCountType = cutlass::gemm::collective::StageCountAuto; // Stage count maximized based on the tile size
using StrideType = cute::Stride<int64_t, cute::_1, cute::_0>;

template <int tile_shape_m, int tile_shape_n, int tile_shape_k, int cluster_shape_m, int cluster_shape_n, int cluster_shape_k>
struct PingpongConfig
{
    using KernelSchedule = cutlass::gemm::KernelPtrArrayTmaWarpSpecializedPingpong;
    using EpilogueSchedule = cutlass::epilogue::PtrArrayTmaWarpSpecializedPingpong;
    using TileShape = Shape<cutlass::Int<tile_shape_m>, cutlass::Int<tile_shape_n>, cutlass::Int<tile_shape_k>>;
    using ClusterShape = Shape<cutlass::Int<cluster_shape_m>, cutlass::Int<cluster_shape_n>, cutlass::Int<cluster_shape_k>>;
};

template <typename ScheduleConfig,
          typename ElementA, typename LayoutA, int AlignmentA,
          typename ElementB, typename LayoutB, int AlignmentB,
          typename ElementC, typename LayoutC, int AlignmentC>
struct GemmGivenSchedule
{
    using TileShape = typename ScheduleConfig::TileShape;               // Threadblock-level tile size
    using ClusterShape = typename ScheduleConfig::ClusterShape;         // Shape of the threadblocks in a cluster
    using KernelSchedule = typename ScheduleConfig::KernelSchedule;     // Kernel to launch
    using EpilogueSchedule = typename ScheduleConfig::EpilogueSchedule; // Epilogue to launch

    using CollectiveEpilogue = typename cutlass::epilogue::collective::CollectiveBuilder<
        cutlass::arch::Sm90, cutlass::arch::OpClassTensorOp,
        TileShape, ClusterShape,
        cutlass::epilogue::collective::EpilogueTileAuto,
        ElementAccumulator, ElementAccumulator,
        ElementC, LayoutC *, AlignmentC,
        ElementC, LayoutC *, AlignmentC,
        EpilogueSchedule,
        cutlass::epilogue::fusion::LinearCombination<ElementC, ElementAccumulator>>::CollectiveOp;

    using CollectiveMainloop = typename cutlass::gemm::collective::CollectiveBuilder<
        ArchTag, OperatorClass,
        ElementA, LayoutA *, AlignmentA,
        ElementB, LayoutB *, AlignmentB,
        ElementAccumulator,
        TileShape, ClusterShape,
        cutlass::gemm::collective::StageCountAutoCarveout<
            static_cast<int>(sizeof(typename CollectiveEpilogue::SharedStorage))>,
        KernelSchedule>::CollectiveOp;

    using GemmKernel = cutlass::gemm::kernel::GemmUniversal<
        ProblemShape,
        CollectiveMainloop,
        CollectiveEpilogue>;

    using Gemm = cutlass::gemm::device::GemmUniversalAdapter<GemmKernel>;
};

template <typename T, int tile_shape_m, int tile_shape_n, int tile_shape_k, int cluster_shape_m, int cluster_shape_n, int cluster_shape_k>
void grouped_gemm_kernel(int num_groups, int M, int* Ns, int K, T* A_storage, T* B_storage, T* C_storage, float alpha, float beta, int device_id, cudaStream_t stream)
{
    // The actual implementation of the grouped GEMM kernel would go here.
    // std::cout << "Grouped GEMM kernel function called." << std::endl;

    // Set gemm instance
    using Element = T;
    using LayoutA = cutlass::layout::RowMajor;
    using LayoutB = cutlass::layout::ColumnMajor;
    using LayoutC = cutlass::layout::RowMajor;
    constexpr int Alignment = 128 / cutlass::sizeof_bits<Element>::value;
    using GemmGivenScheduleT = GemmGivenSchedule<PingpongConfig<tile_shape_m, tile_shape_n, tile_shape_k, cluster_shape_m, cluster_shape_n, cluster_shape_k>,
                                    Element, LayoutA, Alignment, Element, LayoutB, Alignment, Element, LayoutC, Alignment>;
    using GemmKernelPingpong = typename GemmGivenScheduleT::GemmKernel;
    using GemmPingpong = typename GemmGivenScheduleT::Gemm;
    using GemmT = GemmPingpong;
    GemmT gemm;

    // Set arguments for the GEMM operation
    cutlass::KernelHardwareInfo kernel_hw_info = cutlass::KernelHardwareInfo::make_kernel_hardware_info<typename GemmT::GemmKernel>(device_id);
    kernel_hw_info.sm_count -= 2;

    typename GemmT::Arguments arguments;
    decltype(arguments.epilogue.thread) fusion_args;

    // Single alpha and beta for all groups
    fusion_args.alpha = alpha;
    fusion_args.beta = beta;
    fusion_args.alpha_ptr = nullptr;
    fusion_args.beta_ptr = nullptr;
    fusion_args.alpha_ptr_array = nullptr;
    fusion_args.beta_ptr_array = nullptr;
    fusion_args.dAlpha = {cute::_0{}, cute::_0{}, 0};
    fusion_args.dBeta = {cute::_0{}, cute::_0{}, 0};

    size_t tmpsize = 3 * (num_groups * sizeof(T*) + num_groups * sizeof(StrideType)) + num_groups * sizeof(typename ProblemShape::UnderlyingProblemShape);
    void* tmp_ptr = nullptr;
    cudaMallocAsync(&tmp_ptr, tmpsize, stream);
    T** ptr_A;
    T** ptr_B;
    T** ptr_C;
    StrideType * stride_A;
    StrideType * stride_B;
    StrideType * stride_C;
    typename ProblemShape::UnderlyingProblemShape* problem_sizes;
    ptr_A = (T**)tmp_ptr;
    ptr_B = (T**)(static_cast<char*>(tmp_ptr) + num_groups * sizeof(T*));
    ptr_C = (T**)(static_cast<char*>(tmp_ptr) + 2 * num_groups * sizeof(T*));
    stride_A = (StrideType*)(static_cast<char*>(tmp_ptr) + 3 * num_groups * sizeof(T*));
    stride_B = (StrideType*)(static_cast<char*>(tmp_ptr) + 3 * num_groups * sizeof(T*) + num_groups * sizeof(StrideType));
    stride_C = (StrideType*)(static_cast<char*>(tmp_ptr) + 3 * num_groups * sizeof(T*) + 2 * num_groups * sizeof(StrideType));
    problem_sizes = (typename ProblemShape::UnderlyingProblemShape*)(static_cast<char*>(tmp_ptr) + 3 * num_groups * sizeof(T*) + 3 * num_groups * sizeof(StrideType));
    
    precompute_grouped_gemm_arguments<T>(
        stride_A, stride_B, stride_C,
        ptr_A, ptr_B, ptr_C, problem_sizes,
        A_storage, B_storage, C_storage,
        M, Ns, K, num_groups, stream);

    arguments = typename GemmT::Arguments{
        cutlass::gemm::GemmUniversalMode::kGrouped,
        {num_groups, problem_sizes, nullptr},
        {(const T**)(ptr_A), stride_A, (const T**)(ptr_B), stride_B},
        {fusion_args, (const T**)(ptr_C), stride_C, ptr_C, stride_C},
        kernel_hw_info};

    CUTLASS_CHECK(gemm.can_implement(arguments));
    size_t workspace_size = GemmT::get_workspace_size(arguments);
    void* workspace_ptr = nullptr;
    if (workspace_size > 0) {
        cudaMallocAsync(&workspace_ptr, workspace_size, stream);
    }
    // printf("Workspace size: %zu bytes\n", workspace_size);
    CUTLASS_CHECK(gemm.initialize(arguments, workspace_ptr));
    CUTLASS_CHECK(gemm.run(stream));
}

/*
M = output_dim, K = input_dim, Ns = batch sizes for each group
*/
template <typename T>
void dispatch_grouped_gemm_kernel(int num_groups, int M, int* Ns, int K, T* A_storage, T* B_storage, T* C_storage, float alpha, float beta, int device_id, cudaStream_t stream)
{
    grouped_gemm_kernel<T, 128, 32, 128, 2, 1, 1>(num_groups, M, Ns, K, A_storage, B_storage, C_storage, alpha, beta, device_id, stream);
}

[thakkarV]

make sure you are using the CUTLASS build flags are not rolling your own.

[github-actions]

This issue has been labeled inactive-30d due to no recent activity in the past 30 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed. This issue will be labeled inactive-90d if there is no activity in the next 60 days.