TheRock/rocm-systems/projects/rocprofiler-compute/sample/mega_kernel/mega_kernel.hip

// Copyright (c) Advanced Micro Devices, Inc.
// SPDX-License-Identifier:  MIT

/**
 * GPU Mega Kernel Unit Test
 *
 * This file contains a comprehensive unit test kernel that exercises
 * most of the instructions/features available on AMD Instinct GPUs.
 *
 * SUPPORTED ARCHITECTURES:
 * ========================
 * - MI250/MI250X (CDNA2, gfx90a):
 *   - FP32/FP64/FP16/BF16, INT32 operations
 *   - Hardware FP64 atomics (add, min, max)
 *   - MFMA matrix operations (FP16/BF16/INT8)
 *   - Unified s_waitcnt instruction
 *
 * - MI300/MI300X/MI325X (CDNA3, gfx942):
 *   - All MI250 features
 *   - FP8/BF8 conversions (cvt_pk_fp8/bf8, CDNA3 FNUZ-style packed formats)
 *   - TF32 MFMA, Packed FP32
 *   - Sparse MFMA (FP16/BF16/FP8/INT8)
 *   - Async LDS (global_load_async_to_lds - if available)
 *   - Unified s_waitcnt instruction
 *
 * - MI350X/MI355X (CDNA4, gfx950):
 *   - All MI300 features
 *   - FP6/FP4 conversions and MFMA
 *   - Scaled FP8 conversions
 *   - Async LDS operations
 *   - Unified s_waitcnt instruction
 *
 * - Strix / Strix Halo / Krackan (RDNA 3.5 gfx115x) APU iGPU:
 *   - Wave32 native, WMMA (FP16/BF16/INT8), VOPD same-wave dual-issue VALU
 *   - No CDNA-style FP8/BF8 cvt builtins; scalar FP in scalar unit; s_singleuse_vdst hint
 *
 * - RDNA4 Consumer (gfx1200/gfx1201 - RX 9060 Series, RX 9070 Series):
 *   - FP8/BF8 OCP format conversions
 *   - WMMA matrix operations (FP16/BF16/INT8)
 *   - Async LDS operations
 *   - AI inference accelerators
 *   - Unified s_waitcnt instruction (still available)
 *
 * Features tested:
 * - Warp/Wave operations (shuffle, ballot, permute)
 * - FP16/BF16 packed operations
 * - FP32/FP64 operations
 * - Atomic operations (global, LDS, flat)
 * - LDS operations (sync and async)
 * - Memory operations (VMEM comprehensive tests)
 * - Transcendental functions
 * - Integer operations
 * - DOT product operations
 * - MFMA operations (CDNA) / WMMA operations (RDNA4 / RDNA 3.5 where supported)
 *
 * Copyright (c) 2026 - For testing purposes
 */

#include <cmath>
#include <cstdint>
#include <cstdio>
#include <cstring>
#include <hip/hip_bf16.h>
#include <hip/hip_fp16.h>
#include <hip/hip_runtime.h>

#include "atomic_global_buffers.h"
#include "mega_kernel_device_arch.h"
#include "mega_kernel_types.h"

// ============================================================================
// Wait Instruction Compatibility Macros
// ============================================================================
//
// WAITCNT INSTRUCTION CHANGES ON RDNA4 / GFX12+ (e.g. gfx1200/gfx1201):
// ====================================================================
// The unified s_waitcnt instruction is REMOVED on RDNA4 consumer GPUs and replaced
// with split wait instructions (load/store/kmcnt/dscnt/expcnt families).
//
// OLD (pre-GFX12 / unified wait):   NEW (GFX12 RDNA4, e.g. gfx1201):
// ---------------------        -----------------------
// s_waitcnt vmcnt(0)    -->    s_wait_loadcnt 0 + s_wait_storecnt 0
// s_waitcnt lgkmcnt(0)  -->    s_wait_kmcnt 0 + s_wait_dscnt 0
// s_waitcnt expcnt(0)   -->    s_wait_expcnt 0
//
// BUILTIN EQUIVALENTS:
// --------------------
// __builtin_amdgcn_s_waitcnt(0)           --> Old unified wait
// __builtin_amdgcn_s_wait_loadcnt(0)      --> Wait for vector memory loads
// __builtin_amdgcn_s_wait_storecnt(0)     --> Wait for vector memory stores
// __builtin_amdgcn_s_wait_kmcnt(0)        --> Wait for scalar memory
// __builtin_amdgcn_s_wait_dscnt(0)        --> Wait for LDS operations
// __builtin_amdgcn_s_wait_expcnt(0)       --> Wait for exports
//
// USAGE PATTERNS:
// ---------------
// 1. Wait for all vector memory (load + store):
//    WAITCNT_VMEM_ALL()
//
// 2. Wait for LDS + scalar memory:
//    WAITCNT_LDS_SCALAR()
//
// 3. Wait for everything:
//    WAITCNT_ALL()
//
// 4. Inline assembly wait for vector memory:
//    ASM_WAITCNT_VMEM
//
// 5. Inline assembly wait for LDS:
//    ASM_WAITCNT_LDS
//

#if 0  // defined() placeholder for future architecture
    // GFX12+ ONLY: s_waitcnt is REMOVED, use split wait instructions
#    define WAITCNT_VMEM_ALL()                                                           \
        do                                                                               \
        {                                                                                \
            __builtin_amdgcn_s_wait_loadcnt(0);                                          \
            __builtin_amdgcn_s_wait_storecnt(0);                                         \
        } while(0)

    // Wait for LDS + scalar memory (kmcnt + dscnt)
#    define WAITCNT_LDS_SCALAR()                                                         \
        do                                                                               \
        {                                                                                \
            __builtin_amdgcn_s_wait_kmcnt(0);                                            \
            __builtin_amdgcn_s_wait_dscnt(0);                                            \
        } while(0)

    // Wait for all outstanding operations: VMEM (load/store) + scalar + LDS
#    define WAITCNT_ALL()                                                                \
        do                                                                               \
        {                                                                                \
            __builtin_amdgcn_s_wait_loadcnt(0);                                          \
            __builtin_amdgcn_s_wait_storecnt(0);                                         \
            __builtin_amdgcn_s_wait_kmcnt(0);                                            \
            __builtin_amdgcn_s_wait_dscnt(0);                                            \
        } while(0)

    // Inline assembly versions for GFX12+ (split waits with 0x0 syntax)
#    define ASM_WAITCNT_VMEM "s_wait_loadcnt 0x0\n\ts_wait_storecnt 0x0"
#    define ASM_WAITCNT_LDS  "s_wait_kmcnt 0x0\n\ts_wait_dscnt 0x0"
#    define ASM_WAITCNT_ALL                                                              \
        "s_wait_loadcnt 0x0\n\ts_wait_storecnt 0x0\n\ts_wait_kmcnt 0x0\n\ts_wait_dscnt " \
        "0x0"

#else
// All other architectures (CDNA, RDNA 3.5, and pre-GFX12 RDNA): unified s_waitcnt
#    define WAITCNT_VMEM_ALL()   __builtin_amdgcn_s_waitcnt(0)
#    define WAITCNT_LDS_SCALAR() __builtin_amdgcn_s_waitcnt(0)
#    define WAITCNT_ALL()        __builtin_amdgcn_s_waitcnt(0)

// Inline assembly versions for architectures with unified s_waitcnt
#    define ASM_WAITCNT_VMEM "s_waitcnt vmcnt(0)"
#    define ASM_WAITCNT_LDS  "s_waitcnt lgkmcnt(0)"
#    define ASM_WAITCNT_ALL  "s_waitcnt vmcnt(0) lgkmcnt(0)"

#endif

// Helper: Tolerance-based float comparison
__device__ inline bool
approx_equal_f32(float a, float b, float tol = 1e-5f)
{
    return fabsf(a - b) < tol * (fabsf(a) + fabsf(b) + 1.0f);
}

__device__ inline bool
approx_equal_f64(double a, double b, double tol = 1e-10)
{
    return fabs(a - b) < tol * (fabs(a) + fabs(b) + 1.0);
}

// ============================================================================
// Section 1: Warp/Wave Operations
// ============================================================================

__device__ void
test_warp_shuffle(TestResults* results, int lane_id)
{
    int value = lane_id * 10;

    // Test __shfl - broadcast from lane 0
    int  shuffled = __shfl(value, 0, warpSize);
    bool passed   = (shuffled == 0);  // lane 0 has value 0

    // Test __shfl_xor - butterfly reduction pattern
    int xor_result = __shfl_xor(value, 1, warpSize);
    passed         = passed && (xor_result == ((lane_id ^ 1) * 10));

    // Test __shfl_up
    int up_result = __shfl_up(value, 1, warpSize);
    if(lane_id > 0)
    {
        passed = passed && (up_result == (lane_id - 1) * 10);
    }

    // Test __shfl_down
    int down_result = __shfl_down(value, 1, warpSize);
    if(lane_id < warpSize - 1)
    {
        passed = passed && (down_result == (lane_id + 1) * 10);
    }

    if(lane_id == 0)
    {
        atomicAdd(&results->warp_shuffle_passed, passed ? 1 : 0);
    }
}

__device__ void
test_warp_ballot(TestResults* results, int lane_id)
{
    // Test ballot - all even lanes vote true
    unsigned long long ballot_result = __ballot(lane_id % 2 == 0);

    // Expected: 0x5555555555555555 for wave64, 0x55555555 for wave32
    unsigned long long expected =
        (warpSize == 64) ? 0x5555555555555555ULL : 0x55555555ULL;
    bool passed = (ballot_result == expected);

    // Test __any
    int any_result = __any(lane_id == 0);
    passed         = passed && (any_result != 0);

    // Test __all
    int all_result = __all(lane_id >= 0);
    passed         = passed && (all_result != 0);

    if(lane_id == 0)
    {
        atomicAdd(&results->warp_ballot_passed, passed ? 1 : 0);
    }
}

__device__ void
test_warp_permute(TestResults* results, int lane_id)
{
    int value = lane_id * 7;

    // Test ds_bpermute - read from specific lane
    int  target_lane = (lane_id + 1) % warpSize;
    int  permuted    = __builtin_amdgcn_ds_bpermute(target_lane << 2, value);
    bool passed      = (permuted == target_lane * 7);

    // Test ds_permute
    int permuted2 = __builtin_amdgcn_ds_permute(lane_id << 2, value);
    passed        = passed && (permuted2 == lane_id * 7);

    if(lane_id == 0)
    {
        atomicAdd(&results->warp_permute_passed, passed ? 1 : 0);
    }
}

__device__ void
test_warp_reduce(TestResults* results, int lane_id)
{
    // Warp-level sum reduction using shuffle
    int value = 1;  // Each lane contributes 1

    for(int offset = warpSize / 2; offset > 0; offset /= 2)
    {
        value += __shfl_xor(value, offset, warpSize);
    }

    // All lanes should have the sum (warpSize)
    bool passed = (value == warpSize);

    if(lane_id == 0)
    {
        atomicAdd(&results->warp_reduce_passed, passed ? 1 : 0);
    }
}

// ============================================================================
// Section 2: FP8/BF8 Conversion Operations (CDNA / RDNA4; RDNA 3.5 gfx115x bypassed)
// ============================================================================

__device__ void
test_fp8_conversions(TestResults* results, int lane_id)
{
    bool passed = true;

#if defined(__gfx950__)  // MI350 (gfx950/CDNA4): scaled FP8 E4M3 conversions
    float test_val = 1.5f;

    // Convert float to FP8 and back
    uint32_t fp8_packed = 0;

    // Use builtin to convert F32 to packed FP8
    fp8_packed = __builtin_amdgcn_cvt_pk_fp8_f32(test_val, test_val, fp8_packed, false);

    // Convert back to F32
    float2 converted;
    asm volatile("v_cvt_pk_f32_fp8 %0, %1" : "=v"(converted) : "v"(fp8_packed));

    // Check if round-trip is approximately correct (FP8 has limited precision)
    passed = passed && (fabsf(converted.x - test_val) < 0.1f);

#elif defined(__gfx1200__) || defined(__gfx1201__)
    // RDNA4 Consumer (RX 9060 / RX 9070 families): FP8 OCP format support
    float    test_val   = 1.5f;
    uint32_t fp8_packed = 0;

    // Use builtin for FP8 conversion (OCP format on RDNA4)
    fp8_packed = __builtin_amdgcn_cvt_pk_fp8_f32(test_val, test_val, fp8_packed, false);

    // Verify non-zero output
    passed = (fp8_packed != 0);

#elif ARCH_RDNA35_APU
    // RDNA 3.5 APU (Strix): NO FP8 conversion builtins - mark as bypassed
    if(lane_id == 0)
    {
        atomicMin(&results->fp8_convert_passed, -1);
    }
    return;

#elif defined(__gfx942__)
    // MI300 (gfx942/CDNA3): FP8 E4M3 via cvt_pk_fp8_f32 (FNUZ-style packed FP8)
    float    test_val   = 1.5f;
    uint32_t fp8_packed = 0;

    fp8_packed = __builtin_amdgcn_cvt_pk_fp8_f32(test_val, test_val, fp8_packed, false);

    // Simplified test - verify conversion produces non-zero output
    passed = (fp8_packed != 0);

#elif defined(__gfx90a__)
    // MI250 (gfx90a/CDNA2): NO FP8 support - mark as bypassed
    if(lane_id == 0)
    {
        atomicMin(&results->fp8_convert_passed, -1);
    }
    return;

#else
    // Fallback for unknown architectures - mark as bypassed
    if(lane_id == 0)
    {
        atomicMin(&results->fp8_convert_passed, -1);
    }
    return;
#endif

    if(lane_id == 0)
    {
        atomicAdd(&results->fp8_convert_passed, passed ? 1 : 0);
    }
}

__device__ void
test_bf8_conversions(TestResults* results, int lane_id)
{
    bool passed = true;

#if defined(__gfx950__)  // MI350 (gfx950/CDNA4): BF8 E5M2 conversions
    float    test_val   = 2.0f;
    uint32_t bf8_packed = 0;

    // Convert float to BF8
    bf8_packed = __builtin_amdgcn_cvt_pk_bf8_f32(test_val, test_val, bf8_packed, false);

    // Verify non-zero result
    passed = (bf8_packed != 0);

#elif defined(__gfx1200__) || defined(__gfx1201__)
    // RDNA4 Consumer (RX 9060 / RX 9070 families): BF8 OCP format support
    float    test_val   = 2.0f;
    uint32_t bf8_packed = 0;

    // Convert float to BF8
    bf8_packed = __builtin_amdgcn_cvt_pk_bf8_f32(test_val, test_val, bf8_packed, false);

    // Verify non-zero result
    passed = (bf8_packed != 0);

#elif ARCH_RDNA35_APU
    // RDNA 3.5 APU (Strix): NO BF8 conversion builtins - mark as bypassed
    if(lane_id == 0)
    {
        atomicMin(&results->bf8_convert_passed, -1);
    }
    return;

#elif defined(__gfx942__)
    // MI300 (gfx942/CDNA3): BF8 E5M2 support via cvt_pk_bf8_f32
    float    test_val   = 2.0f;
    uint32_t bf8_packed = 0;

    bf8_packed = __builtin_amdgcn_cvt_pk_bf8_f32(test_val, test_val, bf8_packed, false);

    // Verify non-zero result
    passed = (bf8_packed != 0);

#elif defined(__gfx90a__)
    // MI250 (gfx90a/CDNA2): NO BF8 support - mark as bypassed
    if(lane_id == 0)
    {
        atomicMin(&results->bf8_convert_passed, -1);
    }
    return;

#else
    // Fallback for unknown architectures - mark as bypassed
    if(lane_id == 0)
    {
        atomicMin(&results->bf8_convert_passed, -1);
    }
    return;
#endif

    if(lane_id == 0)
    {
        atomicAdd(&results->bf8_convert_passed, passed ? 1 : 0);
    }
}

// ============================================================================
// Section 3: FP16/BF16 Operations
// ============================================================================

__device__ void
test_fp16_operations(TestResults* results, int lane_id)
{
    _Float16 a = (_Float16) 1.5f;
    _Float16 b = (_Float16) 2.5f;

    // Arithmetic
    _Float16 sum  = a + b;
    _Float16 diff = b - a;
    _Float16 prod = a * b;
    _Float16 quot = b / a;

    bool passed = true;
    passed      = passed && (fabsf((float) sum - 4.0f) < 0.01f);
    passed      = passed && (fabsf((float) diff - 1.0f) < 0.01f);
    passed      = passed && (fabsf((float) prod - 3.75f) < 0.01f);
    passed      = passed && (fabsf((float) quot - 1.666f) < 0.01f);

    if(lane_id == 0)
    {
        atomicAdd(&results->fp16_convert_passed, passed ? 1 : 0);
    }
}

__device__ void
test_bf16_operations(TestResults* results, int lane_id)
{
    __hip_bfloat16 a = __float2bfloat16(1.5f);
    __hip_bfloat16 b = __float2bfloat16(2.5f);

    // Convert back and check
    float fa = __bfloat162float(a);
    float fb = __bfloat162float(b);

    bool passed = true;
    passed      = passed && (fabsf(fa - 1.5f) < 0.01f);
    passed      = passed && (fabsf(fb - 2.5f) < 0.01f);

    if(lane_id == 0)
    {
        atomicAdd(&results->bf16_convert_passed, passed ? 1 : 0);
    }
}

// ============================================================================
// Section 4: FP32/FP64 Arithmetic Operations
// ============================================================================

__device__ void
test_fp32_arithmetic(TestResults* results, int lane_id)
{
    float a = 3.14159f;
    float b = 2.71828f;

    float sum  = a + b;
    float diff = a - b;
    float prod = a * b;
    (void) prod;  // Used to verify computation happens

    // FMA operation
    float fma_result   = __fmaf_rn(a, b, 1.0f);
    float expected_fma = a * b + 1.0f;

    bool passed = true;
    passed      = passed && approx_equal_f32(sum, 5.85987f);
    passed      = passed && approx_equal_f32(diff, 0.42331f);
    passed      = passed && approx_equal_f32(fma_result, expected_fma);

    // Test special values
    float inf_val = __int_as_float(0x7F800000);  // +Inf
    float nan_val = __int_as_float(0x7FC00000);  // NaN
    passed        = passed && isinf(inf_val);
    passed        = passed && isnan(nan_val);

    if(lane_id == 0)
    {
        atomicAdd(&results->fp32_arith_passed, passed ? 1 : 0);
    }
}

__device__ void
test_fp64_arithmetic(TestResults* results, int lane_id)
{
    double a = 3.141592653589793;
    double b = 2.718281828459045;

    double sum  = a + b;
    double diff = a - b;
    (void) (a * b);  // Verify mul happens
    (void) (a / b);  // Verify div happens

    // FMA operation
    double fma_result   = __fma_rn(a, b, 1.0);
    double expected_fma = a * b + 1.0;

    bool passed = true;
    passed      = passed && approx_equal_f64(sum, 5.859874482048838);
    passed      = passed && approx_equal_f64(diff, 0.423310825130748);
    passed      = passed && approx_equal_f64(fma_result, expected_fma);

    if(lane_id == 0)
    {
        atomicAdd(&results->fp64_arith_passed, passed ? 1 : 0);
    }
}

// ============================================================================
// Section 5: Integer Operations
// ============================================================================

__device__ void
test_int32_operations(TestResults* results, int lane_id)
{
    int a = 12345;
    int b = 6789;

    int sum  = a + b;
    int diff = a - b;
    int prod = a * b;
    int quot = a / b;
    int rem  = a % b;

    // Bit operations - verify they execute
    volatile int and_val = a & b;
    volatile int or_val  = a | b;
    volatile int xor_val = a ^ b;
    volatile int not_val = ~a;
    int          shl     = a << 4;
    volatile int shr     = a >> 4;

    // Population count and leading zeros - verify they execute
    volatile int popcount = __popc(a);
    volatile int clz_val  = __clz(a);

    // Use volatile results to prevent optimization
    (void) and_val;
    (void) or_val;
    (void) xor_val;
    (void) not_val;
    (void) shr;
    (void) popcount;
    (void) clz_val;

    bool passed = true;
    passed      = passed && (sum == 19134);
    passed      = passed && (diff == 5556);
    passed      = passed && (prod == 83810205);
    passed      = passed && (quot == 1);
    passed      = passed && (rem == 5556);
    passed      = passed && (shl == 197520);

    if(lane_id == 0)
    {
        atomicAdd(&results->int32_arith_passed, passed ? 1 : 0);
    }
}

__device__ void
test_int64_operations(TestResults* results, int lane_id)
{
    // Use simpler values that are less likely to have precision issues
    long long a = 100000000000LL;  // 100 billion
    long long b = 12345678LL;

    // Basic arithmetic
    long long sum  = a + b;
    long long diff = a - b;
    long long prod = (long long) lane_id * 1000000LL;

    bool passed = true;

    // Check sum: 100000000000 + 12345678 = 100012345678
    passed = passed && (sum == 100012345678LL);

    // Check diff: 100000000000 - 12345678 = 99987654322
    passed = passed && (diff == 99987654322LL);

    // Check that product is non-negative (basic sanity)
    passed = passed && (prod >= 0);

    // Use volatile to prevent optimization
    volatile long long sink = sum + diff + prod;
    (void) sink;

    if(lane_id == 0)
    {
        atomicAdd(&results->int64_arith_passed, passed ? 1 : 0);
    }
}

// ============================================================================
// Section 6: Atomic Operations (HW FP64 atomics: CDNA2+ gfx90a/gfx942/gfx950;
// bypassed on RDNA consumer gfx12 and RDNA 3.5 gfx115x — no global/DS FP64 atomics)
// ============================================================================

__device__ void
test_atomic_add_f32(float* global_mem, TestResults* results, int lane_id)
{
    // Each lane atomically adds 1.0
    (void) atomicAdd(global_mem, 1.0f);  // Ignore old value

    __syncthreads();

    if(lane_id == 0)
    {
        // After all threads in block add, value should be >= blockDim.x
        bool passed = (*global_mem >= (float) blockDim.x);
        atomicAdd(&results->atomic_add_f32_passed, passed ? 1 : 0);
    }
}

__device__ void
test_atomic_add_f64([[maybe_unused]] double* global_mem, TestResults* results,
                    int lane_id)
{
#if defined(__gfx90a__) || defined(__gfx942__) ||                                        \
    defined(__gfx950__)  // CDNA architectures: Each lane atomically adds 1.0 using
                         // hardware FP64 atomics
    (void) atomicAdd(global_mem, 1.0);  // Ignore old value

    __syncthreads();

    if(lane_id == 0)
    {
        bool passed = (*global_mem >= (double) blockDim.x);
        atomicAdd(&results->atomic_add_f64_passed, passed ? 1 : 0);
    }
#else
    // RDNA4 consumer GPUs don't have hardware FP64 atomics - mark as bypassed
    if(lane_id == 0)
    {
        atomicMin(&results->atomic_add_f64_passed, -1);
    }
#endif
}

__device__ void
test_atomic_min_max_f32(TestResults* results, int lane_id)
{
    __shared__ unsigned int shared_atomic_min_max_f32;
    float                   my_val = (float) lane_id + 1.0f;

    // ---------------------------------------------------------------------
    // Atomic MIN test: all lanes compete to set the minimum
    // Use shared memory so each block gets an independent scratch value.
    if(threadIdx.x == 0)
    {
        shared_atomic_min_max_f32 = __float_as_uint(1.0e9f);
    }
    __syncthreads();

    (void) atomicMin(&shared_atomic_min_max_f32, __float_as_uint(my_val));

    __syncthreads();

    if(lane_id == 0)
    {
        // Minimum should be 1.0 (from lane 0)
        float min_val    = __uint_as_float(shared_atomic_min_max_f32);
        bool  min_passed = fabsf(min_val - 1.0f) < 1.0e-6f;
        atomicAdd(&results->atomic_min_f32_passed, min_passed ? 1 : 0);
    }

    // ---------------------------------------------------------------------
    // Atomic MAX test: all threads in block compete to set the maximum
    // Reuse the same per-block shared scratch value for the max phase.
    if(threadIdx.x == 0)
    {
        shared_atomic_min_max_f32 = __float_as_uint(0.0f);
    }
    __syncthreads();

    (void) atomicMax(&shared_atomic_min_max_f32, __float_as_uint(my_val));

    __syncthreads();

    if(threadIdx.x == 0)
    {
        float max_val = __uint_as_float(shared_atomic_min_max_f32);
        // All threads in block participate with my_val = lane_id + 1
        // Since lane_id = threadIdx.x % warpSize, and warpSize = 32:
        // Maximum my_val across all threads is warpSize (when lane_id = 31)
        // Verify max_val is at least 1.0 and at most warpSize
        bool max_passed = (max_val >= 1.0f) && (max_val <= (float) warpSize + 0.1f);
        atomicAdd(&results->atomic_max_f32_passed, max_passed ? 1 : 0);
    }
}

__device__ void
test_atomic_int(int* global_mem, TestResults* results, int lane_id)
{
    // Test various integer atomics - use volatile to ensure execution
    (void) atomicAdd(global_mem, 1);
    (void) atomicSub(global_mem + 1, 1);
    (void) atomicAnd(global_mem + 2, 0xFF);
    (void) atomicOr(global_mem + 3, lane_id);
    (void) atomicXor(global_mem + 4, 1);
    (void) atomicMin(global_mem + 5, lane_id);
    (void) atomicMax(global_mem + 6, lane_id);

    __syncthreads();

    if(lane_id == 0)
    {
        bool passed = true;  // Verify no crash and results make sense
        passed      = passed && (global_mem[0] >= (int) blockDim.x);  // All lanes added 1
        atomicAdd(&results->atomic_int_passed, passed ? 1 : 0);
    }
}

// ============================================================================
// Section 6.5: Atomic Operations - Comprehensive Inline ASM Tests
// ============================================================================
// Global buffer sizes and slot layout must match main.cpp / atomic_global_buffers.h.
//
// AMD GPU ATOMIC INSTRUCTION TYPES:
// ==================================
// This section tests all atomic operation types using inline assembly to
// verify the hardware atomic paths are functioning correctly.
//
// ATOMIC INSTRUCTION CATEGORIES:
// ------------------------------
//
// 1. GLOBAL ATOMICS (global_atomic_*):
//    Integer: global_atomic_add, global_atomic_sub, global_atomic_and,
//             global_atomic_or, global_atomic_xor, global_atomic_swap,
//             global_atomic_smin, global_atomic_smax, global_atomic_umin,
//             global_atomic_umax, global_atomic_inc, global_atomic_dec
//    Float:   global_atomic_add_f32, global_atomic_add_f64,
//             global_atomic_min_f64, global_atomic_max_f64
//    Packed:  global_atomic_pk_add_f16, global_atomic_pk_add_bf16
//    CAS:     global_atomic_cmpswap, global_atomic_cmpswap_x2
//
// 2. FLAT ATOMICS (flat_atomic_*):
//    Same operations as global, but with unified addressing
//
// 3. DS (LDS) ATOMICS (ds_*):
//    Integer: ds_add_u32/u64, ds_sub_u32, ds_and_b32/b64, ds_or_b32/b64,
//             ds_xor_b32/b64, ds_min_i32/i64/u32/u64, ds_max_i32/i64/u32/u64,
//             ds_inc_u32, ds_dec_u32
//    Float:   ds_add_f32/f64, ds_min_f32/f64, ds_max_f32/f64
//    CAS:     ds_cmpst_b32/b64
//
// ARCHITECTURE ATOMIC SUPPORT:
// ----------------------------
// | Atomic Type              | gfx90a | gfx942 | gfx950 | gfx115x | gfx1200/gfx1201 |
// |--------------------------|--------|--------|--------|---------|-----------------|
// | INT32/INT64 atomics      |   ✅   |   ✅   |   ✅   |    ✅   |       ✅        |
// | FP32 atomic add          |   ✅   |   ✅   |   ✅   |    ✅   |       ✅        |
// | FP64 atomic add/min/max  |   ✅   |   ✅   |   ✅   |    ❌   |       ❌        |
// | Packed FP16 atomic add   |   ✅   |   ✅   |   ✅   |    ✅   |       ✅        |
// | Packed BF16 atomic add   |   ❌   |   ✅   |   ✅   |    ✅   |       ✅        |
// | DS FP32 atomics          |   ✅   |   ✅   |   ✅   |    ✅   |       ✅        |
// | DS FP64 atomics          |   ✅   |   ✅   |   ✅   |    ❌   |       ❌        |
//
// NAMING CONVENTION:
// ------------------
// CDNA (gfx9x): global_atomic_add, flat_atomic_add (HIP atomics in this sample for gfx115x too)
// RDNA4 (gfx120x): global_atomic_add_u32, flat_atomic_add_u32 (explicit type suffix in asm paths)
//

// Shared memory for DS atomic tests
__shared__ int    ds_atomic_int_test[64];
__shared__ float  ds_atomic_f32_test[64];
__shared__ double ds_atomic_f64_test[32];

__device__ void
test_atomic_operations_asm([[maybe_unused]] int*         global_int_ptr,
                           [[maybe_unused]] float*       global_f32_ptr,
                           [[maybe_unused]] double*      global_f64_ptr,
                           [[maybe_unused]] TestResults* results,
                           [[maybe_unused]] int lane_id, [[maybe_unused]] int tid)
{
    bool global_int_passed = true;
    bool global_f32_passed = true;
    [[maybe_unused]] bool global_f64_passed = true;
    bool flat_int_passed   = true;
    bool ds_int_passed     = true;
    bool ds_f32_passed     = true;
    [[maybe_unused]] bool ds_f64_passed = true;
    bool pk_f16_passed     = true;
    bool pk_bf16_passed    = true;
    bool cas_passed        = true;

    // ========================================================================
    // TEST 1: GLOBAL INTEGER ATOMICS
    // ========================================================================
    // global_atomic_add/sub/and/or/xor/min/max/swap
    //
#if defined(__gfx90a__) || defined(__gfx942__) || defined(__gfx950__) || ARCH_RDNA35_APU
    // CDNA and RDNA 3.5 APU: Use HIP atomics (inline ASM syntax differs)
    {
        int old_val       = atomicAdd(&global_int_ptr[tid % 64], 1);
        global_int_passed = (old_val >= 0);

        atomicMax(&global_int_ptr[(tid % 64) + 64], tid);
    }

#elif defined(__gfx1200__) ||                                                            \
    defined(__gfx1201__)  // RDNA4: global_atomic_add_u32 vdst:opt, vaddr, vdata, off
                          // [th:TH_ATOMIC_RETURN]
    {
        int old_val = 0;
        int add_val = 1;

        // RDNA4 uses explicit type suffix and different return syntax
        asm volatile("global_atomic_add_u32 %0, %1, %2, off "
                     "th:TH_ATOMIC_RETURN\n\t" ASM_WAITCNT_VMEM
                     : "=v"(old_val)
                     : "v"(&global_int_ptr[tid % 64]), "v"(add_val)
                     : "memory");

        global_int_passed = (old_val >= 0);

        // global_atomic_max_i32 (signed max)
        int max_val = tid;
        asm volatile("global_atomic_max_i32 %0, %1, off\n\t" ASM_WAITCNT_VMEM
                     :
                     : "v"(&global_int_ptr[(tid % 64) + 64]), "v"(max_val)
                     : "memory");
    }
#else
    global_int_passed = true;  // Unknown arch - skip
#endif

    // ========================================================================
    // TEST 2: GLOBAL FP32 ATOMICS
    // ========================================================================
    // global_atomic_add_f32 - supported on all target architectures
    //
#if defined(__gfx90a__) || defined(__gfx942__) || defined(__gfx950__) || ARCH_RDNA35_APU
    // CDNA and RDNA 3.5 APU: Use HIP atomics (inline ASM syntax differs)
    {
        float old_val     = atomicAdd(&global_f32_ptr[tid % 64], 1.0f);
        global_f32_passed = global_f32_passed && isfinite(old_val);
    }

#elif defined(__gfx1200__) ||                                                            \
    defined(__gfx1201__)  // RDNA4: global_atomic_add_f32 vdst:opt, vaddr, vdata, off
    {
        float old_val = 0.0f;
        float add_val = 1.0f;

        asm volatile("global_atomic_add_f32 %0, %1, %2, off "
                     "th:TH_ATOMIC_RETURN\n\t" ASM_WAITCNT_VMEM
                     : "=v"(old_val)
                     : "v"(&global_f32_ptr[tid % 64]), "v"(add_val)
                     : "memory");

        global_f32_passed = global_f32_passed && isfinite(old_val);
    }
#else
    global_f32_passed = true;
#endif

    // ========================================================================
    // TEST 3: GLOBAL FP64 ATOMICS (CDNA2+ only; gfx90a/gfx942/gfx950)
    // ========================================================================
    // global_atomic_add_f64, global_atomic_min_f64, global_atomic_max_f64
    // Available on: gfx90a, gfx942, gfx950,
    // NOT available on: gfx1200, gfx1201 (RDNA4 consumer)
    //
#if defined(__gfx90a__) || defined(__gfx942__) || defined(__gfx950__)
    // CDNA: Use HIP atomics for FP64 (inline ASM too complex)
    {
        double old_val    = unsafeAtomicAdd(&global_f64_ptr[tid % 32], 1.0);
        global_f64_passed = global_f64_passed && isfinite(old_val);

        // Min/max also via HIP
        unsafeAtomicMin(&global_f64_ptr[(tid % 32) + 32], 1.0);
    }

#elif 0  // defined() placeholder for future architecture
    // Future arch: FP64 atomics via global_atomic_add_f64
    {
        double old_val = 0.0;
        double add_val = 1.0;

        asm volatile("global_atomic_add_f64 %0, %1, %2, off "
                     "th:TH_ATOMIC_RETURN\n\t" ASM_WAITCNT_VMEM
                     : "=v"(old_val)
                     : "v"(&global_f64_ptr[tid % 32]), "v"(add_val)
                     : "memory");

        global_f64_passed = global_f64_passed && isfinite(old_val);
    }

#elif defined(__gfx1200__) || defined(__gfx1201__)
    // RDNA4 consumer (gfx12): no hardware FP64 global atomics — bypass
    if(lane_id == 0 && tid < 32)
    {
        atomicMin(&results->atomic_global_f64_passed, -1);
    }
    global_f64_passed = true;  // Don't count as failure
#elif ARCH_RDNA35_APU
    // RDNA 3.5 APU (gfx115x): no hardware FP64 global atomics — bypass
    if(lane_id == 0 && tid < 32)
    {
        atomicMin(&results->atomic_global_f64_passed, -1);
    }
    global_f64_passed = true;  // Don't count as failure
#else
    global_f64_passed = true;
#endif

    // ========================================================================
    // TEST 4: FLAT INTEGER ATOMICS
    // ========================================================================
    // flat_atomic_add/sub/min/max - unified addressing
    //
#if defined(__gfx90a__) || defined(__gfx942__) || defined(__gfx950__) || ARCH_RDNA35_APU
    // CDNA and RDNA 3.5 APU: Use HIP atomics (inline ASM syntax differs)
    {
        int old_val     = atomicAdd(&global_int_ptr[(tid % 64) + 128], 1);
        flat_int_passed = (old_val >= 0);
    }

#elif defined(__gfx1200__) ||                                                            \
    defined(__gfx1201__)  // RDNA4: flat_atomic_add_u32 vdst:opt, vaddr, vdata
    {
        int old_val = 0;
        int add_val = 1;

        asm volatile(
            "flat_atomic_add_u32 %0, %1, %2 th:TH_ATOMIC_RETURN\n\t" ASM_WAITCNT_ALL
            : "=v"(old_val)
            : "v"(&global_int_ptr[(tid % 64) + 128]), "v"(add_val)
            : "memory");

        flat_int_passed = (old_val >= 0);
    }
#else
    flat_int_passed = true;
#endif

    // ========================================================================
    // TEST 5: DS (LDS) INTEGER ATOMICS
    // ========================================================================
    // ds_add_u32, ds_min_i32, ds_max_u32, ds_and_b32, ds_or_b32
    // Same syntax on all architectures
    //
    {
        // Initialize LDS
        if(tid < 64)
        {
            ds_atomic_int_test[tid] = 0;
        }
        __syncthreads();

        // Use HIP atomics for LDS instead of inline ASM to avoid pointer issues
        int old_val   = atomicAdd(&ds_atomic_int_test[tid % 64], 1);
        ds_int_passed = (old_val >= 0);

        atomicMin(&ds_atomic_int_test[tid % 64], tid);
        atomicMax((unsigned int*) &ds_atomic_int_test[tid % 64], (unsigned int) tid);

        __syncthreads();
    }

    // ========================================================================
    // TEST 6: DS (LDS) FP32 ATOMICS
    // ========================================================================
    // ds_add_f32, ds_min_f32, ds_max_f32
    //
    {
        if(tid < 64)
        {
            ds_atomic_f32_test[tid] = 0.0f;
        }
        __syncthreads();

        // Use HIP atomics for LDS FP32
        atomicAdd(&ds_atomic_f32_test[tid % 64], 1.0f);

        // Min/max via reinterpret cast
        float min_f32_val = (float) tid * 0.1f;
        atomicMin((unsigned int*) &ds_atomic_f32_test[tid % 64],
                  __float_as_uint(min_f32_val));

        __syncthreads();
        ds_f32_passed = !isnan(ds_atomic_f32_test[tid % 64]);
    }

    // ========================================================================
    // TEST 7: DS (LDS) FP64 ATOMICS (CDNA2+ only; gfx90a/gfx942/gfx950)
    // ========================================================================
    // ds_add_f64, ds_min_f64, ds_max_f64
    //
#if defined(__gfx90a__) || defined(__gfx942__) || defined(__gfx950__)
    {
        if(tid < 32)
        {
            ds_atomic_f64_test[tid] = 0.0;
        }
        __syncthreads();

        // Use unsafeAtomicAdd for LDS FP64 (CDNA architectures)
        unsafeAtomicAdd(&ds_atomic_f64_test[tid % 32], 1.0);

        __syncthreads();
        ds_f64_passed = !isnan(ds_atomic_f64_test[tid % 32]);
    }
#else
    // RDNA consumer (gfx12 + gfx115x APU): no DS FP64 atomics — bypass
    if(lane_id == 0 && tid < 32)
    {
        atomicMin(&results->atomic_ds_f64_passed, -1);
    }
    ds_f64_passed = true;
#endif

    // ========================================================================
    // TEST 8: PACKED FP16 ATOMICS (gfx90a+ CDNA; RDNA 3.5 gfx115x via HIP; gfx12 asm)
    // ========================================================================
    // global_atomic_pk_add_f16 - Atomically add two FP16 values
    //
#if defined(__gfx90a__) || defined(__gfx942__) || defined(__gfx950__) || ARCH_RDNA35_APU
    // CDNA and RDNA 3.5 APU: Use HIP atomics for packed FP16
    {
        // Use regular atomicAdd on reinterpreted pointer
        unsigned int packed_f16 = 0x3C003C00;
        unsigned int old_packed =
            atomicAdd((unsigned int*) &global_f32_ptr[(tid % 64) + 64], packed_f16);
        pk_f16_passed = true;
        (void) old_packed;
    }
#elif defined(__gfx1200__) || defined(__gfx1201__)  // RDNA4: global_atomic_pk_add_f16
    {
        unsigned int packed_f16 = 0x3C003C00;
        unsigned int old_packed = 0;

        asm volatile("global_atomic_pk_add_f16 %0, %1, %2, off "
                     "th:TH_ATOMIC_RETURN\n\t" ASM_WAITCNT_VMEM
                     : "=v"(old_packed)
                     : "v"((unsigned int*) &global_f32_ptr[(tid % 64) + 64]),
                       "v"(packed_f16)
                     : "memory");

        pk_f16_passed = true;
    }
#else
    // Packed FP16 global atomic not available for this ISA — bypass (not a failure).
    if(lane_id == 0)
    {
        atomicMin(&results->atomic_pk_f16_passed, -1);
    }
    pk_f16_passed = false;
#endif

    // ========================================================================
    // TEST 9: PACKED BF16 ATOMICS (gfx942+, gfx115x, gfx1200/gfx1201)
    // ========================================================================
    // global_atomic_pk_add_bf16 - Atomically add two BF16 values
    // Available on: gfx942, gfx950, gfx115x, gfx1200, gfx1201
    // NOT available on: gfx90a
    //
#if defined(__gfx942__) || defined(__gfx950__) || ARCH_RDNA35_APU
    // CDNA3+ and RDNA 3.5 APU: Use HIP atomics for packed BF16
    {
        unsigned int packed_bf16 = 0x3F803F80;
        unsigned int old_packed =
            atomicAdd((unsigned int*) &global_f32_ptr[(tid % 64) + 128], packed_bf16);
        pk_bf16_passed = true;
        (void) old_packed;
    }
#elif defined(__gfx1200__) || defined(__gfx1201__)  // RDNA4: global_atomic_pk_add_bf16
    {
        unsigned int packed_bf16 = 0x3F803F80;
        unsigned int old_packed  = 0;

        asm volatile("global_atomic_pk_add_bf16 %0, %1, %2, off "
                     "th:TH_ATOMIC_RETURN\n\t" ASM_WAITCNT_VMEM
                     : "=v"(old_packed)
                     : "v"((unsigned int*) &global_f32_ptr[(tid % 64) + 128]),
                       "v"(packed_bf16)
                     : "memory");

        pk_bf16_passed = true;
    }
#elif defined(__gfx90a__)
    // MI200: No packed BF16 atomics — bypass (atomicAdd skipped for gfx90a below).
    if(lane_id == 0)
    {
        atomicMin(&results->atomic_pk_bf16_passed, -1);
    }
    pk_bf16_passed = false;
#else
    if(lane_id == 0)
    {
        atomicMin(&results->atomic_pk_bf16_passed, -1);
    }
    pk_bf16_passed = false;
#endif

    // ========================================================================
    // TEST 10: COMPARE-AND-SWAP (CAS) ATOMICS
    // ========================================================================
    // global_atomic_cmpswap - Atomic compare and swap
    //
#if defined(__gfx90a__) || defined(__gfx942__) || defined(__gfx950__) || ARCH_RDNA35_APU
    // CDNA and RDNA 3.5 APU: Use HIP atomicCAS
    {
        int compare_val = 0;
        int swap_val    = tid + 1;
        int old_val = atomicCAS(&global_int_ptr[(tid % 64) + 192], compare_val, swap_val);
        cas_passed  = true;
        (void) old_val;
    }
#elif defined(__gfx1200__) || defined(__gfx1201__)
    {
        int compare_val = 0;
        int swap_val    = tid + 1;
        int old_val     = 0;

        long long cas_data = ((long long) swap_val << 32) | (unsigned int) compare_val;

        asm volatile("global_atomic_cmpswap_b32 %0, %1, %2, off "
                     "th:TH_ATOMIC_RETURN\n\t" ASM_WAITCNT_VMEM
                     : "=v"(old_val)
                     : "v"(&global_int_ptr[(tid % 64) + 192]), "v"(cas_data)
                     : "memory");

        cas_passed = true;
    }
#else
    // CAS path not implemented for this ISA — bypass.
    if(lane_id == 0)
    {
        atomicMin(&results->atomic_cas_passed, -1);
    }
    cas_passed = false;
#endif

    // Record results
    if(lane_id == 0)
    {
        atomicAdd(&results->atomic_global_int_passed, global_int_passed ? 1 : 0);
        atomicAdd(&results->atomic_global_f32_passed, global_f32_passed ? 1 : 0);
#if !defined(__gfx1200__) && !defined(__gfx1201__) && !ARCH_RDNA35_APU
        atomicAdd(&results->atomic_global_f64_passed, global_f64_passed ? 1 : 0);
#endif
        atomicAdd(&results->atomic_flat_int_passed, flat_int_passed ? 1 : 0);
        atomicAdd(&results->atomic_ds_int_passed, ds_int_passed ? 1 : 0);
        atomicAdd(&results->atomic_ds_f32_passed, ds_f32_passed ? 1 : 0);
#if !defined(__gfx1200__) && !defined(__gfx1201__) && !ARCH_RDNA35_APU
        atomicAdd(&results->atomic_ds_f64_passed, ds_f64_passed ? 1 : 0);
#endif
        atomicAdd(&results->atomic_pk_f16_passed, pk_f16_passed ? 1 : 0);
#if !defined(__gfx90a__)
        atomicAdd(&results->atomic_pk_bf16_passed, pk_bf16_passed ? 1 : 0);
#endif
        atomicAdd(&results->atomic_cas_passed, cas_passed ? 1 : 0);
    }
}

// ============================================================================
// Section 7: LDS (Shared Memory) Operations
// ============================================================================

// Use extern shared memory for dynamic sizing based on block size
extern __shared__ char shared_mem[];

__device__ void
test_lds_operations(TestResults* results, int lane_id, int tid)
{
    // Dynamically allocate LDS arrays from shared memory
    // Layout: [float array][int array][double array][atomic counter]
    float*  lds_float  = (float*) shared_mem;
    int*    lds_int    = (int*) &lds_float[blockDim.x];
    double* lds_double = (double*) &lds_int[blockDim.x];
    // Use a separate int at the end for atomic counter (after double array)
    int* lds_atomic_counter = (int*) &lds_double[blockDim.x / 2];

    // Write to LDS - now safe for any block size
    lds_float[tid] = (float) tid * 1.5f;
    lds_int[tid]   = tid * 3;

    // For double, only use if we have enough threads (avoid wasting memory for small
    // blocks)
    if(static_cast<unsigned>(tid) < blockDim.x / 2u)
    {
        lds_double[tid] = (double) tid * 2.5;
    }

    __syncthreads();

    // Read back and verify - track separately from atomic test
    bool lds_rw_passed = true;
    lds_rw_passed = lds_rw_passed && approx_equal_f32(lds_float[tid], (float) tid * 1.5f);
    lds_rw_passed = lds_rw_passed && (lds_int[tid] == tid * 3);

    if(static_cast<unsigned>(tid) < blockDim.x / 2u)
    {
        lds_rw_passed =
            lds_rw_passed && approx_equal_f64(lds_double[tid], (double) tid * 2.5);
    }

    // LDS atomic test - use dedicated counter separate from data arrays
    if(tid == 0)
    {
        *lds_atomic_counter = 0;
    }
    __syncthreads();

    atomicAdd(lds_atomic_counter, 1);
    __syncthreads();

    // Check atomic result - separate from read/write verification
    bool lds_atomic_test_passed = (*lds_atomic_counter >= (int) blockDim.x);

    if(lane_id == 0)
    {
        // Report LDS read/write tests
        atomicAdd(&results->lds_load_passed, lds_rw_passed ? 1 : 0);
        atomicAdd(&results->lds_store_passed, lds_rw_passed ? 1 : 0);
        // Report LDS atomic test separately
        atomicAdd(&results->lds_atomic_passed, lds_atomic_test_passed ? 1 : 0);
    }
}

// ============================================================================
// Section 8: Transcendental Functions
// ============================================================================

__device__ void
test_transcendental(TestResults* results, int lane_id)
{
    float x = 1.0f;

    // Sine/Cosine
    float sin_val  = sinf(x);
    float cos_val  = cosf(x);
    bool  sin_pass = approx_equal_f32(sin_val, 0.8414709848f, 1e-4f);
    bool  cos_pass = approx_equal_f32(cos_val, 0.5403023059f, 1e-4f);

    // Exponential and logarithm
    float exp_val  = expf(x);
    float log_val  = logf(exp_val);
    bool  exp_pass = approx_equal_f32(exp_val, 2.718281828f, 1e-4f);
    bool  log_pass = approx_equal_f32(log_val, x, 1e-4f);

    // Square root
    float sqrt_val  = sqrtf(4.0f);
    bool  sqrt_pass = approx_equal_f32(sqrt_val, 2.0f);

    // Reciprocal
    float rcp_val  = 1.0f / x;
    bool  rcp_pass = approx_equal_f32(rcp_val, 1.0f);

    // Reciprocal square root
    float rsq_val  = rsqrtf(4.0f);
    bool  rsq_pass = approx_equal_f32(rsq_val, 0.5f);

    if(lane_id == 0)
    {
        atomicAdd(&results->trans_sin_passed, sin_pass ? 1 : 0);
        atomicAdd(&results->trans_cos_passed, cos_pass ? 1 : 0);
        atomicAdd(&results->trans_exp_passed, exp_pass ? 1 : 0);
        atomicAdd(&results->trans_log_passed, log_pass ? 1 : 0);
        atomicAdd(&results->trans_sqrt_passed, sqrt_pass ? 1 : 0);
        atomicAdd(&results->trans_rcp_passed, rcp_pass ? 1 : 0);
        atomicAdd(&results->trans_rsq_passed, rsq_pass ? 1 : 0);
    }
}

// ============================================================================
// Section 9: Memory Operations
// ============================================================================

__device__ void
test_global_memory(float* input, float* output, TestResults* results, int lane_id,
                   int tid)
{
    // Coalesced load
    float val = input[tid];

    // Some computation
    val = val * 2.0f + 1.0f;

    // Coalesced store
    output[tid] = val;

    __threadfence();

    if(lane_id == 0)
    {
        atomicAdd(&results->global_load_passed, 1);
        atomicAdd(&results->global_store_passed, 1);
    }
}

// ============================================================================
// Section 9.5: VMEM (Vector Memory) Operations - Comprehensive Memory Tests
// ============================================================================
//
// AMD GPU MEMORY INSTRUCTION TYPES:
// ==================================
// This section tests all vector memory (VMEM) instruction types using inline
// assembly to ensure the hardware memory paths are exercised.
//
// MEMORY INSTRUCTION CATEGORIES:
// ------------------------------
//
// 1. FLAT MEMORY (flat_*):
//    - Unified addressing mode for global/LDS/scratch
//    - Uses 64-bit virtual addresses
//    - Instructions: flat_load_b32, flat_load_b64, flat_load_b128
//                    flat_store_b32, flat_store_b64, flat_store_b128
//    - Data widths: b8 (byte), b16 (short), b32 (dword), b64, b128
//
// 2. GLOBAL MEMORY (global_*):
//    - Direct global memory access with SGPR base + VGPR offset
//    - Instructions: global_load_b32, global_load_b64, global_load_b128
//                    global_store_b32, global_store_b64, global_store_b128
//    - Async variants: global_load_async_to_lds_b* (gfx950+)
//
// 3. BUFFER MEMORY (buffer_*):
//    - Descriptor-based memory access (128-bit resource descriptor)
//    - Instructions: buffer_load_dword, buffer_load_dwordx2/x3/x4
//                    buffer_store_dword, buffer_store_dwordx2/x3/x4
//    - Format variants: buffer_load_format_x/xy/xyz/xyzw
//
// 4. SCRATCH MEMORY (scratch_*):
//    - Private per-thread stack memory
//    - Instructions: scratch_load_b32, scratch_load_b64
//                    scratch_store_b32, scratch_store_b64
//
// 5. LDS (Local Data Share) (ds_*):
//    - Shared memory within workgroup
//    - Instructions: ds_read_b32, ds_read_b64, ds_read_b128
//                    ds_write_b32, ds_write_b64, ds_write_b128
//    - Atomic variants: ds_add_u32, ds_min_i32, ds_max_u32, etc.
//
// ARCHITECTURE DIFFERENCES:
// -------------------------
// | Instruction Type      | gfx90a | gfx942 | gfx950 | gfx115x | gfx1200/gfx1201 |
// |-----------------------|--------|--------|--------|---------|-----------------|
// | flat_load/store_b*    |   ✓    |   ✓    |   ✓    |    ✓    |       ✓         |
// | global_load/store_b*  |   ✓    |   ✓    |   ✓    |    ✓    |       ✓         |
// | buffer_load/store_*   |   ✓    |   ✓    |   ✓    |    ✓    |       ✓         |
// | scratch_load/store_*  |   ✓    |   ✓    |   ✓    |    ✓    |       ✓         |
// | ds_read/write_*       |   ✓    |   ✓    |   ✓    |    ✓    |       ✓         |
// | global_load_async_*   |   ✗    |   ⚠    |   ✓    |    ⚠    |       ✓         |
// | tensor_load_to_lds    |   ✗    |   ✗    |   ✗    |    ✗    |       ✓         |
//
// NAMING CONVENTION CHANGES (GFX12-class RDNA — gfx120x and RDNA 3.5 gfx115x in this sample):
// -----------------------------------------------------------------------------------------
// CDNA (gfx9x):                    RDNA4 / RDNA 3.5 APU (flat/global *_b32 style):
// flat_load_dword                  flat_load_b32
// flat_load_dwordx2                flat_load_b64
// flat_load_dwordx4                flat_load_b128
// global_load_dword                global_load_b32
// buffer_load_dword                buffer_load_b32
//

// Shared memory for LDS tests
__shared__ int vmem_lds_test_buffer[256];

__device__ void
test_vmem_operations([[maybe_unused]] float*       global_ptr,
                     [[maybe_unused]] TestResults* results, [[maybe_unused]] int lane_id,
                     [[maybe_unused]] int tid, [[maybe_unused]] int buffer_size,
                     [[maybe_unused]] hipTextureObject_t tex_obj,
                     [[maybe_unused]] hipSurfaceObject_t surf_obj)
{
    // Host pass of hipcc has no target arch macros; device pass uses in_bounds in guarded VMEM.
    [[maybe_unused]] const bool in_bounds = (buffer_size > 0) && (tid < buffer_size);

    bool flat_passed    = true;
    bool global_passed  = true;
    bool buffer_passed  = true;
    bool scratch_passed = true;
    bool lds_passed     = true;

    // ========================================================================
    // TEST 1: FLAT MEMORY OPERATIONS
    // ========================================================================
    // Flat addressing uses unified virtual addresses that can point to
    // global memory, LDS, or scratch space. The hardware determines the
    // actual memory space based on the address range.
    //
    // CDNA (gfx9x) syntax: flat_load_dword, flat_store_dword
    // RDNA4 (gfx12) syntax: flat_load_b32, flat_store_b32
    //
#if defined(__gfx90a__) || defined(__gfx942__) || defined(__gfx950__)
    // CDNA architectures: flat_load_dword / flat_store_dword
    if(in_bounds)
    {
        float test_val   = (float) tid + 1.0f;
        float loaded_val = 0.0f;

        // Write test value to global memory first (using regular store)
        global_ptr[tid] = test_val;
        __threadfence();

        // Flat load: Load from flat address space
        // flat_load_dword vdst, vaddr [offset:N] [glc] [slc]
        asm volatile("flat_load_dword %0, %1\n\t" ASM_WAITCNT_VMEM
                     : "=v"(loaded_val)
                     : "v"(&global_ptr[tid])
                     : "memory");

        flat_passed = (fabsf(loaded_val - test_val) < 0.001f);

        // Flat store: Store to flat address space
        float store_val = test_val * 2.0f;
        asm volatile("flat_store_dword %0, %1\n\t" ASM_WAITCNT_VMEM
                     :
                     : "v"(&global_ptr[tid]), "v"(store_val)
                     : "memory");

        // Verify the store
        __threadfence();
        flat_passed = flat_passed && (fabsf(global_ptr[tid] - store_val) < 0.001f);
    }

#elif defined(__gfx1200__) || defined(__gfx1201__) || ARCH_RDNA35_APU  // GFX12-style names: flat_load_b32 / flat_store_b32 (gfx12 + gfx115x)
    if(in_bounds)
    {
        float test_val   = (float) tid + 1.0f;
        float loaded_val = 0.0f;

        global_ptr[tid] = test_val;
        __threadfence();

        // RDNA4 uses flat_load_b32 instead of flat_load_dword
        asm volatile("flat_load_b32 %0, %1\n\t" ASM_WAITCNT_VMEM
                     : "=v"(loaded_val)
                     : "v"(&global_ptr[tid])
                     : "memory");

        flat_passed = (fabsf(loaded_val - test_val) < 0.001f);

        float store_val = test_val * 2.0f;
        asm volatile("flat_store_b32 %0, %1\n\t" ASM_WAITCNT_VMEM
                     :
                     : "v"(&global_ptr[tid]), "v"(store_val)
                     : "memory");

        __threadfence();
        flat_passed = flat_passed && (fabsf(global_ptr[tid] - store_val) < 0.001f);
    }
#else
    flat_passed = true;  // Unknown arch - skip
#endif

    // ========================================================================
    // TEST 2: GLOBAL MEMORY OPERATIONS
    // ========================================================================
    // Global memory uses SGPR base address + VGPR offset for addressing.
    // More efficient than flat when the memory type is known.
    //
    // CDNA (gfx9x) syntax: global_load_dword, global_store_dword
    // RDNA4 (gfx12) syntax: global_load_b32, global_store_b32
    //
#if defined(__gfx90a__) || defined(__gfx942__) || defined(__gfx950__)
    // CDNA: global_load_dword / global_store_dword
    if(in_bounds)
    {
        float test_val   = (float) (tid * 3) + 0.5f;
        float loaded_val = 0.0f;

        global_ptr[tid] = test_val;
        __threadfence();

        // Global load with VGPR address (off variant)
        // global_load_dword vdst, vaddr, off [offset:N]
        asm volatile("global_load_dword %0, %1, off\n\t" ASM_WAITCNT_VMEM
                     : "=v"(loaded_val)
                     : "v"(&global_ptr[tid])
                     : "memory");

        global_passed = (fabsf(loaded_val - test_val) < 0.001f);

        float store_val = test_val + 100.0f;
        asm volatile("global_store_dword %0, %1, off\n\t" ASM_WAITCNT_VMEM
                     :
                     : "v"(&global_ptr[tid]), "v"(store_val)
                     : "memory");

        __threadfence();
        global_passed = global_passed && (fabsf(global_ptr[tid] - store_val) < 0.001f);
    }

#elif defined(__gfx1200__) || defined(__gfx1201__) || ARCH_RDNA35_APU  // GFX12-style names: global_load_b32 / global_store_b32 (gfx12 + gfx115x)
    if(in_bounds)
    {
        float test_val   = (float) (tid * 3) + 0.5f;
        float loaded_val = 0.0f;

        global_ptr[tid] = test_val;
        __threadfence();

        // RDNA4 global memory syntax
        asm volatile("global_load_b32 %0, %1, off\n\t" ASM_WAITCNT_VMEM
                     : "=v"(loaded_val)
                     : "v"(&global_ptr[tid])
                     : "memory");

        global_passed = (fabsf(loaded_val - test_val) < 0.001f);

        float store_val = test_val + 100.0f;
        asm volatile("global_store_b32 %0, %1, off\n\t" ASM_WAITCNT_VMEM
                     :
                     : "v"(&global_ptr[tid]), "v"(store_val)
                     : "memory");

        __threadfence();
        global_passed = global_passed && (fabsf(global_ptr[tid] - store_val) < 0.001f);
    }
#else
    global_passed = true;
#endif

    // ========================================================================
    // TEST 3: SCRATCH MEMORY OPERATIONS
    // ========================================================================
    // Scratch memory is private per-thread stack memory.
    // Automatically allocated by the compiler for local variables.
    //
    // Note: Direct scratch ASM is complex due to compiler-managed offsets.
    // We test scratch implicitly via local variable spilling.
    //
    {
        // Force compiler to use scratch by creating large local array
        volatile int local_array[8];
        for(int i = 0; i < 8; i++)
        {
            local_array[i] = tid * 10 + i;
        }
        __threadfence_block();

        // Read back and verify
        int sum = 0;
        for(int i = 0; i < 8; i++)
        {
            sum += local_array[i];
        }

        // Expected sum: sum(tid*10 + i) for i=0..7 = 8*tid*10 + 28
        int expected   = 8 * tid * 10 + 28;
        scratch_passed = (sum == expected);
    }

    // ========================================================================
    // TEST 4: LDS (LOCAL DATA SHARE) OPERATIONS
    // ========================================================================
    // LDS is shared memory within a workgroup.
    //
    // Note: Inline ASM for DS operations has pointer issues on some architectures.
    // We use regular LDS operations via shared memory instead.
    //
    {
        // Initialize LDS (block-local indices; every block must clear all slots)
        if(threadIdx.x < 256)
        {
            vmem_lds_test_buffer[threadIdx.x] = 0;
        }
        __syncthreads();

        int lds_val = tid * 7 + 3;

        // Write to LDS using regular shared memory access
        vmem_lds_test_buffer[tid % 256] = lds_val;
        __syncthreads();

        // Read from LDS
        int lds_read = vmem_lds_test_buffer[tid % 256];
        lds_passed   = (lds_read == lds_val);

        // Test reading from adjacent location
        if(tid % 2 == 0 && tid + 1 < 256)
        {
            int read_val0 = vmem_lds_test_buffer[tid];
            int read_val1 = vmem_lds_test_buffer[tid + 1];

            // Verify both reads
            lds_passed = lds_passed && (read_val0 == (tid * 7 + 3));
            lds_passed = lds_passed && (read_val1 == ((tid + 1) * 7 + 3));
        }

        __syncthreads();
    }

    // ========================================================================
    // TEST 5: BUFFER MEMORY OPERATIONS (descriptor-based)
    // ========================================================================
    // Buffer operations use a 128-bit resource descriptor for memory access.
    // More complex to set up but provides additional features like bounds
    // checking, format conversion, and swizzling.
    //
    // Note: Buffer descriptors require SGPR setup which is complex in inline ASM.
    // We test buffer operations implicitly via the compiler's generated code
    // for texture/buffer access patterns.
    //
    {
        // Buffer operations are typically generated by the compiler for
        // structured buffer access. We verify basic functionality via
        // the global memory tests above, which use similar hardware paths.
        buffer_passed = true;  // Implicit test via global operations
    }

#if ARCH_RDNA35_APU
    bool tex_load_passed  = true;
    bool tex_store_passed = true;

    // ========================================================================
    // TEST 6: TEXTURE LOAD (Strix / RDNA 3.5) - exercises INSTS_TEX_LOAD
    // ========================================================================
    // Use tex1Dfetch to emit image_load (TEX load).
    //
    if(tex_obj != 0 && buffer_size > 0)
    {
        int   idx       = tid % buffer_size;
        float fetched   = tex1Dfetch<float>(tex_obj, idx);
        tex_load_passed = (fabsf(fetched - (float) idx) < 0.001f);
    }
    else
    {
        tex_load_passed = true;  // Bypassed when no texture object
    }

    // ========================================================================
    // TEST 7: TEXTURE STORE (Strix / RDNA 3.5) - exercises INSTS_TEX_STORE
    // ========================================================================
    // Use surf1Dwrite to emit image_store (TEX store).
    //
    if(surf_obj != 0 && buffer_size > 0)
    {
        int   idx       = tid % buffer_size;
        float store_val = (float) tid + 1000.0f;
        surf1Dwrite(store_val, surf_obj, idx);
        tex_store_passed = true;  // Write issued successfully
    }
    else
    {
        tex_store_passed = true;  // Bypassed when no surface object
    }
#endif

    // Record results
    if(lane_id == 0)
    {
        atomicAdd(&results->vmem_flat_passed, flat_passed ? 1 : 0);
        atomicAdd(&results->vmem_global_passed, global_passed ? 1 : 0);
        atomicAdd(&results->vmem_buffer_passed, buffer_passed ? 1 : 0);
        atomicAdd(&results->vmem_scratch_passed, scratch_passed ? 1 : 0);
        atomicAdd(&results->vmem_lds_passed, lds_passed ? 1 : 0);
#if ARCH_RDNA35_APU
        atomicAdd(&results->vmem_tex_load_passed, tex_load_passed ? 1 : 0);
        atomicAdd(&results->vmem_tex_store_passed, tex_store_passed ? 1 : 0);
#else
        // TEX load/store paths are gfx115x-only; mark N/A from every block (not tid==0 only).
        atomicMin(&results->vmem_tex_load_passed, -1);
        atomicMin(&results->vmem_tex_store_passed, -1);
#endif
    }
}

// ============================================================================
// Section 10: DOT Product Operations (CDNA: DOT4/DOT8 asm; RDNA: IU-style opcodes / software sanity)
// ============================================================================

__device__ void
test_dot_products(TestResults* results, int lane_id)
{
    bool passed = true;

    // DOT4/DOT8 Product Instructions:
    // All CDNA architectures (MI250/MI300/MI350) support:
    // - V_DOT4_U32_U8: 4-element unsigned INT8 dot product
    // - V_DOT4_I32_I8: 4-element signed INT8 dot product
    // - V_DOT8_U32_U4: 8-element unsigned INT4 dot product
    // - V_DOT8_I32_I4: 8-element signed INT4 dot product
    //
    // RDNA4 and RDNA 3.5 APUs use different opcode names:
    // - V_DOT4_U32_U8 -> V_DOT4_U32_IU8
    // - V_DOT4_I32_I8 -> V_DOT4_I32_IU8
    // - V_DOT8_U32_U4 -> V_DOT8_U32_IU4
    // - V_DOT8_I32_I4 -> V_DOT8_I32_IU4

#if 0  // defined() placeholder for future architecture
    // Future arch: use new IU suffix opcodes
    unsigned int a = 0x01020304;  // [1, 2, 3, 4] packed as 4 unsigned bytes
    unsigned int b = 0x01010101;  // [1, 1, 1, 1] packed as 4 unsigned bytes

    // v_dot4_u32_iu8 is the GFX12 IU-suffixed variant
    int dot4_result;
    asm volatile("v_dot4_u32_iu8 %0, %1, %2, %3"
                 : "=v"(dot4_result)
                 : "v"(a), "v"(b), "v"(0));

    // Expected: 1*1 + 2*1 + 3*1 + 4*1 = 10
    passed = (dot4_result == 10);

#elif defined(__gfx950__) || defined(__gfx942__) || defined(__gfx90a__)
    // All CDNA architectures: Use V_DOT4_U32_U8 (4-element unsigned INT8 dot product)
    unsigned int a = 0x01020304;  // [1, 2, 3, 4] packed as 4 unsigned bytes
    unsigned int b = 0x01010101;  // [1, 1, 1, 1] packed as 4 unsigned bytes

    // v_dot4_u32_u8 is available on gfx90a, gfx942, and gfx950
    int dot4_result;
    asm volatile("v_dot4_u32_u8 %0, %1, %2, %3"
                 : "=v"(dot4_result)
                 : "v"(a), "v"(b), "v"(0));

    // Expected: 1*1 + 2*1 + 3*1 + 4*1 = 10
    passed = (dot4_result == 10);

#elif defined(__gfx1200__) || defined(__gfx1201__)
    // RDNA4 consumer (gfx120x): ISA uses IU-suffixed DOT4 encodings; use software dot sanity here
    unsigned int a = 0x01020304;  // [1, 2, 3, 4] packed as 4 unsigned bytes
    unsigned int b = 0x01010101;  // [1, 1, 1, 1] packed as 4 unsigned bytes

    int a0 = (a >> 0) & 0xFF, a1 = (a >> 8) & 0xFF, a2 = (a >> 16) & 0xFF,
        a3 = (a >> 24) & 0xFF;
    int b0 = (b >> 0) & 0xFF, b1 = (b >> 8) & 0xFF, b2 = (b >> 16) & 0xFF,
        b3          = (b >> 24) & 0xFF;
    int dot4_result = a0 * b0 + a1 * b1 + a2 * b2 + a3 * b3;
    passed          = (dot4_result == 10);

#elif ARCH_RDNA35_APU
    // RDNA 3.5 APU (gfx115x): DOT4 family per RDNA 3.5 ISA; same software sanity path as gfx120x here
    unsigned int a = 0x01020304;
    unsigned int b = 0x01010101;

    int a0 = (a >> 0) & 0xFF, a1 = (a >> 8) & 0xFF, a2 = (a >> 16) & 0xFF,
        a3 = (a >> 24) & 0xFF;
    int b0 = (b >> 0) & 0xFF, b1 = (b >> 8) & 0xFF, b2 = (b >> 16) & 0xFF,
        b3          = (b >> 24) & 0xFF;
    int dot4_result = a0 * b0 + a1 * b1 + a2 * b2 + a3 * b3;
    passed          = (dot4_result == 10);

#else
    // Fallback for other architectures - software emulation
    unsigned int a  = 0x01020304;
    unsigned int b  = 0x01010101;
    int          a0 = (a >> 0) & 0xFF, a1 = (a >> 8) & 0xFF, a2 = (a >> 16) & 0xFF,
                 a3 = (a >> 24) & 0xFF;
    int          b0 = (b >> 0) & 0xFF, b1 = (b >> 8) & 0xFF, b2 = (b >> 16) & 0xFF,
                 b3          = (b >> 24) & 0xFF;
    int          dot4_result = a0 * b0 + a1 * b1 + a2 * b2 + a3 * b3;
    passed                   = (dot4_result == 10);
#endif

    // Software-emulated DOT8 test: 8 x 8-bit dot product.
    // a8 = [1, 2, 3, 4, 5, 6, 7, 8], b8 = [1, 1, 1, 1, 1, 1, 1, 1]
    // Expected dot product = 1+2+3+4+5+6+7+8 = 36.
    unsigned long long a8          = 0x0807060504030201ULL;
    unsigned long long b8          = 0x0101010101010101ULL;
    int                dot8_result = 0;
#pragma unroll
    for(int i = 0; i < 8; ++i)
    {
        int ai = (int) ((a8 >> (8 * i)) & 0xFF);
        int bi = (int) ((b8 >> (8 * i)) & 0xFF);
        dot8_result += ai * bi;
    }
    int dot8_passed = (dot8_result == 36);

    if(lane_id == 0)
    {
        atomicAdd(&results->dot4_passed, passed ? 1 : 0);
        atomicAdd(&results->dot8_passed, dot8_passed ? 1 : 0);
    }
}

// ============================================================================
// Section 10.3: MFMA (Matrix Fused Multiply-Add) Operations (gfx90a+)
// ============================================================================
//
// MFMA INSTRUCTIONS - CDNA Matrix Cores:
// For maintained MFMA / roofline coverage, see rocprofiler-compute `src/roofline/benchmark/`
// (the older rocm-amdgpu-bench kernels repo is deprecated).
//
// This test includes BOTH inline assembly AND HIP builtin versions to ensure
// comprehensive coverage of MFMA instructions.
//
// MI250 (gfx90a/CDNA2):
//   - v_mfma_f32_32x32x2f32    / __builtin_amdgcn_mfma_f32_32x32x2f32  (FP32: 4096 ops)
//   - v_mfma_f64_16x16x4f64    / __builtin_amdgcn_mfma_f64_16x16x4f64  (FP64: 2048 ops)
//   - v_mfma_f32_32x32x8f16    / __builtin_amdgcn_mfma_f32_32x32x8f16  (FP16: 16384 ops)
//   - v_mfma_f32_32x32x4bf16   / __builtin_amdgcn_mfma_f32_32x32x4bf16 (BF16: 8192 ops)
//   - v_mfma_i32_32x32x8i8     / __builtin_amdgcn_mfma_i32_32x32x8i8   (INT8: 16384 ops)
//
// MI300 (gfx942/CDNA3) - doubled K dimensions + FP8:
//   - v_mfma_f32_32x32x8bf16_1k / __builtin_amdgcn_mfma_f32_32x32x8bf16_1k (BF16: 16384
//   ops)
//   - v_mfma_i32_32x32x16_i8    / __builtin_amdgcn_mfma_i32_32x32x16_i8    (INT8: 32768
//   ops)
//   - v_mfma_f32_32x32x16_fp8   / __builtin_amdgcn_mfma_f32_32x32x16_fp8_fp8 (FP8: 32768
//   ops)
//
// MI350 (gfx950/CDNA4) - scaled F8F6F4:
//   - v_mfma_scale_f32_32x32x64_f8f6f4 / __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4
//   (FP8/BF8/FP6/FP4: 131072 ops)
//

// Vector types for MFMA operations (layout matches CDNA MFMA builtins)
using int32_16vec = __attribute__((__vector_size__(16 * sizeof(int)))) int;
using f32_16vec   = __attribute__((__vector_size__(16 * sizeof(float)))) float;
using f64_4vec    = __attribute__((__vector_size__(4 * sizeof(double)))) double;
using bf16_2vec   = __attribute__((__vector_size__(2 * sizeof(short)))) short;
using bf16_4vec   = __attribute__((__vector_size__(4 * sizeof(short)))) short;
using f16_4vec    = __attribute__((__vector_size__(4 * sizeof(_Float16)))) _Float16;
using int32_8vec  = __attribute__((__vector_size__(8 * sizeof(int)))) int;

/** MFMA sanity: all lanes finite (FP paths) and at least one nonzero accumulator lane. */
__device__ inline bool
mfma_acc_f32_nonempty_finite(const f32_16vec& r)
{
    bool any_nonzero = false;
    for(int i = 0; i < 16; i++)
    {
        if(!isfinite(r[i]))
            return false;
        if(r[i] != 0.0f)
            any_nonzero = true;
    }
    return any_nonzero;
}

__device__ inline bool
mfma_acc_f64_nonempty_finite(const f64_4vec& r)
{
    bool any_nonzero = false;
    for(int i = 0; i < 4; i++)
    {
        if(!isfinite(r[i]))
            return false;
        if(r[i] != 0.0)
            any_nonzero = true;
    }
    return any_nonzero;
}

__device__ inline bool
mfma_acc_i32_nonempty(const int32_16vec& r)
{
    bool any_nonzero = false;
    for(int i = 0; i < 16; i++)
    {
        if(r[i] != 0)
            any_nonzero = true;
    }
    return any_nonzero;
}

/** Pack operand bits for MFMA FP8 builtins that take IEEE binary64 slots at the API. */
__device__ inline double
mfma_fp8_packed_operand_u64(uint64_t bits)
{
    double d;
    memcpy(&d, &bits, sizeof(bits));
    return d;
}

// mfma_mode: 0=both (ASM+builtin), 1=asm only, 2=builtin only
__device__ void
test_mfma_operations([[maybe_unused]] TestResults* results, [[maybe_unused]] int lane_id,
                     [[maybe_unused]] int mfma_mode)
{
#if defined(__gfx90a__)
    // ========================================================================
    // MI250 (gfx90a/CDNA2) MFMA Tests
    // mfma_mode: 0=both, 1=asm only, 2=builtin only
    // Note: MFMA inline ASM requires complex accumulator register constraints.
    // We use builtins as the primary test method since they properly handle
    // the register allocation. ASM option is preserved for compatibility.
    // ========================================================================
    bool f32_passed  = true;
    bool f64_passed  = true;
    bool f16_passed  = true;
    bool bf16_passed = true;
    bool i8_passed   = true;

    // ---- FP32 MFMA: v_mfma_f32_32x32x2f32 (4096 ops) ----
    {
        if(mfma_mode != 1)
        {
            f32_16vec result = { 0 };
            result = __builtin_amdgcn_mfma_f32_32x32x2f32(1.0f, 1.0f, result, 0, 0, 0);
            f32_passed = mfma_acc_f32_nonempty_finite(result);
        }
        else
        {
            f32_passed = true;
        }
    }

    // ---- FP64 MFMA: v_mfma_f64_16x16x4f64 (2048 ops) ----
    {
        if(mfma_mode != 1)
        {
            f64_4vec result = { 0 };
            result = __builtin_amdgcn_mfma_f64_16x16x4f64(1.0, 1.0, result, 0, 0, 0);
            f64_passed = mfma_acc_f64_nonempty_finite(result);
        }
        else
        {
            f64_passed = true;
        }
    }

    // ---- FP16 MFMA: v_mfma_f32_32x32x8f16 (16384 ops) ----
    {
        if(mfma_mode != 1)
        {
            const _Float16 one_h = (_Float16)1.0f;
            f16_4vec          a      = { one_h, one_h, one_h, one_h };
            f32_16vec         result = { 0 };
            result                   = __builtin_amdgcn_mfma_f32_32x32x8f16(a, a, result, 0, 0, 0);
            f16_passed               = mfma_acc_f32_nonempty_finite(result);
        }
        else
        {
            f16_passed = true;
        }
    }

    // ---- BF16 MFMA: v_mfma_f32_32x32x4bf16 (8192 ops) ----
    {
        if(mfma_mode != 1)
        {
            const short       bf16_one = (short) (0x3f800000u >> 16);
            bf16_2vec         a        = { bf16_one, bf16_one };
            f32_16vec         result   = { 0 };
            result                     = __builtin_amdgcn_mfma_f32_32x32x4bf16(a, a, result, 0, 0, 0);
            bf16_passed                = mfma_acc_f32_nonempty_finite(result);
        }
        else
        {
            bf16_passed = true;
        }
    }

    // ---- INT8 MFMA: v_mfma_i32_32x32x8i8 (16384 ops) ----
    {
        if(mfma_mode != 1)
        {
            const int     i8_pack = 0x01010101;
            int32_16vec   result  = { 0 };
            result                = __builtin_amdgcn_mfma_i32_32x32x8i8(i8_pack, i8_pack, result,
                                                           0, 0, 0);
            i8_passed             = mfma_acc_i32_nonempty(result);
        }
        else
        {
            i8_passed = true;
        }
    }

    if(lane_id == 0)
    {
        atomicAdd(&results->mfma_f32_passed, f32_passed ? 1 : 0);
        atomicAdd(&results->mfma_f64_passed, f64_passed ? 1 : 0);
        atomicAdd(&results->mfma_f16_passed, f16_passed ? 1 : 0);
        atomicAdd(&results->mfma_bf16_passed, bf16_passed ? 1 : 0);
        atomicAdd(&results->mfma_i8_passed, i8_passed ? 1 : 0);
        atomicMin(&results->mfma_f8_passed, -1);
    }

#elif defined(__gfx942__) || defined(__gfx940__) || defined(__gfx941__)
    // ========================================================================
    // MI300 (gfx942/CDNA3) MFMA Tests - Enhanced with FP8 support
    // mfma_mode: 0=both, 1=asm only, 2=builtin only
    // ========================================================================
    bool f32_passed  = true;
    bool f64_passed  = true;
    bool f16_passed  = true;
    bool bf16_passed = true;
    bool i8_passed   = true;
    bool f8_passed   = true;

    // ---- FP32 MFMA: v_mfma_f32_32x32x2f32 (4096 ops) ----
    {
        if(mfma_mode != 1)
        {
            f32_16vec result = { 0 };
            result = __builtin_amdgcn_mfma_f32_32x32x2f32(1.0f, 1.0f, result, 0, 0, 0);
            f32_passed = mfma_acc_f32_nonempty_finite(result);
        }
        else
        {
            f32_passed = true;
        }
    }

    // ---- FP64 MFMA: v_mfma_f64_16x16x4f64 (2048 ops) ----
    {
        if(mfma_mode != 1)
        {
            f64_4vec result = { 0 };
            result = __builtin_amdgcn_mfma_f64_16x16x4f64(1.0, 1.0, result, 0, 0, 0);
            f64_passed = mfma_acc_f64_nonempty_finite(result);
        }
        else
        {
            f64_passed = true;
        }
    }

    // ---- FP16 MFMA: v_mfma_f32_32x32x8f16 (16384 ops) ----
    {
        if(mfma_mode != 1)
        {
            const _Float16 one_h = (_Float16)1.0f;
            f16_4vec          a      = { one_h, one_h, one_h, one_h };
            f32_16vec         result = { 0 };
            result                   = __builtin_amdgcn_mfma_f32_32x32x8f16(a, a, result, 0, 0, 0);
            f16_passed               = mfma_acc_f32_nonempty_finite(result);
        }
        else
        {
            f16_passed = true;
        }
    }

    // ---- BF16 MFMA: v_mfma_f32_32x32x8bf16_1k (16384 ops) - MI300 enhanced ----
    {
        if(mfma_mode != 1)
        {
            const short bf16_one       = (short) (0x3f800000u >> 16);
            bf16_4vec   a              = { bf16_one, bf16_one, bf16_one, bf16_one };
            f32_16vec   result         = { 0 };
            result                     = __builtin_amdgcn_mfma_f32_32x32x8bf16_1k(a, a, result, 0, 0, 0);
            bf16_passed                = mfma_acc_f32_nonempty_finite(result);
        }
        else
        {
            bf16_passed = true;
        }
    }

    // ---- INT8 MFMA: v_mfma_i32_32x32x16_i8 (32768 ops) - MI300 doubled K ----
    {
        if(mfma_mode != 1)
        {
            const long long i8_pack = 0x0101010101010101LL;
            int32_16vec     result  = { 0 };
            result =
                __builtin_amdgcn_mfma_i32_32x32x16_i8(i8_pack, i8_pack, result, 0, 0, 0);
            i8_passed = mfma_acc_i32_nonempty(result);
        }
        else
        {
            i8_passed = true;
        }
    }

    // ---- FP8 MFMA: v_mfma_f32_32x32x16_fp8_fp8 (32768 ops) - MI300 new ----
    {
        if(mfma_mode != 1)
        {
            const uint64_t fp8_bits =
                (uint64_t) (unsigned) (lane_id + 1)
                | ((uint64_t) (unsigned) (lane_id + 1) << 32);
            const double   fp8_a = mfma_fp8_packed_operand_u64(fp8_bits);
            const double   fp8_b = mfma_fp8_packed_operand_u64(fp8_bits);
            f32_16vec      result = { 0 };
            result =
                __builtin_amdgcn_mfma_f32_32x32x16_fp8_fp8(fp8_a, fp8_b, result, 0, 0, 0);
            f8_passed = mfma_acc_f32_nonempty_finite(result);
        }
        else
        {
            f8_passed = true;
        }
    }

    if(lane_id == 0)
    {
        atomicAdd(&results->mfma_f32_passed, f32_passed ? 1 : 0);
        atomicAdd(&results->mfma_f64_passed, f64_passed ? 1 : 0);
        atomicAdd(&results->mfma_f16_passed, f16_passed ? 1 : 0);
        atomicAdd(&results->mfma_bf16_passed, bf16_passed ? 1 : 0);
        atomicAdd(&results->mfma_i8_passed, i8_passed ? 1 : 0);
        atomicAdd(&results->mfma_f8_passed, f8_passed ? 1 : 0);
    }

#elif defined(__gfx950__)
    // ========================================================================
    // MI350 (gfx950/CDNA4) MFMA Tests - Adds scaled F8F6F4
    // mfma_mode: 0=both, 1=asm only, 2=builtin only
    // ========================================================================
    bool f32_passed  = true;
    bool f64_passed  = true;
    bool f16_passed  = true;
    bool bf16_passed = true;
    bool i8_passed   = true;
    bool f8_passed   = true;

    // ---- FP32 MFMA: v_mfma_f32_32x32x2f32 ----
    {
        if(mfma_mode != 1)
        {
            f32_16vec result = { 0 };
            result = __builtin_amdgcn_mfma_f32_32x32x2f32(1.0f, 1.0f, result, 0, 0, 0);
            f32_passed = mfma_acc_f32_nonempty_finite(result);
        }
        else
        {
            f32_passed = true;
        }
    }

    // ---- FP64 MFMA: v_mfma_f64_16x16x4f64 ----
    {
        if(mfma_mode != 1)
        {
            f64_4vec result = { 0 };
            result = __builtin_amdgcn_mfma_f64_16x16x4f64(1.0, 1.0, result, 0, 0, 0);
            f64_passed = mfma_acc_f64_nonempty_finite(result);
        }
        else
        {
            f64_passed = true;
        }
    }

    // ---- FP16 MFMA: v_mfma_f32_32x32x8f16 ----
    {
        if(mfma_mode != 1)
        {
            const _Float16 one_h = (_Float16)1.0f;
            f16_4vec          a      = { one_h, one_h, one_h, one_h };
            f32_16vec         result = { 0 };
            result                   = __builtin_amdgcn_mfma_f32_32x32x8f16(a, a, result, 0, 0, 0);
            f16_passed               = mfma_acc_f32_nonempty_finite(result);
        }
        else
        {
            f16_passed = true;
        }
    }

    // ---- BF16 MFMA: v_mfma_f32_32x32x8bf16_1k ----
    {
        if(mfma_mode != 1)
        {
            const short bf16_one       = (short) (0x3f800000u >> 16);
            bf16_4vec   a              = { bf16_one, bf16_one, bf16_one, bf16_one };
            f32_16vec   result         = { 0 };
            result                     = __builtin_amdgcn_mfma_f32_32x32x8bf16_1k(a, a, result, 0, 0, 0);
            bf16_passed                = mfma_acc_f32_nonempty_finite(result);
        }
        else
        {
            bf16_passed = true;
        }
    }

    // ---- INT8 MFMA: v_mfma_i32_32x32x16_i8 ----
    {
        if(mfma_mode != 1)
        {
            const long long i8_pack = 0x0101010101010101LL;
            int32_16vec     result  = { 0 };
            result =
                __builtin_amdgcn_mfma_i32_32x32x16_i8(i8_pack, i8_pack, result, 0, 0, 0);
            i8_passed = mfma_acc_i32_nonempty(result);
        }
        else
        {
            i8_passed = true;
        }
    }

    // ---- Scaled F8F6F4 MFMA: v_mfma_scale_f32_32x32x64_f8f6f4 (131072 ops) ----
    {
        if(mfma_mode != 1)
        {
            const int  pkt       = 0x01010101 * (lane_id + 1);
            int32_8vec a         = { pkt, pkt, pkt, pkt, pkt, pkt, pkt, pkt };
            f32_16vec  result    = { 0 };
            result               = __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(a, a, result,
                                                                   0, 0, 0, 0, 0, 0);
            f8_passed            = mfma_acc_f32_nonempty_finite(result);
        }
        else
        {
            f8_passed = true;
        }
    }

    if(lane_id == 0)
    {
        atomicAdd(&results->mfma_f32_passed, f32_passed ? 1 : 0);
        atomicAdd(&results->mfma_f64_passed, f64_passed ? 1 : 0);
        atomicAdd(&results->mfma_f16_passed, f16_passed ? 1 : 0);
        atomicAdd(&results->mfma_bf16_passed, bf16_passed ? 1 : 0);
        atomicAdd(&results->mfma_i8_passed, i8_passed ? 1 : 0);
        atomicAdd(&results->mfma_f8_passed, f8_passed ? 1 : 0);
    }

#else
    // MFMA not supported on RDNA4 consumer GPUs - mark as bypassed
    if(lane_id == 0)
    {
        atomicMin(&results->mfma_f32_passed, -1);
        atomicMin(&results->mfma_f64_passed, -1);
        atomicMin(&results->mfma_f16_passed, -1);
        atomicMin(&results->mfma_bf16_passed, -1);
        atomicMin(&results->mfma_i8_passed, -1);
        atomicMin(&results->mfma_f8_passed, -1);
    }
#endif
}

// ============================================================================
// Section 10.5: Dual-Issue VALU Instructions (gfx942+)
// ============================================================================
//
// DUAL-ISSUE VALU FEATURE (gfx942+):
// ===================================
// AMD GPUs support dual-issue VALU in different forms:
//
// TWO TYPES OF DUAL-ISSUE:
// ------------------------
// 1. SAME-WAVE CO-ISSUE (VOPD/VOPD3) - MI300 (gfx942), MI350 (gfx950):
//    - Single instruction encodes two VALU operations (X and Y opcodes)
//    - X opcode executes on Core MACC, Y opcode executes on Side MACC
//    - Example: VOPD encoding packs v_add_f32 + v_mul_f32 in one instruction
//    - VOPD3 extends this with 3 source operands per instruction
//    - Consumer RDNA (gfx115x / gfx120x) may expose related VOPD-style dual-op patterns
//      where the ISA and compiler support them; this sample does not claim perf verification.
//
// 2. CROSS-WAVE CO-ISSUE - MI350 (gfx950) ONLY:
//    - Two waves each contribute one VALU instruction
//    - Instructions execute simultaneously on different MACC units
//    - One instruction uses DP-MACC (Arbiter-B), other uses Core/Side-MACC (Arbiter-A)
//    - Requires dual VALU arbiters (new hardware in gfx950+)
//
// DUAL-ISSUE REQUIREMENTS:
// ------------------------
// For cross-wave co-issue to occur, BOTH instructions must satisfy:
//   - Both waves must be Wave32 (NOT Wave64)
//   - Both must be single-cycle VALU instructions
//   - One uses DP-MACC (FP64/DGEMM), other uses Core/Side MACC
//   - Neither instruction uses DPP (Data-Parallel Primitives)
//   - Non-DPFP instruction is NOT VOPD or VOPD3
//   - Neither is a Transcendental (sin, cos, exp, log, etc.) or WMMA instruction
//   - Source operands must not conflict (no register bank conflicts)
//
// INSTRUCTIONS THAT CAN CO-EXECUTE WITH DP-MACC (Arbiter-B):
// ----------------------------------------------------------
//   - Single-cycle FP32/FP16/BF16 operations (v_add_f32, v_mul_f32, v_fma_f32, etc.)
//   - Single-cycle integer operations (v_add_i32, v_mul_lo_u32, etc.)
//   - Comparison operations
//   - Bitwise operations
//
// INSTRUCTIONS THAT BLOCK DUAL-ISSUE:
// -----------------------------------
//   - VOPD/VOPD3 instructions (already dual-op)
//   - DPP instructions (data-parallel primitives)
//   - DGEMM/SGEMM (matrix multiply)
//   - XDL/WMMA instructions
//   - TMACC instructions
//   - Multi-cycle instructions (V_SWAP*, V_MQSAD*, V_MUL_U64, etc.)
//   - Transcendental operations (v_sin_f32, v_cos_f32, v_exp_f32, etc.)
//   - PIPEFLUSH operations
//
// HOW TO MAXIMIZE DUAL-ISSUE IN YOUR CODE:
// ----------------------------------------
// 1. Mix FP64 operations with FP32/FP16 operations
//    Example: While one wave does FP64 accumulation, another does FP32 normalization
//
// 2. Avoid back-to-back transcendentals - they block co-execution
//    Bad:  sin(x); cos(x);   // Neither can dual-issue
//    Good: sin(x); a + b;    // a+b might co-issue with sin from another wave
//
// 3. Use Wave32 dispatch for maximum dual-issue opportunity
//
// 4. Minimize register conflicts between concurrent operations
//
// DUAL-ISSUE PERFORMANCE TEST:
// ----------------------------
// The following test demonstrates instruction patterns that enable dual-issue.
// On MI350, mixing FP64 and FP32 ops from different waves can achieve 2x IPC.
//
// NOTE: Actual dual-issue happens at hardware scheduling level between waves.
// This test provides instruction patterns that are dual-issue eligible.
//

__device__ void
test_dual_issue_patterns(TestResults* results, int lane_id, int tid)
{
    (void) lane_id;
    // This function demonstrates dual-issue eligible instruction patterns
    // for MI350 (gfx950). The patterns are designed to maximize throughput.

#if defined(__gfx950__)
    // ========================================================================
    // PATTERN 1: FP64 + FP32 Co-execution
    // ========================================================================
    // FP64 ops use DP-MACC (Arbiter-B), FP32 ops use Core MACC (Arbiter-A)
    // These can potentially co-execute with FP64 ops from another wave

    double dp_accum = 1.0;
    float  sp_accum = 1.0f;

// Mix of FP64 and FP32 operations - eligible for cross-wave dual-issue
// When wave A executes FP64 FMA, wave B's FP32 add can co-execute
#    pragma unroll 4
    for(int i = 0; i < 4; i++)
    {
        dp_accum = dp_accum * 1.00001 + 0.00001;  // DP-MACC (Arbiter-B)
        sp_accum = sp_accum * 1.001f + 0.001f;    // Core MACC (Arbiter-A)
    }

    // ========================================================================
    // PATTERN 2: Integer + FP32 Co-execution
    // ========================================================================
    // Integer VALU ops and FP32 ops can potentially co-execute

    int   int_accum = lane_id;
    float fp_accum  = (float) lane_id;

#    pragma unroll 4
    for(int i = 0; i < 4; i++)
    {
        int_accum = int_accum * 3 + 1;       // Integer MACC
        fp_accum  = fp_accum * 1.5f + 0.5f;  // FP32 MACC
    }

    // ========================================================================
    // PATTERN 3: Avoid Dual-Issue Blockers
    // ========================================================================
    // Transcendentals CANNOT co-execute - avoid clustering them

    // BAD PATTERN (blocks dual-issue):
    // float s = sinf(x);
    // float c = cosf(x);  // Both use Trans pipe, no co-execution

    // BETTER PATTERN (allows some dual-issue):
    float x = (float) lane_id * 0.1f;
    float s = sinf(x);          // Trans pipe - no co-execute
    float y = x * 2.0f + 1.0f;  // Core MACC - might co-execute with other wave
    float c = cosf(x);          // Trans pipe - no co-execute
    float z = y * 0.5f - 0.5f;  // Core MACC - might co-execute with other wave

    // Use results to prevent optimization
    volatile float sink =
        (float) dp_accum + sp_accum + (float) int_accum + fp_accum + s + c + y + z;
    (void) sink;

#elif defined(__gfx942__)
    // ========================================================================
    // MI300 (gfx942): Same-Wave Co-Issue via VOPD/VOPD3
    // ========================================================================
    // MI300 supports VOPD encoding which executes 2 operations from the SAME
    // wave in parallel. X instruction uses Core MACC, Y instruction uses Side MACC.
    //
    // VOPD Operations (executed as dual ops in one instruction):
    // - V_DUAL_FMAC_F32, V_DUAL_MUL_F32, V_DUAL_ADD_F32, V_DUAL_SUB_F32
    // - V_DUAL_MOV_B32, V_DUAL_CNDMASK_B32
    // - V_DUAL_MAX_NUM_F32, V_DUAL_MIN_NUM_F32
    // - V_DUAL_ADD_NC_U32, V_DUAL_LSHLREV_B32, etc.
    //
    // VOPD3 extends this with 3 source operands per instruction.

    // Pattern: FP32 operations that compiler can potentially encode as VOPD
    float a = (float) lane_id * 0.1f;
    float b = (float) lane_id * 0.2f;
    float c = 1.0f;
    float d = 2.0f;

// These paired operations may be encoded as VOPD by the compiler
#    pragma unroll 4
    for(int i = 0; i < 4; i++)
    {
        // Pair 1: Two independent FMAs - potential VOPD
        c = c * a + 0.001f;  // V_DUAL_FMAC_F32 (X op - Core MACC)
        d = d * b + 0.002f;  // V_DUAL_FMAC_F32 (Y op - Side MACC)
    }

    // Pattern: Mixed integer and FP32 - potential VOPD
    int   e = lane_id;
    float f = (float) lane_id;

#    pragma unroll 4
    for(int i = 0; i < 4; i++)
    {
        e = e + 1;     // V_DUAL_ADD_NC_U32 (X op)
        f = f + 1.0f;  // V_DUAL_ADD_F32 (Y op)
    }

    // Use results to prevent optimization
    volatile float sink = c + d + f + (float) e;
    (void) sink;

#elif defined(__gfx90a__)
    // MI250 (gfx90a/CDNA2): Limited dual-issue capability
    // Basic operations for compatibility
    float accum = 1.0f;
#    pragma unroll 4
    for(int i = 0; i < 4; i++)
    {
        accum = accum * 1.001f + 0.001f;
    }
    volatile float sink = accum;
    (void) sink;

#else
    // Other architectures - basic fallback
    volatile float sink = 1.0f;
    (void) sink;
#endif

    // Note: Dual-issue is a hardware scheduling decision.
    // This test provides dual-issue ELIGIBLE patterns.
    // Actual dual-issue depends on wave scheduling and register availability.

    // Dual-issue is a scheduler/performance property: we emit eligible instruction
    // patterns but do not verify co-issue in hardware (would need perf counters).
    if(tid == 0)
    {
        results->dual_issue_passed = -1;
    }
}

// ============================================================================
// Section 11: Lane/Thread Info Operations
// ============================================================================

__device__ void
test_lane_thread_info(TestResults* results, int lane_id)
{
    // Get lane ID using builtin
    unsigned int hw_lane_id;
    if(warpSize == 32)
    {
        hw_lane_id = __builtin_amdgcn_mbcnt_lo(-1, 0);
    }
    else
    {
        hw_lane_id = __builtin_amdgcn_mbcnt_hi(-1, __builtin_amdgcn_mbcnt_lo(-1, 0));
    }

    bool passed = (hw_lane_id == (unsigned) lane_id);

    // Test wavefront size
    int ws = __builtin_amdgcn_wavefrontsize();
    passed = passed && (ws == 32 || ws == 64);

    if(lane_id == 0)
    {
        atomicAdd(&results->lane_id_passed, passed ? 1 : 0);
        atomicAdd(&results->wavefront_passed, 1);
    }
}

// ============================================================================
// Section 12: Packed FP16/BF16 Operations
// ============================================================================

__device__ void
test_packed_operations(TestResults* results, int lane_id)
{
    bool f16_passed  = true;
    bool bf16_passed = true;

    // ========================================================================
    // TEST 1: Packed FP16 Operations
    // ========================================================================
    // Create packed FP16 values (2 x FP16 in 32 bits)
    __half2 h2_a = __float2half2_rn(1.5f);
    __half2 h2_b = __float2half2_rn(2.0f);

    // Packed add
    __half2 h2_sum = __hadd2(h2_a, h2_b);

    // Extract and verify
    float low  = __low2float(h2_sum);
    float high = __high2float(h2_sum);
    f16_passed = f16_passed && approx_equal_f32(low, 3.5f, 0.01f);
    f16_passed = f16_passed && approx_equal_f32(high, 3.5f, 0.01f);

    // Packed multiply
    __half2 h2_prod = __hmul2(h2_a, h2_b);
    low             = __low2float(h2_prod);
    f16_passed      = f16_passed && approx_equal_f32(low, 3.0f, 0.01f);

    // ========================================================================
    // TEST 2: Packed BF16 Operations
    // ========================================================================
    // Create packed BF16 values (2 x BF16 in 32 bits)
    // BF16 has same exponent range as FP32 but reduced mantissa precision

    // Convert floats to BF16
    __hip_bfloat16 bf16_a = __float2bfloat16(1.5f);
    __hip_bfloat16 bf16_b = __float2bfloat16(2.5f);

    // Create packed BF16 (using __hip_bfloat162 if available)
#if defined(__HIP_PLATFORM_AMD__)
    // Pack two BF16 values
    __hip_bfloat162 bf16_2_a;
    bf16_2_a.x = bf16_a;
    bf16_2_a.y = bf16_a;

    __hip_bfloat162 bf16_2_b;
    bf16_2_b.x = bf16_b;
    bf16_2_b.y = bf16_b;

    // Packed BF16 add (if available)
    __hip_bfloat162 bf16_2_sum = __hadd2(bf16_2_a, bf16_2_b);

    // Convert back to float and verify
    float bf16_low  = __bfloat162float(bf16_2_sum.x);
    float bf16_high = __bfloat162float(bf16_2_sum.y);

    // Expected: 1.5 + 2.5 = 4.0 (with BF16 precision tolerance)
    bf16_passed = bf16_passed && approx_equal_f32(bf16_low, 4.0f, 0.02f);
    bf16_passed = bf16_passed && approx_equal_f32(bf16_high, 4.0f, 0.02f);

    // Packed BF16 multiply
    __hip_bfloat162 bf16_2_prod = __hmul2(bf16_2_a, bf16_2_b);
    bf16_low                    = __bfloat162float(bf16_2_prod.x);

    // Expected: 1.5 * 2.5 = 3.75 (with BF16 precision tolerance)
    bf16_passed = bf16_passed && approx_equal_f32(bf16_low, 3.75f, 0.02f);
#else
    // Fallback: Manual BF16 packed operations
    // Convert to BF16, do arithmetic, convert back
    float a_val = __bfloat162float(bf16_a);
    float b_val = __bfloat162float(bf16_b);

    float sum_val  = a_val + b_val;
    float prod_val = a_val * b_val;

    bf16_passed = bf16_passed && approx_equal_f32(sum_val, 4.0f, 0.02f);
    bf16_passed = bf16_passed && approx_equal_f32(prod_val, 3.75f, 0.02f);
#endif

    if(lane_id == 0)
    {
        atomicAdd(&results->pk_f16_passed, f16_passed ? 1 : 0);
        atomicAdd(&results->pk_bf16_passed, bf16_passed ? 1 : 0);
        atomicAdd(&results->pk_add_passed, (f16_passed && bf16_passed) ? 1 : 0);
    }
}

// ============================================================================
// Section 13: Async LDS Operations (MI350/RDNA4 specific - global_load/store_async)
// ============================================================================
//
// Async LDS Support by Architecture:
// - MI250 (gfx90a): NO async LDS support
// - MI300 (gfx942): LIMITED async LDS support (may not have all builtins)
// - MI350 (gfx950): FULL async LDS support (global_load/store_async_to/from_lds)
// - RDNA4: FULL async LDS support + DS atomic async barriers
//
// GFX12+ / MI350-SPECIFIC DS ATOMIC BARRIER BUILTINS:
// --------------------------------------------
// These are used for synchronization with async LDS operations:
// - __builtin_amdgcn_ds_atomic_async_barrier_arrive_b64: Signal barrier arrival
// - __builtin_amdgcn_ds_atomic_barrier_arrive_rtn_b64: Signal barrier with return value
//
// USAGE PATTERN:
// 1. Issue async LDS load/store operations
// 2. Signal barrier arrival using ds_atomic_async_barrier_arrive
// 3. Wait for barrier completion (all threads arrived)
// 4. Use s_wait_asynccnt + s_wait_dscnt to ensure completion
//

// Shared memory for async LDS test
__shared__ int async_lds_buffer[64];
#if 0  // defined() placeholder for future architecture
__shared__ long async_barrier_counter;  // Placeholder type for DS atomic barrier experiments
#else
__shared__ long long async_barrier_counter;
#endif

__device__ void
test_async_lds_operations([[maybe_unused]] float* global_src, float* global_dst,
                          TestResults* results, int lane_id, int tid)
{
    bool passed = true;

#if __has_builtin(__builtin_amdgcn_global_load_async_to_lds_b32)
    // MI350 (gfx950) / RDNA4: Full async LDS load/store support
    // global_load_async_to_lds_b8/b32/b64/b128: Load from global to LDS asynchronously
    // global_store_async_from_lds_b8/b32/b64/b128: Store from LDS to global
    // asynchronously

    // Initialize global source with test pattern
    if(tid < 64)
    {
        // First, copy test data to global memory
        global_src[tid] = (float) (tid * 10 + 1);
    }

    // Initialize async barrier counter per block for barrier experiments
    if(threadIdx.x == 0)
    {
        async_barrier_counter = 0;
    }
    __syncthreads();

    // Perform async load from global to LDS
    // Only first thread in each block issues the async load for this test
    if(threadIdx.x == 0)
    {
        // Cast pointers to appropriate address spaces
        // address_space(1) = global, address_space(3) = local/LDS
        typedef int __attribute__((ext_vector_type(4))) vint4;

        // Load 16 bytes (4 ints) asynchronously from global to LDS
        __builtin_amdgcn_global_load_async_to_lds_b128(
            ((__attribute__((address_space(1))) vint4*) global_src),
            ((__attribute__((address_space(3))) vint4*) async_lds_buffer), 0 /* offset */,
            0 /* cache policy */);
    }

#    if 0  // defined() placeholder for future architecture
    // GFX12+/MI350: Use DS atomic async barrier for synchronization
    // This signals that this thread has reached the barrier point
    // The barrier counter is atomically incremented
    if (__has_builtin(__builtin_amdgcn_ds_atomic_async_barrier_arrive_b64)) {
        __builtin_amdgcn_ds_atomic_async_barrier_arrive_b64(
            ((__attribute__((address_space(3))) long*)&async_barrier_counter));
    }

    // Wait for async operations to complete (s_wait_asynccnt on targets that expose it)
    __builtin_amdgcn_s_wait_asynccnt(0);  // Wait for async copy operations
#    else
    // Other architectures: Use unified s_waitcnt
    WAITCNT_ALL();  // Wait for all memory operations
#    endif
    __syncthreads();

    // Verify the data was loaded correctly to LDS
    if(tid < 4)
    {
        int expected = (int) (tid * 10 + 1);
        // Note: We loaded floats but stored as ints, reinterpret
        float* lds_as_float = (float*) async_lds_buffer;
        passed = passed && approx_equal_f32(lds_as_float[tid], (float) expected, 0.1f);
    }
    __syncthreads();

    // Now test async store from LDS to global
    if(tid == 0)
    {
        // Modify LDS data
        async_lds_buffer[0] = 12345;
        async_lds_buffer[1] = 67890;
        async_lds_buffer[2] = 11111;
        async_lds_buffer[3] = 22222;
    }
    __syncthreads();

    if(tid == 0)
    {
        typedef int __attribute__((ext_vector_type(4))) vint4;

        // Store 16 bytes from LDS to global asynchronously
        __builtin_amdgcn_global_store_async_from_lds_b128(
            ((__attribute__((address_space(1))) vint4*) (global_dst)),
            ((__attribute__((address_space(3))) vint4*) async_lds_buffer), 0 /* offset */,
            0 /* cache policy */);
    }

#    if 0  // defined() placeholder for future architecture
    // : Use DS atomic async barrier for store synchronization
    if (__has_builtin(__builtin_amdgcn_ds_atomic_async_barrier_arrive_b64)) {
        __builtin_amdgcn_ds_atomic_async_barrier_arrive_b64(
            ((__attribute__((address_space(3))) long*)&async_barrier_counter));
    }

    // Wait for async store operations to complete
    __builtin_amdgcn_s_wait_asynccnt(0);  // Wait for async copy operations
#    else
    // Other architectures: Use unified s_waitcnt
    WAITCNT_ALL();
#    endif
    __syncthreads();

    // Verify the store (on subsequent kernel or by reading back)
    // For now, just check that no crash occurred
    passed = passed && true;

#else
    // Async LDS builtins not available - mark as bypassed
    (void) global_src;
    (void) global_dst;
    (void) tid;
    (void) passed;

    if(lane_id == 0)
    {
        atomicMin(&results->async_lds_load_passed, -1);
        atomicMin(&results->async_lds_store_passed, -1);
    }
    return;
#endif

    if(lane_id == 0)
    {
        atomicAdd(&results->async_lds_load_passed, passed ? 1 : 0);
        atomicAdd(&results->async_lds_store_passed, passed ? 1 : 0);
    }
}

// ============================================================================
// Section 15.5: WMMA (Wave Matrix Multiply Accumulate) Operations (RDNA4 / RDNA 3.5)
// ============================================================================
//
// WMMA INSTRUCTIONS - RDNA4 Matrix Cores:
// =======================================
// Reference: https://gpuopen.com/learn/using_matrix_core_amd_rdna4/
//
// RDNA4 (gfx12/) features 3rd-generation Matrix Cores with significantly
// improved GEMM performance compared to RDNA3:
//
// PERFORMANCE COMPARISON (FLOPS/clock/CU):
// ----------------------------------------
// | Format | RDNA2 (RX 6950) | RDNA3 (RX 7900) | RDNA4 (RX 9070) |
// |--------|-----------------|-----------------|-----------------|
// | FP16   | 256             | 512             | 1024            |
// | BF16   | N/A             | 512             | 1024            |
// | I8     | 512             | 512             | 2048            |
//
// WMMA INTRINSIC FORMAT:
// ----------------------
// GFX12 (gfx1200/gfx1201/gfx1150/gfx1151): __builtin_amdgcn_wmma_*_16x16x16_*_w32_gfx12
// (3-arg)
// Pre-gfx12 WMMA: __builtin_amdgcn_wmma_* (no _gfx12 suffix; different signatures per
// BuiltinsAMDGPU.td)
//
// Available for gfx12 (RDNA4 / RDNA 3.5): _w32_gfx12 suffix, 3 args (a, b, c)
// - __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12  (FP16 → FP32)
// - __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12 (BF16 → FP32)
// - __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12  (INT8 → INT32)
//
// Available for older WMMA entry points only (no suffix): 8-arg FP16/BF16, 7-arg INT8
// - __builtin_amdgcn_wmma_f32_16x16x32_f16, __builtin_amdgcn_wmma_f32_16x16x32_bf16
// - __builtin_amdgcn_wmma_i32_16x16x64_iu8
//
// MATRIX LAYOUT IN VGPR:
// ----------------------
// WMMA operates on 16x16 matrices distributed across a Wave32:
// - Each lane holds 8 elements (8 * 32 = 256 = 16 * 16 elements)
// - Matrix A is column-major (transposed)
// - Matrices B, C, D are row-major
// - laneWrapped = lane_id % 16 (column index within matrix; lane_id = threadIdx.x % warpSize)
// - laneGroup = lane_id / 16 (upper/lower half of matrix)
//
// WMMA USAGE PATTERN:
// -------------------
// 1. Load matrix A (column-major) into fragment
// 2. Load matrix B (row-major) into fragment
// 3. Execute WMMA intrinsic with all lanes in wavefront
// 4. Store result matrix D (row-major) from fragment
//
// IMPORTANT NOTES:
// ----------------
// - RDNA4 uses Wave32 ONLY for WMMA (no Wave64 WMMA)
// - VGPR layout changed vs RDNA3 - NOT backward compatible
// - All 32 lanes must participate in the WMMA operation
// - Matrix dimensions are fixed at 16x16
//

// Vector types for WMMA operations (RDNA4)
// Fragment types: 8 elements per lane for 16x16 matrix
#if defined(__gfx1200__) || defined(__gfx1201__) || ARCH_RDNA35_APU
using wmma_f16_frag  = __attribute__((__vector_size__(8 * sizeof(_Float16)))) _Float16;
using wmma_bf16_frag = __attribute__((__vector_size__(8 * sizeof(short)))) short;
using wmma_f32_frag  = __attribute__((__vector_size__(8 * sizeof(float)))) float;
using wmma_i32_frag  = __attribute__((__vector_size__(8 * sizeof(int)))) int;
using wmma_i8_frag   = __attribute__((
    __vector_size__(2 * sizeof(int)))) int;  // 2 packed i32 for 8 i8 elements

__device__ inline bool
wmma_frag_f32_nonempty_finite(const wmma_f32_frag& c)
{
    bool any_nonzero = false;
    for(int i = 0; i < 8; i++)
    {
        if(!isfinite(c[i]))
            return false;
        if(c[i] != 0.0f)
            any_nonzero = true;
    }
    return any_nonzero;
}

__device__ inline bool
wmma_frag_i32_nonempty(const wmma_i32_frag& c)
{
    bool any_nonzero = false;
    for(int i = 0; i < 8; i++)
    {
        if(c[i] != 0)
            any_nonzero = true;
    }
    return any_nonzero;
}
#endif

__device__ void
test_wmma_operations([[maybe_unused]] TestResults* results, [[maybe_unused]] int lane_id,
                     [[maybe_unused]] int tid)
{
#if 0  // defined() placeholder for future architecture
    // ========================================================================
    //  () WMMA Tests - use  builtins (no _gfx12 suffix)
    // BuiltinsAMDGPU.td: __builtin_amdgcn_wmma_f32_16x16x32_f16, _bf16, i32_16x16x64_iu8
    // ========================================================================
    bool f16_passed = true;
    bool bf16_passed = true;
    bool i8_passed = true;

    const int laneWrapped = tid % 16;
    const int laneGroup = tid / 16;
    //  FP16/BF16: 16 elements per lane for A and B; INT8: 8 int (32 i8) per lane
    using wmma_f16_frag_ = __attribute__((__vector_size__(16 * sizeof(_Float16)))) _Float16;
    using wmma_bf16_frag_ = __attribute__((__vector_size__(16 * sizeof(short)))) short;
    using wmma_i8_frag_ = __attribute__((__vector_size__(8 * sizeof(int)))) int;

    // ---- FP16: __builtin_amdgcn_wmma_f32_16x16x32_f16(transA, A_16, transB, B_16, scale, C_8, bool, bool) ----
    {
        wmma_f16_frag_ a_frag;
        wmma_f16_frag_ b_frag;
        wmma_f32_frag c_frag = {0};
        for (int ele = 0; ele < 16; ele++) {
            a_frag[ele] = (_Float16)((laneWrapped + (ele + laneGroup * 16) % 32) % 16 + 1) * 0.1f;
            b_frag[ele] = (_Float16)(((ele + laneGroup * 16) % 32 + laneWrapped) % 16 + 1) * 0.1f;
        }
        for (int i = 0; i < 8; i++) c_frag[i] = 0.0f;
#    if __has_builtin(__builtin_amdgcn_wmma_f32_16x16x32_f16)
        c_frag = __builtin_amdgcn_wmma_f32_16x16x32_f16(false, a_frag, false, b_frag, (short)0, c_frag, false, false);
#    endif
        volatile float sink = 0;
        for (int i = 0; i < 8; i++) sink += c_frag[i];
        f16_passed = (sink != 0.0f || tid >= 32);
        (void)sink;
    }

    // ---- BF16: __builtin_amdgcn_wmma_f32_16x16x32_bf16(transA, A_16, transB, B_16, scale, C_8, bool, bool) ----
    {
        wmma_bf16_frag_ a_frag;
        wmma_bf16_frag_ b_frag;
        wmma_f32_frag c_frag = {0};
        for (int ele = 0; ele < 16; ele++) {
            float va = (float)((laneWrapped + ele) % 16 + 1) * 0.1f;
            float vb = (float)((ele + laneWrapped) % 16 + 1) * 0.1f;
            a_frag[ele] = (short)(*(unsigned int*)&va >> 16);
            b_frag[ele] = (short)(*(unsigned int*)&vb >> 16);
        }
        for (int i = 0; i < 8; i++) c_frag[i] = 0.0f;
#    if __has_builtin(__builtin_amdgcn_wmma_f32_16x16x32_bf16)
        c_frag = __builtin_amdgcn_wmma_f32_16x16x32_bf16(false, a_frag, false, b_frag, (short)0, c_frag, false, false);
#    endif
        volatile float sink = 0;
        for (int i = 0; i < 8; i++) sink += c_frag[i];
        bf16_passed = (sink != 0.0f || tid >= 32);
        (void)sink;
    }

    // ---- INT8: __builtin_amdgcn_wmma_i32_16x16x64_iu8(negA, A_8i, negB, B_8i, C_8i, bool, bool) ----
    {
        wmma_i8_frag_ a_frag = {0};
        wmma_i8_frag_ b_frag = {0};
        wmma_i32_frag c_frag = {0};
        for (int ele = 0; ele < 8; ele++) {
            unsigned int pa = 0, pb = 0;
            for (int i = 0; i < 4; i++) {
                int ia = (laneWrapped + ele * 4 + i) % 16 + 1;
                int ib = (ele * 4 + i + laneWrapped) % 16 + 1;
                pa |= (ia & 0xFF) << (i * 8);
                pb |= (ib & 0xFF) << (i * 8);
            }
            a_frag[ele] = (int)pa;
            b_frag[ele] = (int)pb;
        }
#    if __has_builtin(__builtin_amdgcn_wmma_i32_16x16x64_iu8)
        c_frag = __builtin_amdgcn_wmma_i32_16x16x64_iu8(false, a_frag, false, b_frag, c_frag, false, false);
#    endif
        volatile int sink = 0;
        for (int i = 0; i < 8; i++) sink += c_frag[i];
        i8_passed = (sink != 0 || tid >= 32);
        (void)sink;
    }

    if (lane_id == 0) {
        atomicAdd(&results->wmma_f16_passed, f16_passed ? 1 : 0);
        atomicAdd(&results->wmma_bf16_passed, bf16_passed ? 1 : 0);
        atomicAdd(&results->wmma_i8_passed, i8_passed ? 1 : 0);
    }

#elif ARCH_RDNA35_APU
    // ========================================================================
    // RDNA 3.5 APU (gfx115x - Strix) WMMA Tests
    // Uses GFX11 WMMA builtins (wmma-256b-insts) with _Vector<16> inputs
    // These are DIFFERENT from GFX12 builtins which use _Vector<8> inputs
    // ========================================================================
    bool f16_passed  = true;
    bool bf16_passed = true;
    bool i8_passed   = true;

    const int laneWrapped = lane_id % 16;
    const int laneGroup   = lane_id / 16;

    // GFX11 WMMA uses 16-element input vectors (wmma-256b-insts)
    using wmma_f16_frag_gfx11 =
        __attribute__((__vector_size__(16 * sizeof(_Float16)))) _Float16;
    using wmma_bf16_frag_gfx11 =
        __attribute__((__vector_size__(16 * sizeof(short)))) short;
    using wmma_i8_frag_gfx11 = __attribute__((
        __vector_size__(4 * sizeof(int)))) int;  // 4 x int32 for 16 x int8 packed

    // ---- FP16 WMMA: __builtin_amdgcn_wmma_f32_16x16x16_f16_w32 ----
    {
        wmma_f16_frag_gfx11 a_frag;
        wmma_f16_frag_gfx11 b_frag;
        wmma_f32_frag       c_frag = { 0 };

        for(int ele = 0; ele < 16; ele++)
        {
            a_frag[ele] =
                (_Float16) ((laneWrapped + ele + laneGroup * 8) % 16 + 1) * 0.1f;
            b_frag[ele] =
                (_Float16) ((ele + laneGroup * 8 + laneWrapped) % 16 + 1) * 0.1f;
        }

        for(int i = 0; i < 8; i++)
            c_frag[i] = 0.0f;

#    if __has_builtin(__builtin_amdgcn_wmma_f32_16x16x16_f16_w32)
        c_frag = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32(a_frag, b_frag, c_frag);
#    endif

        f16_passed = wmma_frag_f32_nonempty_finite(c_frag);
    }

    // ---- BF16 WMMA: __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32 ----
    {
        wmma_bf16_frag_gfx11 a_frag;
        wmma_bf16_frag_gfx11 b_frag;
        wmma_f32_frag        c_frag = { 0 };

        for(int ele = 0; ele < 16; ele++)
        {
            float val   = (float) ((laneWrapped + ele) % 16 + 1) * 0.1f;
            a_frag[ele] = (short) (*(unsigned int*) &val >> 16);
            b_frag[ele] = (short) (*(unsigned int*) &val >> 16);
        }

        for(int i = 0; i < 8; i++)
            c_frag[i] = 0.0f;

#    if __has_builtin(__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32)
        c_frag = __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32(a_frag, b_frag, c_frag);
#    endif

        bf16_passed = wmma_frag_f32_nonempty_finite(c_frag);
    }

    // ---- INT8 WMMA: __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32 ----
    {
        wmma_i8_frag_gfx11 a_frag;  // 4 x int32 for 16 x int8 packed
        wmma_i8_frag_gfx11 b_frag;
        wmma_i32_frag      c_frag = { 0 };

        // Pack 4 INT8 values into each of 4 int32 elements (16 total)
        for(int ele = 0; ele < 4; ele++)
        {
            unsigned int packed_a = 0;
            unsigned int packed_b = 0;
            for(int i = 0; i < 4; i++)
            {
                int val_a = ((laneWrapped + ele * 4 + i) % 16 + 1);
                int val_b = ((ele * 4 + i + laneWrapped) % 16 + 1);
                packed_a |= (val_a & 0xFF) << (i * 8);
                packed_b |= (val_b & 0xFF) << (i * 8);
            }
            a_frag[ele] = (int) packed_a;
            b_frag[ele] = (int) packed_b;
        }

        for(int i = 0; i < 8; i++)
            c_frag[i] = 0;

#    if __has_builtin(__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32)
        c_frag = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32(false, a_frag, false, b_frag,
                                                            c_frag, false);
#    endif

        i8_passed = wmma_frag_i32_nonempty(c_frag);
    }

    if(lane_id == 0)
    {
        atomicAdd(&results->wmma_f16_passed, f16_passed ? 1 : 0);
        atomicAdd(&results->wmma_bf16_passed, bf16_passed ? 1 : 0);
        atomicAdd(&results->wmma_i8_passed, i8_passed ? 1 : 0);
    }

#elif defined(__gfx1200__) || defined(__gfx1201__)
    // ========================================================================
    // RDNA4 Consumer (gfx1200/gfx1201 - RX 9060 / RX 9070 families) WMMA Tests
    // Uses GFX12 WMMA builtins (wmma-128b-insts) with _Vector<8> inputs
    // ========================================================================
    bool f16_passed  = true;
    bool bf16_passed = true;
    bool i8_passed   = true;

    const int laneWrapped     = lane_id % 16;
    const int laneGroup       = lane_id / 16;
    const int WMMA_DATA_WIDTH = 8;

    // ---- FP16 WMMA: __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12 ----
    {
        wmma_f16_frag a_frag;
        wmma_f16_frag b_frag;
        wmma_f32_frag c_frag = { 0 };

        for(int ele = 0; ele < WMMA_DATA_WIDTH; ele++)
        {
            a_frag[ele] =
                (_Float16) ((laneWrapped + ele + laneGroup * WMMA_DATA_WIDTH) % 16 + 1) *
                0.1f;
            b_frag[ele] =
                (_Float16) ((ele + laneGroup * WMMA_DATA_WIDTH + laneWrapped) % 16 + 1) *
                0.1f;
        }

        for(int i = 0; i < 8; i++)
            c_frag[i] = 0.0f;

#    if __has_builtin(__builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12)
        c_frag = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12(a_frag, b_frag, c_frag);
#    endif

        f16_passed = wmma_frag_f32_nonempty_finite(c_frag);
    }

    // ---- BF16 WMMA: __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12 ----
    {
        wmma_bf16_frag a_frag;
        wmma_bf16_frag b_frag;
        wmma_f32_frag  c_frag = { 0 };

        for(int ele = 0; ele < WMMA_DATA_WIDTH; ele++)
        {
            float val   = (float) ((laneWrapped + ele) % 16 + 1) * 0.1f;
            a_frag[ele] = (short) (*(unsigned int*) &val >> 16);
            b_frag[ele] = (short) (*(unsigned int*) &val >> 16);
        }

        for(int i = 0; i < 8; i++)
            c_frag[i] = 0.0f;

#    if __has_builtin(__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12)
        c_frag =
            __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12(a_frag, b_frag, c_frag);
#    endif

        bf16_passed = wmma_frag_f32_nonempty_finite(c_frag);
    }

    // ---- INT8 WMMA: __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12 ----
    {
        wmma_i8_frag  a_frag;  // 2 x int32 for 8 x int8 packed
        wmma_i8_frag  b_frag;
        wmma_i32_frag c_frag = { 0 };

        // Pack 4 INT8 values into each of 2 int32 elements
        for(int ele = 0; ele < 2; ele++)
        {
            unsigned int packed_a = 0;
            unsigned int packed_b = 0;
            for(int i = 0; i < 4; i++)
            {
                int val_a = ((laneWrapped + ele * 4 + i) % 16 + 1);
                int val_b = ((ele * 4 + i + laneWrapped) % 16 + 1);
                packed_a |= (val_a & 0xFF) << (i * 8);
                packed_b |= (val_b & 0xFF) << (i * 8);
            }
            a_frag[ele] = (int) packed_a;
            b_frag[ele] = (int) packed_b;
        }

        for(int i = 0; i < 8; i++)
            c_frag[i] = 0;

#    if __has_builtin(__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12)
        c_frag = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12(false, a_frag, false,
                                                                  b_frag, c_frag, false);
#    endif

        i8_passed = wmma_frag_i32_nonempty(c_frag);
    }

    if(lane_id == 0)
    {
        atomicAdd(&results->wmma_f16_passed, f16_passed ? 1 : 0);
        atomicAdd(&results->wmma_bf16_passed, bf16_passed ? 1 : 0);
        atomicAdd(&results->wmma_i8_passed, i8_passed ? 1 : 0);
    }

#else
    // WMMA not available on CDNA architectures (use MFMA instead)
    // Mark as bypassed
    if(lane_id == 0)
    {
        atomicMin(&results->wmma_f16_passed, -1);
        atomicMin(&results->wmma_bf16_passed, -1);
        atomicMin(&results->wmma_i8_passed, -1);
    }
#endif
}

// ============================================================================
// Section 16: RDNA4 Consumer-Specific Features (RX 9060 / RX 9070 families / gfx1200/gfx1201)
// ============================================================================
//
// RDNA4 CONSUMER GPU FEATURES (RX 9070 XT, RX 9070, RX 9060 series):
// ==================================================================
//
// ARCHITECTURE OVERVIEW:
// ----------------------
// - GFX12 architecture (gfx1200/gfx1201)
// - WGP (Workgroup Processor) based - 2 CUs per WGP
// - Wave32 native, Wave64 emulated
// - FP8/BF8 OCP (Open Compute Platform) format support
//
// KEY DIFFERENCES FROM CDNA (MI series):
// --------------------------------------
// - No hardware FP64 atomics (limited FP64 in general)
// - Graphics-focused instruction set (more pixel/vertex ops)
// - Different memory hierarchy (infinity cache)
// - WGP instead of CU-centric design
//
// RDNA4 SPECIFIC INSTRUCTIONS:
// ----------------------------
// - FP8/BF8 conversions (OCP format)
// - Async LDS operations
//
// This test verifies RDNA4-specific features when running on consumer GPUs.

__device__ void
test_rdna4_features([[maybe_unused]] TestResults* results, [[maybe_unused]] int lane_id)
{
#if defined(__gfx1200__) || defined(__gfx1201__) || ARCH_RDNA35_APU
    bool passed = true;

    // Test 1: Wave32 operations (RDNA4 is Wave32 native)
    // Wave32 ballot mask test
    uint32_t ballot_mask = __ballot(1);
    passed               = passed && (ballot_mask != 0);

    // Test 2: Subgroup operations
    int lane_val     = lane_id;
    int first_active = __shfl(lane_val, 0);
    passed           = passed && (first_active == 0);

    // Test 3: FP16 packed operations (RDNA4 has optimized FP16)
    half2 h2_val    = __float2half2_rn(1.5f);
    float converted = __half2float(__low2half(h2_val));
    passed          = passed && (fabsf(converted - 1.5f) < 0.01f);

    // Test 4: WGP local barrier (implicit in __syncthreads on RDNA4)
    __syncthreads();

    if(lane_id == 0)
    {
        atomicAdd(&results->lane_id_passed, passed ? 1 : 0);
    }

#else
    // Not RDNA4 - skip
    (void) results;
    (void) lane_id;
#endif
}

// ============================================================================
// Main Mega Kernel
// ============================================================================

extern "C" __global__ void
mega_kernel(
    TestResults* results, float* global_float, double* global_double, int* global_int,
    float* input_buffer, float* output_buffer,
    float*             async_lds_src,  // Source buffer for async LDS load test
    float*             async_lds_dst,  // Destination buffer for async LDS store test
    int                buffer_size,
    int                mfma_mode,  // 0=both, 1=asm only, 2=builtin only
    hipTextureObject_t tex_obj,    // Texture for TEX load test (Strix); 0 if not used
    hipSurfaceObject_t surf_obj    // Surface for TEX store test (Strix); 0 if not used
)
{
    int tid     = threadIdx.x + blockIdx.x * blockDim.x;
    int lane_id = threadIdx.x % warpSize;

    // Use active flag instead of early return to avoid __syncthreads() deadlock
    // When buffer_size is not a multiple of blockDim.x, some threads in the last
    // block would return early while others hit barriers, causing deadlock
    bool active = (tid < buffer_size);

    // results (and other global buffers) are initialized on the host before launch
    // (main.cpp). Do not memset here: with grid_size > 1, only tid==0 would run it and
    // __syncthreads() is block-local, so other blocks could start updating results via
    // atomics before memset completes.

    // =========== Run all tests ===========

    // 1. Warp Operations
    test_warp_shuffle(results, lane_id);
    test_warp_ballot(results, lane_id);
    test_warp_permute(results, lane_id);
    test_warp_reduce(results, lane_id);

    // 2. FP8/BF8 Conversions
    test_fp8_conversions(results, lane_id);
    test_bf8_conversions(results, lane_id);

    // 3. FP16/BF16 Operations
    test_fp16_operations(results, lane_id);
    test_bf16_operations(results, lane_id);

    // 4. FP32/FP64 Arithmetic
    test_fp32_arithmetic(results, lane_id);
    test_fp64_arithmetic(results, lane_id);

    // 5. Integer Operations
    test_int32_operations(results, lane_id);
    test_int64_operations(results, lane_id);

    // 6. Atomic Operations
    test_atomic_add_f32(global_float, results, lane_id);
    __syncthreads();
    test_atomic_add_f64(global_double, results, lane_id);
    __syncthreads();
    test_atomic_min_max_f32(results, lane_id);
    __syncthreads();
    test_atomic_int(global_int, results, lane_id);
    __syncthreads();

    // 6.5. Comprehensive Atomic Operations (Inline ASM)
    test_atomic_operations_asm(global_int, global_float, global_double, results, lane_id,
                               tid);
    __syncthreads();

    // 7. LDS Operations
    test_lds_operations(results, lane_id, threadIdx.x);

    // 8. Transcendental Functions
    test_transcendental(results, lane_id);

    // 9. Memory Operations (guard with active flag to avoid out-of-bounds access)
    if(active)
    {
        test_global_memory(input_buffer, output_buffer, results, lane_id, tid);
    }

    // 9.5. VMEM (Vector Memory) Comprehensive Tests
    // Always entered so all threads hit __syncthreads inside; global_ptr uses tid only when
    // tid < buffer_size (see in_bounds in test_vmem_operations).
    test_vmem_operations(input_buffer, results, lane_id, tid, buffer_size, tex_obj,
                         surf_obj);
    __syncthreads();

    // 10. DOT Products
    test_dot_products(results, lane_id);

    // 10.3 MFMA (Matrix Fused Multiply-Add) Operations
    test_mfma_operations(results, lane_id, mfma_mode);

    // 10.5 Dual-Issue VALU Patterns (gfx942+)
    test_dual_issue_patterns(results, lane_id, tid);

    // 11. Lane/Thread Info
    test_lane_thread_info(results, lane_id);

    // 12. Packed Operations
    test_packed_operations(results, lane_id);

    // 13. Async LDS Operations (MI350/RDNA4)
    test_async_lds_operations(async_lds_src, async_lds_dst, results, lane_id, tid);
    __syncthreads();

    // 14. RDNA4 Consumer-Specific Features (RX 9060 / RX 9070 families)
    test_rdna4_features(results, lane_id);
    __syncthreads();

    // 15. WMMA (Wave Matrix Multiply Accumulate) Operations (RDNA4 / RDNA 3.5)
    test_wmma_operations(results, lane_id, tid);
    __syncthreads();
}