Multi GPU INT4/INT8/FP16 benchmark from long gemini session

mi50_multi_precision_bench.cpp

#include <hip/hip_runtime.h>
#include <hip/hip_fp16.h>
#include <iostream>
#include <vector>
#include <thread>
#include <chrono>
#include <iomanip>
#include <atomic>
#include <string>

enum PrecisionMode {
    PREC_INT4,
    PREC_INT8,
    PREC_FP16
};

// ---------------------------------------------------------------------------
// Compute Kernels
// ---------------------------------------------------------------------------

__global__ void __launch_bounds__(256, 2) mi50_int4_kernel(uint64_t iterations, int* dummy_out) {
    int src0 = 0x12345678;
    int src1 = 0x87654321;
    int acc0 = 0, acc1 = 0, acc2 = 0, acc3 = 0;
    int acc4 = 0, acc5 = 0, acc6 = 0, acc7 = 0;

    #pragma unroll 1
    for (uint64_t i = 0; i < iterations; ++i) {
        #pragma unroll
        for (int k = 0; k < 16; ++k) {
            acc0 = __builtin_amdgcn_sdot8(src0, src1, acc0, false);
            acc1 = __builtin_amdgcn_sdot8(src0, src1, acc1, false);
            acc2 = __builtin_amdgcn_sdot8(src0, src1, acc2, false);
            acc3 = __builtin_amdgcn_sdot8(src0, src1, acc3, false);
            acc4 = __builtin_amdgcn_sdot8(src0, src1, acc4, false);
            acc5 = __builtin_amdgcn_sdot8(src0, src1, acc5, false);
            acc6 = __builtin_amdgcn_sdot8(src0, src1, acc6, false);
            acc7 = __builtin_amdgcn_sdot8(src0, src1, acc7, false);
        }
    }

    if (threadIdx.x == 0 && blockIdx.x == 0) {
        *dummy_out = acc0 + acc1 + acc2 + acc3 + acc4 + acc5 + acc6 + acc7;
    }
}

__global__ void __launch_bounds__(256, 2) mi50_int8_kernel(uint64_t iterations, int* dummy_out) {
    int src0 = 0x12345678;
    int src1 = 0x87654321;
    int acc0 = 0, acc1 = 0, acc2 = 0, acc3 = 0;
    int acc4 = 0, acc5 = 0, acc6 = 0, acc7 = 0;

    #pragma unroll 1
    for (uint64_t i = 0; i < iterations; ++i) {
        #pragma unroll
        for (int k = 0; k < 16; ++k) {
            acc0 = __builtin_amdgcn_sdot4(src0, src1, acc0, false);
            acc1 = __builtin_amdgcn_sdot4(src0, src1, acc1, false);
            acc2 = __builtin_amdgcn_sdot4(src0, src1, acc2, false);
            acc3 = __builtin_amdgcn_sdot4(src0, src1, acc3, false);
            acc4 = __builtin_amdgcn_sdot4(src0, src1, acc4, false);
            acc5 = __builtin_amdgcn_sdot4(src0, src1, acc5, false);
            acc6 = __builtin_amdgcn_sdot4(src0, src1, acc6, false);
            acc7 = __builtin_amdgcn_sdot4(src0, src1, acc7, false);
        }
    }

    if (threadIdx.x == 0 && blockIdx.x == 0) {
        *dummy_out = acc0 + acc1 + acc2 + acc3 + acc4 + acc5 + acc6 + acc7;
    }
}

__global__ void __launch_bounds__(256, 2) mi50_fp16_kernel(uint64_t iterations, half2* dummy_out) {
    half2 src0 = __float2half2_rn(1.001f);
    half2 src1 = __float2half2_rn(0.999f);
    half2 acc0 = __float2half2_rn(0.1f);
    half2 acc1 = __float2half2_rn(0.2f);
    half2 acc2 = __float2half2_rn(0.3f);
    half2 acc3 = __float2half2_rn(0.4f);
    half2 acc4 = __float2half2_rn(0.5f);
    half2 acc5 = __float2half2_rn(0.6f);
    half2 acc6 = __float2half2_rn(0.7f);
    half2 acc7 = __float2half2_rn(0.8f);

    #pragma unroll 1
    for (uint64_t i = 0; i < iterations; ++i) {
        #pragma unroll
        for (int k = 0; k < 16; ++k) {
            acc0 = __hfma2(src0, src1, acc0);
            acc1 = __hfma2(src0, src1, acc1);
            acc2 = __hfma2(src0, src1, acc2);
            acc3 = __hfma2(src0, src1, acc3);
            acc4 = __hfma2(src0, src1, acc4);
            acc5 = __hfma2(src0, src1, acc5);
            acc6 = __hfma2(src0, src1, acc6);
            acc7 = __hfma2(src0, src1, acc7);
        }
    }

    if (threadIdx.x == 0 && blockIdx.x == 0) {
        *dummy_out = acc0 + acc1 + acc2 + acc3 + acc4 + acc5 + acc6 + acc7;
    }
}

// ---------------------------------------------------------------------------
// Worker Infrastructure
// ---------------------------------------------------------------------------

struct GpuResult {
    int dev_id;
    std::string name;
    float kernel_ms;
    double tops;
    double total_ops;
    bool success;
};

void worker_thread(int dev_id, PrecisionMode mode, uint64_t iterations, GpuResult* result,
                   std::atomic<int>& ready_count, std::atomic<bool>& start_flag) {
    result->dev_id = dev_id;
    result->success = false;

    if (hipSetDevice(dev_id) != hipSuccess) return;

    hipDeviceProp_t prop;
    if (hipGetDeviceProperties(&prop, dev_id) != hipSuccess) return;
    result->name = prop.name;

    const int num_blocks = 240; // 60 CUs * 4 blocks
    const int threads_per_block = 256;
    const double total_threads = (double)num_blocks * threads_per_block;

    double ops_per_thread_per_iter = 0.0;
    if (mode == PREC_INT4)      ops_per_thread_per_iter = 128.0 * 16.0; // 2048 ops
    else if (mode == PREC_INT8) ops_per_thread_per_iter = 128.0 * 8.0;  // 1024 ops
    else if (mode == PREC_FP16) ops_per_thread_per_iter = 128.0 * 4.0;  // 512 FLOPs

    result->total_ops = ops_per_thread_per_iter * (double)iterations * total_threads;

    void* d_out = nullptr;
    if (hipMalloc(&d_out, sizeof(double)) != hipSuccess) return;

    // Warmup pass
    if (mode == PREC_INT4) {
        mi50_int4_kernel<<<num_blocks, threads_per_block>>>(10000ULL, (int*)d_out);
    } else if (mode == PREC_INT8) {
        mi50_int8_kernel<<<num_blocks, threads_per_block>>>(10000ULL, (int*)d_out);
    } else if (mode == PREC_FP16) {
        mi50_fp16_kernel<<<num_blocks, threads_per_block>>>(10000ULL, (half2*)d_out);
    }
    
    if (hipDeviceSynchronize() != hipSuccess) {
        hipFree(d_out);
        return;
    }

    hipEvent_t start, stop;
    hipEventCreate(&start);
    hipEventCreate(&stop);

    ready_count++;

    while (!start_flag.load(std::memory_order_acquire)) {
        std::this_thread::yield();
    }

    // Synchronized launch
    hipEventRecord(start);
    if (mode == PREC_INT4) {
        mi50_int4_kernel<<<num_blocks, threads_per_block>>>(iterations, (int*)d_out);
    } else if (mode == PREC_INT8) {
        mi50_int8_kernel<<<num_blocks, threads_per_block>>>(iterations, (int*)d_out);
    } else if (mode == PREC_FP16) {
        mi50_fp16_kernel<<<num_blocks, threads_per_block>>>(iterations, (half2*)d_out);
    }
    hipEventRecord(stop);
    hipDeviceSynchronize();

    float milliseconds = 0;
    hipEventElapsedTime(&milliseconds, start, stop);

    double seconds = milliseconds / 1000.0;
    result->kernel_ms = milliseconds;
    result->tops = (result->total_ops / seconds) / 1e12;
    result->success = true;

    hipEventDestroy(start);
    hipEventDestroy(stop);
    hipFree(d_out);
}

double run_benchmark_pass(int device_count, PrecisionMode mode, const std::string& label, uint64_t iterations) {
    std::cout << "=============================================================================\n";
    std::cout << " RUNNING " << label << " BENCHMARK PASS (" << iterations << " iterations)\n";
    std::cout << "=============================================================================\n";

    std::vector<GpuResult> results(device_count);
    std::vector<std::thread> threads;
    std::atomic<int> ready_count(0);
    std::atomic<bool> start_flag(false);

    for (int i = 0; i < device_count; ++i) {
        threads.emplace_back(worker_thread, i, mode, iterations, &results[i], 
                             std::ref(ready_count), std::ref(start_flag));
    }

    while (ready_count.load(std::memory_order_relaxed) < device_count) {
        std::this_thread::yield();
    }

    auto wall_start = std::chrono::high_resolution_clock::now();
    start_flag.store(true, std::memory_order_release);

    for (auto& t : threads) {
        t.join();
    }

    auto wall_stop = std::chrono::high_resolution_clock::now();
    double wall_seconds = std::chrono::duration<double>(wall_stop - wall_start).count();

    std::cout << std::left << std::setw(10) << "Logical ID"
              << std::setw(30) << "Device Name"
              << std::setw(16) << "Time (ms)"
              << std::setw(16) << (mode == PREC_FP16 ? "Per-GPU TFLOPS" : "Per-GPU TOPS") << "\n";
    std::cout << std::string(72, '-') << "\n";

    double sum_ops = 0.0;
    int active_gpus = 0;

    for (const auto& r : results) {
        if (r.success) {
            std::cout << std::left << std::setw(10) << r.dev_id
                      << std::setw(30) << r.name
                      << std::setw(16) << std::fixed << std::setprecision(2) << r.kernel_ms
                      << std::setw(16) << std::setprecision(2) << r.tops << "\n";
            sum_ops += r.total_ops;
            active_gpus++;
        }
    }

    std::cout << std::string(72, '-') << "\n";
    double aggregate_throughput = (sum_ops / wall_seconds) / 1e12;

    std::cout << "Active GPUs:            " << active_gpus << " / " << device_count << "\n";
    std::cout << "Concurrent Wall Time:   " << std::fixed << std::setprecision(3) << wall_seconds * 1000.0 << " ms\n";
    std::cout << "AGGREGATE SYSTEM " << label << ": " << std::setprecision(2) << aggregate_throughput 
              << (mode == PREC_FP16 ? " TFLOPS" : " TOPS") << "\n\n";

    return aggregate_throughput;
}

int main() {
    int device_count = 0;
    if (hipGetDeviceCount(&device_count) != hipSuccess || device_count == 0) {
        std::cerr << "No valid HIP devices found." << std::endl;
        return 1;
    }

    std::cout << "Dispatched multi-precision benchmark across " << device_count << " visible GPU device(s)...\n\n";

    const uint64_t iterations = 1000000ULL;

    double agg_int4 = run_benchmark_pass(device_count, PREC_INT4, "INT4", iterations);
    double agg_int8 = run_benchmark_pass(device_count, PREC_INT8, "INT8", iterations);
    double agg_fp16 = run_benchmark_pass(device_count, PREC_FP16, "FP16", iterations);

    std::cout << "=============================================================================\n";
    std::cout << "                       SYSTEM SUMMARY PERFORMANCE                            \n";
    std::cout << "=============================================================================\n";
    std::cout << " INT4 Aggregate Compute:  " << std::fixed << std::setprecision(2) << agg_int4 << " TOPS\n";
    std::cout << " INT8 Aggregate Compute:  " << std::fixed << std::setprecision(2) << agg_int8 << " TOPS\n";
    std::cout << " FP16 Aggregate Compute:  " << std::fixed << std::setprecision(2) << agg_fp16 << " TFLOPS\n";
    std::cout << "=============================================================================\n";

    return 0;
}

Compile:

hermann@7600x:~$ hipcc -O3 -std=c++11 --offload-arch=gfx906 mi50_multi_precision_bench.cpp -o mi50_multi_precision_bench
hermann@7600x:~$ 

It is important that iGPU is disabled either in bios, or through udev rules or other means.
On my system only the 8 MI50s are reported.
The one without MI50 in name (later) and with 225W power cap is my only 32GB MI50, the others are 16GB:

hermann@7600x:~$ rocm-smi 


============================================ ROCm System Management Interface ============================================
====================================================== Concise Info ======================================================
Device  Node  IDs              Temp    Power     Partitions          SCLK     MCLK     Fan     Perf  PwrCap  VRAM%  GPU%  
              (DID,     GUID)  (Edge)  (Socket)  (Mem, Compute, ID)                                                       
==========================================================================================================================
0       1     0x66a1,   45820  32.0°C  30.0W     N/A, N/A, 0         1725Mhz  1000Mhz  14.51%  high  250.0W  0%     0%    
1       2     0x66a1,   29947  34.0°C  99.0W     N/A, N/A, 0         1725Mhz  1000Mhz  14.51%  high  250.0W  0%     0%    
2       3     0x66a1,   10040  35.0°C  40.0W     N/A, N/A, 0         1725Mhz  1000Mhz  14.51%  high  250.0W  0%     0%    
3       4     0x66a1,   8048   36.0°C  40.0W     N/A, N/A, 0         1725Mhz  1000Mhz  14.51%  high  250.0W  0%     0%    
4       5     0x66a1,   19635  36.0°C  42.0W     N/A, N/A, 0         1725Mhz  1000Mhz  14.51%  high  250.0W  0%     0%    
5       6     0x66a1,   48627  36.0°C  35.0W     N/A, N/A, 0         1725Mhz  1000Mhz  14.51%  high  250.0W  0%     0%    
6       7     0x66a1,   10295  36.0°C  32.0W     N/A, N/A, 0         1725Mhz  1000Mhz  14.51%  high  250.0W  0%     0%    
7       8     0x66a1,   59764  34.0°C  89.0W     N/A, N/A, 0         1725Mhz  1000Mhz  14.51%  high  225.0W  0%     0%    
==========================================================================================================================
================================================== End of ROCm SMI Log ===================================================
hermann@7600x:~$ 

Nice INT4+INT8 TOPS, and nice FP16 TFLOPS.
It started with gemini working on INT4 benchmark demonstrating 100 TOPS INT4.
Because I had problem with one MI50 I showed gemini rocm-smi output.
After the problem was resolved, gemini asked whether I wanted to extend to 800 TOPS multi-GPU benchmark.
Later it asked me whether I wanted to see benchmark for INT8 or FP16 as well.
I said that I wanted a combined benchmark for INT4/INT8/FP16, and the result is this gist.
Really nice summary from below:

=============================================================================
                       SYSTEM SUMMARY PERFORMANCE                            
=============================================================================
 INT4 Aggregate Compute:  799.58 TOPS
 INT8 Aggregate Compute:  399.87 TOPS
 FP16 Aggregate Compute:  207.80 TFLOPS
=============================================================================

hermann@7600x:~$ sudo rocm-smi --setperflevel high
...
GPU[6]		: Performance level set to high
GPU[7]		: Performance level set to high
==========================================================================================
================================== End of ROCm SMI Log ===================================
hermann@7600x:~$ HSA_OVERRIDE_GFX_VERSION=9.0.6 ./mi50_multi_precision_bench
Dispatched multi-precision benchmark across 8 visible GPU device(s)...

=============================================================================
 RUNNING INT4 BENCHMARK PASS (1000000 iterations)
=============================================================================
Logical IDDevice Name                   Time (ms)       Per-GPU TOPS    
------------------------------------------------------------------------
0         AMD Instinct MI50/MI60        1258.34         100.00          
1         AMD Instinct MI50/MI60        1258.34         100.00          
2         AMD Instinct MI50/MI60        1258.34         100.00          
3         AMD Instinct MI50/MI60        1258.34         100.00          
4         AMD Instinct MI50/MI60        1258.39         99.99           
5         AMD Instinct MI50/MI60        1258.34         100.00          
6         AMD Instinct MI50/MI60        1258.33         100.00          
7         AMD Radeon Graphics           1258.33         100.00          
------------------------------------------------------------------------
Active GPUs:            8 / 8
Concurrent Wall Time:   1258.953 ms
AGGREGATE SYSTEM INT4: 799.58 TOPS

=============================================================================
 RUNNING INT8 BENCHMARK PASS (1000000 iterations)
=============================================================================
Logical IDDevice Name                   Time (ms)       Per-GPU TOPS    
------------------------------------------------------------------------
0         AMD Instinct MI50/MI60        1258.45         49.99           
1         AMD Instinct MI50/MI60        1258.34         50.00           
2         AMD Instinct MI50/MI60        1258.34         50.00           
3         AMD Instinct MI50/MI60        1258.34         50.00           
4         AMD Instinct MI50/MI60        1258.57         49.99           
5         AMD Instinct MI50/MI60        1258.35         50.00           
6         AMD Instinct MI50/MI60        1258.33         50.00           
7         AMD Radeon Graphics           1258.33         50.00           
------------------------------------------------------------------------
Active GPUs:            8 / 8
Concurrent Wall Time:   1258.709 ms
AGGREGATE SYSTEM INT8: 399.87 TOPS

=============================================================================
 RUNNING FP16 BENCHMARK PASS (1000000 iterations)
=============================================================================
Logical IDDevice Name                   Time (ms)       Per-GPU TFLOPS  
------------------------------------------------------------------------
0         AMD Instinct MI50/MI60        1210.49         25.99           
1         AMD Instinct MI50/MI60        1210.50         25.99           
2         AMD Instinct MI50/MI60        1210.50         25.99           
3         AMD Instinct MI50/MI60        1210.50         25.99           
4         AMD Instinct MI50/MI60        1210.51         25.99           
5         AMD Instinct MI50/MI60        1210.51         25.99           
6         AMD Instinct MI50/MI60        1210.49         25.99           
7         AMD Radeon Graphics           1210.49         25.99           
------------------------------------------------------------------------
Active GPUs:            8 / 8
Concurrent Wall Time:   1211.065 ms
AGGREGATE SYSTEM FP16: 207.80 TFLOPS

=============================================================================
                       SYSTEM SUMMARY PERFORMANCE                            
=============================================================================
 INT4 Aggregate Compute:  799.58 TOPS
 INT8 Aggregate Compute:  399.87 TOPS
 FP16 Aggregate Compute:  207.80 TFLOPS
=============================================================================
hermann@7600x:~$ 

Including two more gfx906 type GPUs (MI50 is gfx906 as well) via OpenMP did show 1.0003 POPS (peta operations per second)!
https://github.com/Hermann-SW/1.0003-POPS

=============================================================================
                   TOTAL 10-GPU CLUSTER PERFORMANCE                          
=============================================================================
 Combined INT4 Compute:  1000.37 TOPS
 Combined INT8 Compute:  500.05 TOPS
 Combined FP16 Compute:  260.13 TFLOPS
=============================================================================