performance evaluation using matmul.py

[zjin-lcf]

Hello

Running the matmul.py (https://github.com/spcl/dace/tree/main/samples/optimization) with the FP32 data types on an H100 GPU shows the following performance result. The GEMM performance seems very low.
Do users need to add additional settings to run the script for evaluating the performance of GEMM properly ? Thank you for your instructions.

Matrix multiplication 8192x8192x8192 (version: optimize_gpu)
1.970089 iter/s, 101.518 sec, 2.2 TFLOPS
Difference: 2.9815700486324204e-07

The related code is

        for _ in range(num_iter):   # num_iter = 200
             sdfg(A=A, B=B, C=C, M=m, N=n, K=k)

[tbennun]

Hi,

Are you benchmarking based on our instrumentation interface? If not, you may be measuring the Python overhead of calling the matrix multiplication too. See this Jupyter notebook for more information:
https://nbviewer.jupyter.org/github/spcl/dace/blob/main/tutorials/benchmarking.ipynb

[zjin-lcf]

Thank you for the comments and links. The example in matmul.py says "Sample showcasing transformations applied on a naive matrix multiplication program, yielding performance that competes with Intel MKL and NVIDIA CUBLAS.". It will be helpful to add your interface to the example to show the performance.

[zjin-lcf]

@tbennun

I listed the CUDA code generated by dace below after I tried to instrument the Python program (matmul.py). I have a few questions. How can a user instrument the Python program to measure only the matmul kernel time ?

In addition, I observe that verification result is "inf" when the data type for A, B, and C is float16 in the example. Can you reproduce the result ?

Thank you

    {

        __dace_runkernel_freduce_init_map_0_1_0(__state, gpu_C, M, N);


    }
    {

        auto __dace_tbegin_0_0 = std::chrono::high_resolution_clock::now();
        __dace_runkernel_matmul_40_0_0_4(__state, gpu_A, gpu_B, gpu_C, K, M, N);
        DACE_GPU_CHECK(cudaMemcpyAsync(C, gpu_C, (M * N) * sizeof(dace::float16), cudaMemcpyDeviceToHost, __state->gpu_context->streams[0]));
        DACE_GPU_CHECK(cudaStreamSynchronize(__state->gpu_context->streams[0]));
        auto __dace_tend_0_0 = std::chrono::high_resolution_clock::now();
        ...
    }