Demonstrate 77 TFLOPS FP16 (__half) on NVIDIA RTX 5060 GPU

benchmark_fp16.cu

#include <iostream>
#include <vector>
#include <cuda_runtime.h>
#include <cublas_v2.h>

#define CUDA_CHECK(status) \
    if (status != cudaSuccess) { \
        std::cerr << "CUDA Error: " << cudaGetErrorString(status) \
                  << " at line " << __LINE__ << std::endl; \
        exit(EXIT_FAILURE); \
    }

#define CUBLAS_CHECK(status) \
    if (status != CUBLAS_STATUS_SUCCESS) { \
        std::cerr << "cuBLAS Error: " << status \
                  << " at line " << __LINE__ << std::endl; \
        exit(EXIT_FAILURE); \
    }

void run_fp16_benchmark(cublasHandle_t handle, int M, int N, int K, int warmup = 10, int bench = 100) {
    size_t size_A = (size_t)M * K * sizeof(__half);
    size_t size_B = (size_t)K * N * sizeof(__half);
    size_t size_C = (size_t)M * N * sizeof(__half);

    __half *d_A, *d_B, *d_C;
    CUDA_CHECK(cudaMalloc(&d_A, size_A));
    CUDA_CHECK(cudaMalloc(&d_B, size_B));
    CUDA_CHECK(cudaMalloc(&d_C, size_C));

    CUDA_CHECK(cudaMemset(d_A, 0x3C, size_A));
    CUDA_CHECK(cudaMemset(d_B, 0x3C, size_B));
    CUDA_CHECK(cudaMemset(d_C, 0, size_C));

    // For CUBLAS_COMPUTE_16F, alpha and beta must be __half
    __half alpha = __float2half(1.0f);
    __half beta  = __float2half(0.0f);

    // Warmup
    for (int i = 0; i < warmup; ++i) {
        CUBLAS_CHECK(cublasGemmEx(
            handle, 
            CUBLAS_OP_N, CUBLAS_OP_N,
            M, N, K,
            &alpha, 
            d_A, CUDA_R_16F, M,
            d_B, CUDA_R_16F, K,
            &beta, 
            d_C, CUDA_R_16F, M,
            CUBLAS_COMPUTE_16F,
            CUBLAS_GEMM_DEFAULT_TENSOR_OP
        ));
    }
    CUDA_CHECK(cudaDeviceSynchronize());

    cudaEvent_t start, stop;
    CUDA_CHECK(cudaEventCreate(&start));
    CUDA_CHECK(cudaEventCreate(&stop));

    // Timed Loop
    CUDA_CHECK(cudaEventRecord(start));
    for (int i = 0; i < bench; ++i) {
        CUBLAS_CHECK(cublasGemmEx(
            handle, 
            CUBLAS_OP_N, CUBLAS_OP_N,
            M, N, K,
            &alpha, 
            d_A, CUDA_R_16F, M,
            d_B, CUDA_R_16F, K,
            &beta, 
            d_C, CUDA_R_16F, M,
            CUBLAS_COMPUTE_16F,
            CUBLAS_GEMM_DEFAULT_TENSOR_OP
        ));
    }
    CUDA_CHECK(cudaEventRecord(stop));
    CUDA_CHECK(cudaEventSynchronize(stop));

    float total_ms = 0;
    CUDA_CHECK(cudaEventElapsedTime(&total_ms, start, stop));

    double avg_ms = total_ms / bench;
    double total_flops = 2.0 * (double)M * (double)N * (double)K;
    double tflops = (total_flops / (avg_ms * 1e-3)) / 1e12;

    std::cout << "--------------------------------------------------------\n";
    std::cout << " Matrix Size (M x N x K): " << M << " x " << N << " x " << K << "\n";
    std::cout << " Average Time per GEMM  : " << avg_ms << " ms\n";
    std::cout << " Sustained FP16 TFLOPS  : " << tflops << " TFLOPS\n";
    std::cout << "--------------------------------------------------------\n";

    CUDA_CHECK(cudaFree(d_A));
    CUDA_CHECK(cudaFree(d_B));
    CUDA_CHECK(cudaFree(d_C));
}

int main() {
    int deviceCount = 0;
    CUDA_CHECK(cudaGetDeviceCount(&deviceCount));

    cudaDeviceProp prop;
    CUDA_CHECK(cudaGetDeviceProperties(&prop, 0));
    std::cout << "Running FP16 Benchmark on GPU 0: " << prop.name
              << " (Compute Capability: " << prop.major << "." << prop.minor << ")\n";

    cublasHandle_t handle;
    CUBLAS_CHECK(cublasCreate(&handle));

    std::vector<int> dimensions = {4096, 8192, 16384};
    for (int dim : dimensions) {
        run_fp16_benchmark(handle, dim, dim, dim);
    }

    cublasDestroy(handle);
    return 0;
}

Created in a long gemini session:

$ nvcc -O3 -arch=sm_120a benchmark_fp16.cu -o benchmark_fp16 -lcublas
$ ./benchmark_fp16
Running FP16 Benchmark on GPU 0: NVIDIA GeForce RTX 5060 (Compute Capability: 12.0)
--------------------------------------------------------
 Matrix Size (M x N x K): 4096 x 4096 x 4096
 Average Time per GEMM  : 1.88534 ms
 Sustained FP16 TFLOPS  : 72.8988 TFLOPS
--------------------------------------------------------
--------------------------------------------------------
 Matrix Size (M x N x K): 8192 x 8192 x 8192
 Average Time per GEMM  : 14.2612 ms
 Sustained FP16 TFLOPS  : 77.0979 TFLOPS
--------------------------------------------------------
--------------------------------------------------------
 Matrix Size (M x N x K): 16384 x 16384 x 16384
 Average Time per GEMM  : 116.165 ms
 Sustained FP16 TFLOPS  : 75.7207 TFLOPS
--------------------------------------------------------
$

In very long session gemini was not able to demonstrate >200 TFPLOPS with FP4.
But it helped to make cutlass demo show 279 TFLOPS FP4 (quarter petaflops):

hermann@7950x:~/cuda/cutlass/examples/79_blackwell_geforce_gemm$ ./79a_blackwell_geforce_nvfp4_nvfp4_gemm --m=6144 --n=6144 --k=6144
  Disposition: Passed
  Problem Size: 6144x6144x6144
  Avg runtime: 1.66186 ms
  GFLOPS: 279120
hermann@7950x:~/cuda/cutlass/examples/79_blackwell_geforce_gemm$ 

Compiled with:

hermann@7950x:~/cuda/cutlass/examples/79_blackwell_geforce_gemm$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2025 NVIDIA Corporation
Built on Fri_Feb_21_20:23:50_PST_2025
Cuda compilation tools, release 12.8, V12.8.93
Build cuda_12.8.r12.8/compiler.35583870_0
hermann@7950x:~/cuda/cutlass/examples/79_blackwell_geforce_gemm$ nvcc -O3 -arch=sm_120a --expt-relaxed-constexpr -std=c++17   79a_blackwell_geforce_nvfp4_bf16_gemm.cu   -o 79a_blackwell_geforce_nvfp4_bf16_gemm   -DCUTLASS_ARCH_MMA_SM120_SUPPORTED   -Xcudafe --diag_suppress=20012   -I../../include   -I../../tools/util/include   -I../../examples/common
hermann@7950x:~/cuda/cutlass/examples/79_blackwell_geforce_gemm$

More than 1 POPS (1000 TOPS) shown here (10 old server GPUs):
https://github.com/Hermann-SW/1.0003-POPS/blob/main/README.md#10003-pops