[Kernel] Add PDL support for DeepGEMM kernel

[jeejeelee]

Purpose

Test Plan

Test Result

---

Essential Elements of an Effective PR Description Checklist

• The purpose of the PR, such as "Fix some issue (link existing issues this PR will resolve)".
• The test plan, such as providing test command.
• The test results, such as pasting the results comparison before and after, or e2e results
• (Optional) The necessary documentation update, such as updating supported_models.md and examples for a new model.

[gemini-code-assist]

Code Review

This pull request implements support for Programmatic Dependent Launch (PDL) and Grid Dependency Control (GDC) for DeepGEMM kernels on Hopper and Blackwell architectures. It introduces new environment variables, configuration utilities, and integrates GDC primitives into CUDA and Triton kernels. Review feedback identified critical issues: a race condition in the CUDA quantization kernel caused by signaling completion in an early exit block, incorrect argument packing for the cudaLaunchKernelEx API, and a logic error in the Triton kernel where completion is signaled before the work starts.

[gemini-code-assist]

The call to griddepcontrol.launch_dependents within this early exit block is problematic because mn_idx is not uniform across the CTA (it varies with row_local). If some threads in a block exit early and signal completion while others are still performing quantization and writing to global memory, it creates a race condition where dependent kernels in the stream may start reading incomplete data. Since the grid launch logic ensures that at least one thread in every block is in bounds (blockIdx.y * ry < tma_aligned_mn), the signal at the end of the kernel (line 405) will correctly handle the CTA-level completion signal. The early exit should simply return without signaling.

  if (mn_idx >= tma_aligned_mn) {
    return;
  }

[gemini-code-assist]

cudaLaunchKernelEx is a C API that expects an array of pointers to kernel arguments (void** args) as its third parameter. Passing the arguments directly as variadic parameters will result in a compilation error. You must pack the arguments into a void* array and pass that to the function.

    cudaLaunchConfig_t config = {};                                          \
    config.gridDim = dim3(static_cast<unsigned int>(blocks_x),               \
                          static_cast<unsigned int>(blocks_y));              \
    config.blockDim = dim3(num_threads);                                     \
    config.dynamicSmemBytes = 0;                                             \
    config.stream = stream;                                                  \
    cudaLaunchAttribute attrs[1];                                            \
    attrs[0].id = cudaLaunchAttributeProgrammaticStreamSerialization;        \
    attrs[0].val.programmaticStreamSerializationAllowed = 1;                 \
    config.numAttrs = 1;                                                     \
    config.attrs = attrs;                                                    \
    const T* input_ptr = static_cast<const T*>(input.data_ptr());            \
    void* output_q_ptr = output_q.data_ptr();                                \
    unsigned int* output_s_ptr = reinterpret_cast<unsigned int*>(output_s_packed.data_ptr()); \
    int p_groups = static_cast<int>(padded_groups_per_row);                  \
    int groups = static_cast<int>(groups_per_row);                           \
    int mn_val = static_cast<int>(mn);                                       \
    int q_extent = static_cast<int>(output_q_mn_extent);                     \
    int tma_mn = static_cast<int>(tma_aligned_mn);                           \
    float eps_f = static_cast<float>(eps);                                   \
    float min_f = static_cast<float>(min_8bit);                              \
    float max_f = static_cast<float>(max_8bit);                              \
    void* kernel_args[] = {                                                  \
        &input_ptr, &output_q_ptr, &output_s_ptr, &p_groups,                 \
        &groups, &mn_val, &q_extent, &tma_mn, &num_scale_elems,              \
        &eps_f, &min_f, &max_f                                               \
    };                                                                       \
    cudaLaunchKernelEx(                                                      \
        &config,                                                             \
        (const void*)per_token_group_quant_8bit_packed_register_kernel<T, DST_DTYPE, 128, \
                                                          KX, RY>,           \
        kernel_args);

[mergify]

This pull request has merge conflicts that must be resolved before it can be
merged. Please rebase the PR, @jeejeelee.

https://docs.github.com/en/pull-requests/collaborating-with-pull-requests/working-with-forks/syncing-a-fork

[gau-nernst]

Shall we only enable PDL for the imported deepgemm module? Instead of all DeepGEMM modules. Something like

if envs.VLLM_USE_DEEP_GEMM_PDL and current_platform.is_arch_support_pdl():
    _dg.set_pdl(True)

Btw, do we even need VLLM_USE_DEEP_GEMM_PDL flag? I think it's safe to always enable it?