Benchmarks: Micro benchmark - add nvbench based kernel-launch & sleep-kernel #750

Copilot · 2026-01-23T00:05:47Z

Using @latest for third-party GitHub Actions is a supply-chain risk and can lead to non-reproducible CI behavior. Pin this action to a specific tagged version or commit SHA.

Suggested change

uses: lukka/get-cmake@latest

uses: lukka/get-cmake@v3.20.0

Copilot · 2026-01-23T00:05:46Z

Collapsing __all__ to long comma-separated lines reduces readability and likely violates typical line-length formatting used elsewhere in the project. Consider reverting to one-entry-per-line (or a more structured wrap) to keep diffs smaller and maintenance easier.

Suggested change

'BlasLtBaseBenchmark', 'ComputationCommunicationOverlap', 'CpuMemBwLatencyBenchmark', 'CpuHplBenchmark',

'CpuStreamBenchmark', 'CublasBenchmark', 'CublasLtBenchmark', 'CudaGemmFlopsBenchmark', 'CudaMemBwBenchmark',

'CudaNcclBwBenchmark', 'CudnnBenchmark', 'DiskBenchmark', 'DistInference', 'HipBlasLtBenchmark', 'GPCNetBenchmark',

'GemmFlopsBenchmark', 'GpuBurnBenchmark', 'GpuCopyBwBenchmark', 'GpuStreamBenchmark', 'IBBenchmark',

'IBLoopbackBenchmark', 'KernelLaunch', 'MemBwBenchmark', 'MicroBenchmark', 'MicroBenchmarkWithInvoke',

'ORTInferenceBenchmark', 'RocmGemmFlopsBenchmark', 'RocmMemBwBenchmark', 'ShardingMatmul',

'TCPConnectivityBenchmark', 'TensorRTInferenceBenchmark', 'DirectXGPUEncodingLatency', 'DirectXGPUCopyBw',

'DirectXGPUMemBw', 'DirectXGPUCoreFlops', 'NvBandwidthBenchmark', 'NvbenchKernelLaunch', 'NvbenchSleepKernel'

'BlasLtBaseBenchmark',

'ComputationCommunicationOverlap',

'CpuMemBwLatencyBenchmark',

'CpuHplBenchmark',

'CpuStreamBenchmark',

'CublasBenchmark',

'CublasLtBenchmark',

'CudaGemmFlopsBenchmark',

'CudaMemBwBenchmark',

'CudaNcclBwBenchmark',

'CudnnBenchmark',

'DiskBenchmark',

'DistInference',

'HipBlasLtBenchmark',

'GPCNetBenchmark',

'GemmFlopsBenchmark',

'GpuBurnBenchmark',

'GpuCopyBwBenchmark',

'GpuStreamBenchmark',

'IBBenchmark',

'IBLoopbackBenchmark',

'KernelLaunch',

'MemBwBenchmark',

'MicroBenchmark',

'MicroBenchmarkWithInvoke',

'ORTInferenceBenchmark',

'RocmGemmFlopsBenchmark',

'RocmMemBwBenchmark',

'ShardingMatmul',

'TCPConnectivityBenchmark',

'TensorRTInferenceBenchmark',

'DirectXGPUEncodingLatency',

'DirectXGPUCopyBw',

'DirectXGPUMemBw',

'DirectXGPUCoreFlops',

'NvBandwidthBenchmark',

'NvbenchKernelLaunch',

'NvbenchSleepKernel',

-Original file line number
+Diff line change
@@ Expand Up / @@ -151,6 +151,9 @@ cython_debug/ @@
     *.userosscache
     *.sln.docstates
+    # Build temporary files
+    compile_commands.json
     # Build results
     [Dd]ebug/
     [Dd]ebugPublic/
@@ Expand Down @@

-Original file line number
+Diff line change
@@ Expand Up / @@ -33,3 +33,6 @@ @@
     [submodule "third_party/nvbandwidth"]
     	path = third_party/nvbandwidth
     	url = https://github.com/NVIDIA/nvbandwidth.git
+    [submodule "third_party/nvbench"]
+    	path = third_party/nvbench
+    	url = https://github.com/NVIDIA/nvbench.git

-Original file line number
+Diff line change
@@ Expand Up / @@ -61,6 +61,27 @@ RUN apt-get update && \ @@
         apt-get clean && \
         rm -rf /var/lib/apt/lists/* /tmp/*
+    # Install CMake 3.30.4 for nvbench compatibility
+    RUN apt-get update && \
+        apt-get remove -y cmake cmake-data && \
+        apt-get autoremove -y && \
+        cd /tmp && \
+        ARCH=$(uname -m) && \
+        case ${ARCH} in \
+            "aarch64") CMAKE_ARCH="aarch64" ;; \
+            "x86_64") CMAKE_ARCH="x86_64" ;; \
+            "arm64") CMAKE_ARCH="aarch64" ;; \
+            *) CMAKE_ARCH="x86_64" ;; \
+        esac && \
+        echo "Detected architecture: ${ARCH}, using CMAKE_ARCH: ${CMAKE_ARCH}" && \
+        wget -q https://github.com/Kitware/CMake/releases/download/v3.30.4/cmake-3.30.4-linux-${CMAKE_ARCH}.tar.gz && \
+        tar -xzf cmake-3.30.4-linux-${CMAKE_ARCH}.tar.gz && \
+        mv cmake-3.30.4-linux-${CMAKE_ARCH} /opt/cmake && \
+        ln -sf /opt/cmake/bin/* /usr/local/bin/ && \
+        rm -rf cmake-3.30.4-linux-${CMAKE_ARCH}* && \
+        apt-get clean && \
+        rm -rf /var/lib/apt/lists/*
     ARG NUM_MAKE_JOBS=
     ARG TARGETPLATFORM
     ARG TARGETARCH
@@ Expand Down Expand Up / @@ -161,7 +182,7 @@ ADD dockerfile/etc /opt/microsoft/ @@
     WORKDIR ${SB_HOME}
     ADD third_party third_party
-    RUN make -C third_party cuda_with_msccl
+    RUN make -C third_party cuda_with_msccl cuda_nvbench
     ADD . .
     RUN python3 -m pip install --upgrade setuptools==70.3.0 && \
@@ Expand Down @@

-Original file line number
+Diff line change
@@ Expand Up / @@ -62,6 +62,28 @@ RUN apt-get update && \ @@
         apt-get clean && \
         rm -rf /var/lib/apt/lists/* /tmp/*
+    # Install CMake 3.30.4 for nvbench compatibility
+    RUN apt-get update && \
+        apt-get remove -y cmake cmake-data && \
+        apt-get autoremove -y && \
+        cd /tmp && \
+        ARCH=$(uname -m) && \
+        case ${ARCH} in \
+            "aarch64") CMAKE_ARCH="aarch64" ;; \
+            "x86_64") CMAKE_ARCH="x86_64" ;; \
+            "arm64") CMAKE_ARCH="aarch64" ;; \
+            *) CMAKE_ARCH="x86_64" ;; \
+        esac && \
+        echo "Detected architecture: ${ARCH}, using CMAKE_ARCH: ${CMAKE_ARCH}" && \
+        wget -q https://github.com/Kitware/CMake/releases/download/v3.30.4/cmake-3.30.4-linux-${CMAKE_ARCH}.tar.gz && \
+        tar -xzf cmake-3.30.4-linux-${CMAKE_ARCH}.tar.gz && \
+        mv cmake-3.30.4-linux-${CMAKE_ARCH} /opt/cmake && \
+        ln -sf /opt/cmake/bin/* /usr/local/bin/ && \
+        rm -rf cmake-3.30.4-linux-${CMAKE_ARCH}* && \
+        apt-get clean && \
+        rm -rf /var/lib/apt/lists/*
     ARG NUM_MAKE_JOBS=
     ARG TARGETPLATFORM
     ARG TARGETARCH
@@ Expand Down Expand Up / @@ -162,7 +184,7 @@ ADD dockerfile/etc /opt/microsoft/ @@
     WORKDIR ${SB_HOME}
     ADD third_party third_party
-    RUN make -C third_party cuda_with_msccl
+    RUN make -C third_party cuda_with_msccl cuda_nvbench
     ADD . .
     RUN python3 -m pip install --upgrade setuptools==78.1.0 && \
@@ Expand Down @@

-Original file line number
+Diff line change
@@ Expand Up / @@ -62,6 +62,27 @@ RUN apt-get update && \ @@
         apt-get clean && \
         rm -rf /var/lib/apt/lists/* /tmp/*
+    # Install CMake 3.30.4 for nvbench compatibility
+    RUN apt-get update && \
+        apt-get remove -y cmake cmake-data && \
+        apt-get autoremove -y && \
+        cd /tmp && \
+        ARCH=$(uname -m) && \
+        case ${ARCH} in \
+            "aarch64") CMAKE_ARCH="aarch64" ;; \
+            "x86_64") CMAKE_ARCH="x86_64" ;; \
+            "arm64") CMAKE_ARCH="aarch64" ;; \
+            *) CMAKE_ARCH="x86_64" ;; \
+        esac && \
+        echo "Detected architecture: ${ARCH}, using CMAKE_ARCH: ${CMAKE_ARCH}" && \
+        wget -q https://github.com/Kitware/CMake/releases/download/v3.30.4/cmake-3.30.4-linux-${CMAKE_ARCH}.tar.gz && \
+        tar -xzf cmake-3.30.4-linux-${CMAKE_ARCH}.tar.gz && \
+        mv cmake-3.30.4-linux-${CMAKE_ARCH} /opt/cmake && \
+        ln -sf /opt/cmake/bin/* /usr/local/bin/ && \
+        rm -rf cmake-3.30.4-linux-${CMAKE_ARCH}* && \
+        apt-get clean && \
+        rm -rf /var/lib/apt/lists/*
     ARG NUM_MAKE_JOBS=
     ARG TARGETPLATFORM
     ARG TARGETARCH
@@ Expand Down Expand Up / @@ -151,7 +172,7 @@ ADD dockerfile/etc /opt/microsoft/ @@
     WORKDIR ${SB_HOME}
     ADD third_party third_party
-    RUN make -C third_party cuda
+    RUN make -C third_party cuda cuda_nvbench
     ADD . .
     RUN python3 -m pip install --upgrade setuptools==78.1.0 && \
@@ Expand Down @@

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Benchmarks: Micro benchmark - add nvbench based kernel-launch & sleep-kernel #750

Uh oh!

Diff view

Diff view

There are no files selected for viewing

Copilot AI Jan 23, 2026

Uh oh!

Copilot AI Jan 23, 2026

Uh oh!

Uh oh!

Uh oh!

-Original file line number
+Diff line change
@@ Expand Up / @@ -98,7 +98,7 @@ RUN cd /tmp && \ @@
     # Install Intel MLC
     RUN cd /tmp && \
-        wget -q https://downloadmirror.intel.com/793041/mlc_v3.11.tgz -O mlc.tgz && \
+        wget -q https://downloadmirror.intel.com/866182/mlc_v3.12.tgz -O mlc.tgz && \
         tar xzf mlc.tgz Linux/mlc && \
         cp ./Linux/mlc /usr/local/bin/ && \
         rm -rf ./Linux mlc.tgz
@@ Expand Down @@

-Original file line number
+Diff line change
@@ Expand Up @@
     | gpu-burn/gpu_[0-9]_pass | yes/no   | The result of the gpu-burn test for each GPU (1: yes, 0: no).                      |
     | gpu-burn/abort          | yes/no   | Whether or not GPU-burn test aborted before returning GPU results (1: yes, 0: no). |
+    ### `nvbench-sleep-kernel`
+    #### Introduction
+    Measure GPU kernel execution time using NVBench's sleep kernel benchmark. This benchmark creates CUDA kernels that sleep for specified durations (in microseconds) and measures the actual execution time, providing insights into GPU scheduling overhead and timing accuracy.
+    The benchmark supports multiple duration specification formats:
+    - Single value: `"50"` - Test single duration of 50μs
+    - List format: `"[25,50,75]"` - Test multiple specific durations
+    - Range format: `"[25:75]"` - Test all values from 25μs to 75μs
+    - Range with step: `"[0:50:10]"` - Test from 0μs to 50μs in steps of 10μs
+    Performed by [NVBench](https://github.com/NVIDIA/nvbench) sleep kernel benchmark.
+    #### Metrics
+    | Name                                    | Unit      | Description                                           |
+    |-----------------------------------------|-----------|-------------------------------------------------------|
+    | nvbench-sleep-kernel/duration_us_{X}_cpu_time     | time (μs) | CPU-measured time for duration X microseconds.       |
+    | nvbench-sleep-kernel/duration_us_{X}_gpu_time     | time (μs) | GPU-measured time for duration X microseconds.       |
+    | nvbench-sleep-kernel/duration_us_{X}_batch_gpu_time | time (μs) | GPU batch execution time for duration X microseconds. |
+    Where `{X}` is the sleep duration in microseconds (e.g., 25, 50, 75).
+    ### `nvbench-kernel-launch`
+    #### Introduction
+    Measure GPU kernel launch overhead and execution time using NVBench's kernel launch benchmark. This benchmark evaluates the time required to launch kernels on the GPU and measures both CPU-side and GPU-side timing for kernel execution.
+    The benchmark provides insights into:
+    - Kernel launch latency
+    - CPU/GPU synchronization overhead
+    - Batch execution performance
+    Performed by [NVBench](https://github.com/NVIDIA/nvbench) kernel launch benchmark.
+    #### Metrics
+    | Name                                | Unit      | Description                                    |
+    |-------------------------------------|-----------|------------------------------------------------|
+    | nvbench-kernel-launch/cpu_time      | time (μs) | CPU-measured kernel execution time.            |
+    | nvbench-kernel-launch/gpu_time      | time (μs) | GPU-measured kernel execution time.            |
+    | nvbench-kernel-launch/batch_gpu_time | time (μs) | GPU batch execution time.                     |
     ### `cpu-hpl`
     #### Introduction
@@ Expand Down @@

-Original file line number
+Diff line change
@@ -0,0 +1,30 @@
+    # Copyright (c) Microsoft Corporation.
+    # Licensed under the MIT License.
+    """Example of NVBench Kernel Launch benchmark."""
+    from superbench.benchmarks import BenchmarkRegistry, Platform
+    from superbench.common.utils import logger
+    if __name__ == '__main__':
+        context = BenchmarkRegistry.create_benchmark_context(
+            'nvbench-kernel-launch',
+            platform=Platform.CUDA,
+            parameters=(
+                '--timeout 30 '
+                '--min-samples 10 '
+                '--min-time 1.0 '
+                '--max-noise 0.1 '
+                '--stopping-criterion stdrel '
+                '--throttle-threshold 80 '
+                '--throttle-recovery-delay 1.0'
+            )
+        )
+        benchmark = BenchmarkRegistry.launch_benchmark(context)
+        if benchmark:
+            logger.info(
+                'benchmark: {}, return code: {}, result: {}'.format(
+                    benchmark.name, benchmark.return_code, benchmark.result
+                )
+            )

-Original file line number
+Diff line change
@@ -0,0 +1,45 @@
+    cmake_minimum_required(VERSION 3.18)
+    project(nvbench_benchmarks LANGUAGES CUDA)
+    # Check if we have a recent enough CMake for nvbench (which requires 3.30.4)
+    if(CMAKE_VERSION VERSION_LESS "3.30.4")
+      message(STATUS "CMake version ${CMAKE_VERSION} is less than 3.30.4 (required by nvbench), skipping nvbench benchmarks")
+      return()
+    endif()
+    find_package(CUDAToolkit QUIET)
+    if (CUDAToolkit_FOUND)
+      include(../cuda_common.cmake)
+      # Try to find nvbench, but don't require it
+      find_package(nvbench CONFIG QUIET)
+      if (nvbench_FOUND)
+        message(STATUS "Found nvbench, building nvbench benchmarks")
+        # list all your CUDA benchmark source files here
+        set(NVBENCH_SOURCES
+          kernel_launch.cu
+          sleep_kernel.cu
+          # add more *.cu as needed
+        )
+        foreach(src ${NVBENCH_SOURCES})
+          # strip ".cu" → NAME_WE
+          get_filename_component(basename ${src} NAME_WE)
+          set(target nvbench_${basename})
+          add_executable(${target} ${src})
+          target_compile_features(${target} PUBLIC cuda_std_17)
+          target_link_libraries(${target}
+            PRIVATE nvbench::nvbench nvbench::main
+          )
+          install(TARGETS ${target} RUNTIME DESTINATION bin)
+        endforeach()
+      else()
+        message(STATUS "nvbench not found, skipping nvbench benchmarks.")
+        message(STATUS "To build nvbench benchmarks, first build the submodule in third_party/nvbench")
+      endif()
+    else()
+      message(STATUS "CUDA not found, skipping nvbench benchmarks.")
+    endif()

-Original file line number
+Diff line change
@@ -0,0 +1,9 @@
+    #include <nvbench/nvbench.cuh>
+    __global__ void empty_kernel() {}
+    void kernel_launch(nvbench::state &state) {
+        state.exec([](nvbench::launch &launch) { empty_kernel<<<1, 1, 0, launch.get_stream()>>>(); });
+    }
+    NVBENCH_BENCH(kernel_launch);

Benchmarks: Micro benchmark - add nvbench based kernel-launch & sleep-kernel #750

Are you sure you want to change the base?

Uh oh!

Benchmarks: Micro benchmark - add nvbench based kernel-launch & sleep-kernel #750

Uh oh!

Uh oh!

Diff view

Diff view

There are no files selected for viewing

Copilot AI Jan 23, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI Jan 23, 2026

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!