What is PDL mean?

[MARD1NO]

I found there is a params bool launch_with_pdl = false in gemm universal adapter, what does it mean? It seems it is a Hopper feature

[thakkarV]

@alihassanijr

[alihassanijr]

PDL stands for Programmatic Dependent Launch.

It is a feature new to Hopper through which one kernel that is currently being executed can allow dependent kernels to launch ahead of its completion and ramp up while the current kernel ramps down. Current kernel can issue cudaTriggerProgrammaticLaunchCompletion() when it is about to ramp down and finish execution, and the next kernel can gate its conflicting memory accesses with cudaGridDependencySynchronize(), which just blocks execution until the first kernel flushes its memory.

PDL requires both a kernel launch attribute and those device instructions built into the kernel, but only the former currently exists in CUTLASS.

[MARD1NO]

And it seems cudaTriggerProgrammaticLaunchCompletion and cudaGridDependencySynchronize() (or its PTX) is not used in cutlass code currently, it may not work ?

Yes, I don't think any existing cutlass kernels have those instructions today. It is also very important to make sure all kernels launched with PDL and their next kernel have these instructions, especially cudaGridDependencySynchronize(), otherwise that can lead to incorrect results.

I test the pdl gemm exmaple kernel in my framework, when select a proper overlap ratio, I can observe the gemm kernel execution time is faster. But I do not observe any kernel overlap execution in nsys, does this meet the expectations?...

[alihassanijr]

I just wonder if primary kernel has no work that can coincide with the secondary kernel

No, it's not that the primary kernel doesn't do any work. Primary kernel typically issues cudaTriggerProgrammaticLaunchCompletion() when it's about to finish its work. For example if you look at the non-persistent warp specialized GEMM for hopper, you can see it's being called before the epilogue store, meaning the MMAs are done, and as it's about to write registers into gmem, it unblocks the next kernel from launching. However, that does not say anything about the next kernel occupying SMs.

I test the pdl gemm exmaple kernel in my framework, when select a proper overlap ratio, I can observe the gemm kernel execution time is faster. But I do not observe any kernel overlap execution in nsys, does this meet the expectations?...

Did you build CUTLASS with -DCUTLASS_ENABLE_GDC_FOR_SM90=1?

[MARD1NO]

I just wonder if primary kernel has no work that can coincide with the secondary kernel

No, it's not that the primary kernel doesn't do any work. Primary kernel typically issues cudaTriggerProgrammaticLaunchCompletion() when it's about to finish its work. For example if you look at the non-persistent warp specialized GEMM for hopper, you can see it's being called before the epilogue store, meaning the MMAs are done, and as it's about to write registers into gmem, it unblocks the next kernel from launching. However, that does not say anything about the next kernel occupying SMs.

I test the pdl gemm exmaple kernel in my framework, when select a proper overlap ratio, I can observe the gemm kernel execution time is faster. But I do not observe any kernel overlap execution in nsys, does this meet the expectations?...

Did you build CUTLASS with -DCUTLASS_ENABLE_GDC_FOR_SM90=1?

Oh, I think I misinterpreted the PDL and example 63_hopper_gemm_with_weight_prefetch.

In non-persistent warp specialized GEMM for hopper, you mean the mma warp finish the computation and can unblocks the next kernel, and the next kernel should add wait_on_dependent_grids when it need to consume previous kernel's output. there exist some opportunity to overlap the gemm epilogue and next kernel's perparation.

In 63_hopper_gemm_with_weight_prefetch, I think it is used in like RMSNorm + Gemm case, when RMSNorm Kernel still runing, the gemm kernel  can prefetch the weight into L2 Cache, I construct this case, and ensure the previous kernel use pdl launch and invoke cudaTriggerProgrammaticLaunchCompletion before write normed result to global memory, but I do not observe the overlap execution. But its execution time is faster than the origin kernel(it means I set overlap ratio, prefetch ratio as -1.0)

And I'm sure I build this kernel with -DCUTLASS_ENABLE_GDC_FOR_SM90=1 and gemm.run with launch_pdl=True.

Take a simple pdl demo as an example:

#include "cuda_runtime.h"


__global__ void copy1(float* __restrict__ a, float* __restrict__ b, 
                      int32_t n) {
    float tmp = 0.0f; 
                    
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n; i += blockDim.x * gridDim.x) {
        tmp += b[i];
    }

    cudaTriggerProgrammaticLaunchCompletion();

    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n; i += blockDim.x * gridDim.x) {
        b[i] = a[i] + tmp;
    }

}

__global__ void copy2(const float* __restrict__ b, float* __restrict__ c, 
                      const float* __restrict__ d, int32_t n) {
    float result = 0.0f; 
                        
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n; i += blockDim.x * gridDim.x) {
        result += d[i];
    }

    cudaGridDependencySynchronize();

    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n; i += blockDim.x * gridDim.x) {
        c[i] = b[i] + result + 2.0f;
    }

}


int main() {
    // const int n = 4096 * 4096; 
    const int n = 64 * 4096; 

    float *a, *b, *c, *d;
    cudaMalloc(&a, n * sizeof(float));
    cudaMalloc(&b, n * sizeof(float));
    cudaMalloc(&c, n * sizeof(float));
    cudaMalloc(&d, n * sizeof(float));

    const int32_t grid = 16; 
    const int32_t blockdim = 256; 


    cudaStream_t stream; 
    cudaStreamCreate(&stream);

    for (int i = 0; i < 10; i++) {
        cudaLaunchConfig_t kernelConfig2 = {0};                                                         
        kernelConfig2.gridDim = grid;   
        kernelConfig2.blockDim = blockdim;                                                                
        kernelConfig2.dynamicSmemBytes = 0;                                                   
        kernelConfig2.stream = stream;                                                                
        cudaLaunchAttribute attribute2[1];                                                            
        attribute2[0].id = cudaLaunchAttributeProgrammaticStreamSerialization;                        
        attribute2[0].val.programmaticStreamSerializationAllowed = 1;                                 
        kernelConfig2.attrs = attribute2;                                                              
        kernelConfig2.numAttrs = 1;  

        // PDL Launch
        // copy1<<<grid, blockdim, 0, stream>>>(a, b, n);  
        // cudaLaunchKernelEx(&kernelConfig2, &copy2, b, c, d, n);  

        // No PDL Launch
        copy1<<<grid, blockdim, 0, stream>>>(a, b, n);  
        copy2<<<grid, blockdim, 0, stream>>>(b, c, d, n);  
    }
    
    cudaDeviceSynchronize(); 

    cudaFree(a); 
    cudaFree(b); 
    cudaFree(c); 
    cudaFree(d);

    cudaStreamDestroy(stream);
}

When the elem_cnt is 4096 * 4096, I can observe kernel execution overlap. But in 64 * 4096 cases, I do not observe overlap, But the end2end time is faster with pdl launch:

PDL launch nsys: 40us

origin launch: 50us

Is this expected or I make some mistake?...

here is my nsys profile command:

nvcc -arch=sm_90a -O3 test3.cu -o test

export CUDA_VISIBLE_DEVICES=0
nsys profile --force-overwrite=true --cuda-graph-trace=node --stats=true \
    -o test_overlap ./test

[alihassanijr]

In non-persistent warp specialized GEMM for hopper, you mean the mma warp finish the computation and can unblocks the next kernel, and the next kernel should add wait_on_dependent_grids when it need to consume previous kernel's output. there exist some opportunity to overlap the gemm epilogue and next kernel's perparation.

Yes exactly.

In 63_hopper_gemm_with_weight_prefetch, I think it is used in like RMSNorm + Gemm case, when RMSNorm Kernel still runing, the gemm kernel can prefetch the weight into L2 Cache, I construct this case, and ensure the previous kernel use pdl launch and invoke cudaTriggerProgrammaticLaunchCompletion before write normed result to global memory, but I do not observe the overlap execution. But its execution time is faster than the origin kernel(it means I set overlap ratio, prefetch ratio as -1.0)

Yes that is the idea. Regarding runtime, I think there's a bug in the example. L2 cache lines should be cleared, otherwise as long as you prefetch once, the operands could very well become L2 resident and throw off the entire benchmark.

How much of a difference in effective BW / runtime are you observing?

Is this expected or I make some mistake?...

Sorry I'm confused, so if you are seeing overlap then what's the issue?

Keep in mind that the overlap is not in any way guaranteed, the same way that concurrency of launched kernels is not guaranteed.

[MARD1NO]

Thanks for your explanation, I think it is clear. :D