[QST] Performance overhead: Frequent cuTensorMapEncodeTiled and cudaGetDriverEntryPoint calls during Blackwell GEMM loops

While benchmarking GEMM kernels using the CUTLASS 3.x API on NVIDIA Blackwell architecture, I observed significant overhead between kernel launches. Profiling shows a high frequency of cudaGetDriverEntryPoint and cuTensorMapEncodeTiled calls.

gemm kernel: 91us
loop 1: 625us

code

template <typename T>
struct KernelConfigM128 {
    using OutputType = T;
    using MmaTileShape = Shape<_128, _256, _256>;
    using ClusterShape = Shape<int, int, _1>;
    using EpilogueTile = Shape<_128, _64>;  // Avoid register spilling
    using EpilogueSchedule = cutlass::epilogue::TmaWarpSpecialized1Sm;
    using MainloopSchedule = cutlass::gemm::KernelTmaWarpSpecialized1SmNvf4Sm100;
    const static dim3 preferred_cluster;
    const static dim3 fallback_cluster;
};
template <typename T>
const dim3 KernelConfigM128<T>::preferred_cluster(1, 4, 1);
template <typename T>
const dim3 KernelConfigM128<T>::fallback_cluster(1, 2, 1);

template <typename T>
struct KernelConfigM256 {
    using OutputType = T;
    using MmaTileShape = Shape<_256, _256, _256>;
    using ClusterShape = Shape<int, int, _1>;
    using EpilogueTile = Shape<_128, _64>;  // Avoid register spilling
    using EpilogueSchedule = cutlass::epilogue::TmaWarpSpecialized2Sm;
    using MainloopSchedule = cutlass::gemm::KernelTmaWarpSpecialized2SmNvf4Sm100;
    const static dim3 preferred_cluster;
    const static dim3 fallback_cluster;
};
template <typename T>
const dim3 KernelConfigM256<T>::preferred_cluster(2, 4, 1);
template <typename T>
const dim3 KernelConfigM256<T>::fallback_cluster(2, 1, 1);

template <typename T>
struct KernelConfigDefault {
    using OutputType = T;
    using MmaTileShape = Shape<_256, _256, _256>;
    using ClusterShape = Shape<int, int, _1>;
    using EpilogueTile = Shape<_128, _64>;  // Avoid register spilling
    using EpilogueSchedule = cutlass::epilogue::TmaWarpSpecialized2Sm;
    using MainloopSchedule = cutlass::gemm::KernelTmaWarpSpecialized2SmNvf4Sm100;
    const static dim3 preferred_cluster;
    const static dim3 fallback_cluster;
};
template <typename T>
const dim3 KernelConfigDefault<T>::preferred_cluster(4, 4, 1);
template <typename T>
const dim3 KernelConfigDefault<T>::fallback_cluster(2, 1, 1);

template <typename T, bool kEnableSilu, bool kEnableBias>
void runNvfp4Gemm(
    void* output,
    const void* input_a,
    const void* input_b,
    const void* scales_a,
    const void* scales_b,
    const void* bias,
    float alpha,
    int M,
    int N,
    int K,
    cudaStream_t stream) {
    typename T::Gemm gemm;
    auto arguments = args_from_options<T, kEnableSilu, kEnableBias>(
        output, input_a, input_b, scales_a, scales_b, bias, alpha, M, N, K);

    size_t workspace_size = T::Gemm::get_workspace_size(arguments);
    cutlass::device_memory::allocation<uint8_t> workspace(workspace_size);

    CUTLASS_CHECK(gemm.can_implement(arguments));
    CUTLASS_CHECK(gemm.initialize(arguments, workspace.get(), stream));

    CUTLASS_CHECK(gemm.run(arguments, workspace.get(), stream));
}

Environment
GPU: thor-u
CUTLASS Version: v4.3.5
CUDA Toolkit: 12.8
OS: Linux

Start with runNvfp4Gemm in the issue body and profile the calls around gemm.initialize and gemm.run, focusing on cudaGetDriverEntryPoint and cuTensorMapEncodeTiled. Compare the 91us kernel time with the 625us loop time in the stated CUTLASS v4.3.5, CUDA 12.8 and thor-u environment; done means identifying the source of the overhead and documenting a reproducible finding or next step.