INT16 add and mul as well as FP16 fma benchmark for gfx900 GPUS

gfx900_multi_precision_bench.cpp

/*
   f=gfx900_multi_precision_bench
   hipcc -O3 --offload-arch=gfx900 $f.cpp -o $f
*/
#include <hip/hip_runtime.h>
#include <hip/hip_fp16.h>
#include <iostream>
#include <iomanip>
#include <vector>

// Define 2-element 16-bit vector types for Clang/LLVM
typedef short v2i16 __attribute__((ext_vector_type(2)));
typedef unsigned short v2u16 __attribute__((ext_vector_type(2)));
typedef __fp16 v2f16 __attribute__((ext_vector_type(2)));

#define HIP_CHECK(status) \
    if (status != hipSuccess) { \
        std::cerr << "HIP Error: " << hipGetErrorString(status) << " at line " << __LINE__ << std::endl; \
        exit(EXIT_FAILURE); \
    }

// ----------------------------------------------------------------------------
// Kernel 1: Packed 16-bit Integer Addition (v_pk_add_i16) -> 2 Ops/Instruction
// ----------------------------------------------------------------------------
__global__ void kernel_pk_add_i16(v2i16* d_out, v2i16 val, int inner_loop_count) {
    v2i16 acc0 = {0, 0}, acc1 = {0, 0}, acc2 = {0, 0}, acc3 = {0, 0};
    v2i16 acc4 = {0, 0}, acc5 = {0, 0}, acc6 = {0, 0}, acc7 = {0, 0};
    v2i16 src0 = val;

    for (int i = 0; i < inner_loop_count; ++i) {
        acc0 += src0;
        acc1 += src0;
        acc2 += src0;
        acc3 += src0;
        acc4 += src0;
        acc5 += src0;
        acc6 += src0;
        acc7 += src0;
    }

    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    d_out[idx] = acc0 + acc1 + acc2 + acc3 + acc4 + acc5 + acc6 + acc7;
}

// ----------------------------------------------------------------------------
// Kernel 2: Packed 16-bit Low Multiplication (v_pk_mul_lo_u16) -> 2 Ops/Instruction
// ----------------------------------------------------------------------------
__global__ void kernel_pk_mul_u16(v2i16* d_out, v2i16 val, int inner_loop_count) {
    v2i16 acc0 = {1, 1}, acc1 = {1, 1}, acc2 = {1, 1}, acc3 = {1, 1};
    v2i16 acc4 = {1, 1}, acc5 = {1, 1}, acc6 = {1, 1}, acc7 = {1, 1};
    v2i16 src0 = val;

    for (int i = 0; i < inner_loop_count; ++i) {
        acc0 *= src0;
        acc1 *= src0;
        acc2 *= src0;
        acc3 *= src0;
        acc4 *= src0;
        acc5 *= src0;
        acc6 *= src0;
        acc7 *= src0;
    }

    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    d_out[idx] = acc0 + acc1 + acc2 + acc3 + acc4 + acc5 + acc6 + acc7;
}

// ----------------------------------------------------------------------------
// Kernel 3: Packed FP16 Fused Multiply-Add (v_pk_fma_f16) -> 4 FLOPs/Instruction
// ----------------------------------------------------------------------------
__global__ void kernel_pk_fma_f16(v2f16* d_out, v2f16 val, int inner_loop_count) {
    v2f16 acc0 = {1.0f16, 1.0f16}, acc1 = {1.0f16, 1.0f16};
    v2f16 acc2 = {1.0f16, 1.0f16}, acc3 = {1.0f16, 1.0f16};
    v2f16 acc4 = {1.0f16, 1.0f16}, acc5 = {1.0f16, 1.0f16};
    v2f16 acc6 = {1.0f16, 1.0f16}, acc7 = {1.0f16, 1.0f16};

    v2f16 src0 = val;
    v2f16 src1 = {1.001f16, 0.999f16};

    for (int i = 0; i < inner_loop_count; ++i) {
        // Generates v_pk_fma_f16 (acc = acc * src0 + src1)
        acc0 = __builtin_elementwise_fma(acc0, src0, src1);
        acc1 = __builtin_elementwise_fma(acc1, src0, src1);
        acc2 = __builtin_elementwise_fma(acc2, src0, src1);
        acc3 = __builtin_elementwise_fma(acc3, src0, src1);
        acc4 = __builtin_elementwise_fma(acc4, src0, src1);
        acc5 = __builtin_elementwise_fma(acc5, src0, src1);
        acc6 = __builtin_elementwise_fma(acc6, src0, src1);
        acc7 = __builtin_elementwise_fma(acc7, src0, src1);
    }

    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    d_out[idx] = acc0 + acc1 + acc2 + acc3 + acc4 + acc5 + acc6 + acc7;
}

// ----------------------------------------------------------------------------
// Helper Function: Run Benchmark & Print Standard Statistics
// ----------------------------------------------------------------------------
template <typename T>
void run_and_report_benchmark(
    const char* kernel_name,
    void (*kernel_func)(T*, T, int),
    T* d_out,
    T val,
    int inner_loop_count,
    int grid_size,
    int block_size,
    double target_clock_mhz,
    double ops_per_instruction // 2.0 for add/mul, 4.0 for FMA
) {
    hipEvent_t start, stop;
    HIP_CHECK(hipEventCreate(&start));
    HIP_CHECK(hipEventCreate(&stop));

    // Warmup
    hipLaunchKernelGGL(kernel_func, dim3(grid_size), dim3(block_size), 0, 0, d_out, val, 100);
    HIP_CHECK(hipDeviceSynchronize());

    // Timed run
    HIP_CHECK(hipEventRecord(start, 0));
    hipLaunchKernelGGL(kernel_func, dim3(grid_size), dim3(block_size), 0, 0, d_out, val, inner_loop_count);
    HIP_CHECK(hipEventRecord(stop, 0));
    HIP_CHECK(hipEventSynchronize(stop));

    float time_ms = 0.0f;
    HIP_CHECK(hipEventElapsedTime(&time_ms, start, stop));

    double total_threads = static_cast<double>(grid_size) * block_size;
    double ops_per_thread = static_cast<double>(inner_loop_count) * 8.0 * ops_per_instruction;
    double total_ops = total_threads * ops_per_thread;

    double time_sec = time_ms / 1000.0;
    double giops = (total_ops / 1e9) / time_sec;
    double tops = giops / 1000.0;

    // Peak calculation: 4096 ALU lanes * ops_per_instruction * clock_mhz * 1e6
    double theoretical_peak_giops = 4096.0 * ops_per_instruction * (target_clock_mhz * 1e6) / 1e9;
    double percent_of_peak = (giops / theoretical_peak_giops) * 100.0;

    const char* unit_label = (ops_per_instruction > 2.0) ? "GFLOPS" : "GIOPS";
    const char* top_label  = (ops_per_instruction > 2.0) ? "TFLOPS" : "TOPS";

    std::cout << "\n======================================================\n";
    std::cout << " Kernel: " << kernel_name << "\n";
    std::cout << "------------------------------------------------------\n";
    std::cout << "  Execution Time        : " << std::fixed << std::setprecision(3) << time_ms << " ms\n";
    std::cout << "  Throughput            : " << std::fixed << std::setprecision(2) << giops << " " << unit_label << " (" << tops << " " << top_label << ")\n";
    std::cout << "  Theoretical Peak      : " << std::fixed << std::setprecision(2) << theoretical_peak_giops << " " << unit_label << " (" << target_clock_mhz << " MHz)\n";
    std::cout << "  Efficiency            : " << std::fixed << std::setprecision(2) << percent_of_peak << " % of Peak\n";
    std::cout << "======================================================\n";

    HIP_CHECK(hipEventDestroy(start));
    HIP_CHECK(hipEventDestroy(stop));
}

int main() {
    hipDeviceProp_t prop;
    HIP_CHECK(hipGetDeviceProperties(&prop, 0));

    std::cout << "Device Name             : " << prop.name << "\n";
    std::cout << "GCN Arch                : " << prop.gcnArchName << "\n";
    std::cout << "Compute Units           : " << prop.multiProcessorCount << "\n";

    double target_clock_mhz = prop.clockRate/1000.0;

    const int block_size = 256;
    const int grid_size = prop.multiProcessorCount * 16;
    const int inner_loop_count = 50'000'000;

    v2i16* d_out_i16 = nullptr;
    v2f16* d_out_f16 = nullptr;
    size_t data_size = grid_size * block_size * sizeof(v2i16);
    HIP_CHECK(hipMalloc(&d_out_i16, data_size));
    HIP_CHECK(hipMalloc(&d_out_f16, data_size));

    v2i16 init_val_i16 = {2, 3};
    v2f16 init_val_f16 = {1.0001f16, 1.0002f16};

    // Benchmark Kernels 1 & 2 (2 ops/instruction)
    run_and_report_benchmark("v_pk_add_i16 (Packed Add)", kernel_pk_add_i16, d_out_i16, init_val_i16, inner_loop_count, grid_size, block_size, target_clock_mhz, 2.0);
    run_and_report_benchmark("v_pk_mul_lo_u16 (Packed Mul)", kernel_pk_mul_u16, d_out_i16, init_val_i16, inner_loop_count, grid_size, block_size, target_clock_mhz, 2.0);

    // Benchmark Kernel 3: FP16 Packed FMA (4 FLOPs/instruction -> ~25.3-26.7 TFLOPS peak)
    run_and_report_benchmark("v_pk_fma_f16 (Packed FP16 FMA)", kernel_pk_fma_f16, d_out_f16, init_val_f16, inner_loop_count, grid_size, block_size, target_clock_mhz, 4.0);

    HIP_CHECK(hipFree(d_out_i16));
    HIP_CHECK(hipFree(d_out_f16));
    return 0;
}

For execution on next gen GPUs (gfx906) compile command needs to be adapted:

hermann@Radeon-vii:~$ hipcc -O3 --offload-arch=gfx906 $f.cpp -o $f
hermann@Radeon-vii:~$ sudo rocm-smi --setperflevel high


============================ ROCm System Management Interface ============================
================================= Set Performance Level ==================================
GPU[0]		: Performance level set to high
==========================================================================================
================================== End of ROCm SMI Log ===================================
hermann@Radeon-vii:~$ ./gfx900_multi_precision_bench 
Device Name             : AMD Radeon VII
GCN Arch                : gfx906:sramecc+:xnack-
Compute Units           : 60

======================================================
 Kernel: v_pk_add_i16 (Packed Add)
------------------------------------------------------
  Execution Time        : 15048.847 ms
  Throughput            : 13064.66 GIOPS (13.06 TOPS)
  Theoretical Peak      : 14753.79 GIOPS (1801.00 MHz)
  Efficiency            : 88.55 % of Peak
======================================================

======================================================
 Kernel: v_pk_mul_lo_u16 (Packed Mul)
------------------------------------------------------
  Execution Time        : 14955.788 ms
  Throughput            : 13145.95 GIOPS (13.15 TOPS)
  Theoretical Peak      : 14753.79 GIOPS (1801.00 MHz)
  Efficiency            : 89.10 % of Peak
======================================================

======================================================
 Kernel: v_pk_fma_f16 (Packed FP16 FMA)
------------------------------------------------------
  Execution Time        : 15318.190 ms
  Throughput            : 25669.87 GFLOPS (25.67 TFLOPS)
  Theoretical Peak      : 29507.58 GFLOPS (1801.00 MHz)
  Efficiency            : 86.99 % of Peak
======================================================
hermann@Radeon-vii:~$