RDNA4 Instinct Emulation Library

Unlock the full potential of AMD RDNA4 consumer GPUs by emulating MI300 Instinct-class capabilities.

This library provides production-grade, high-performance compute primitives that bring data center-class features to consumer hardware like the Radeon RX 9070 XT.

---

🚀 Features

Feature	Description	Status
FP16 WMMA GEMM	16x16x16 tensor core matrix multiplication	✅ Working
FP8 WMMA GEMM	Native 8-bit floating point with E4M3 format	✅ Working
FP4 GEMM	Experimental 4-bit with on-the-fly dequantization	✅ Working
FP32 GEMM	Standard precision via VALU path	✅ Working
FP64 GEMM	Double precision via VALU path	✅ Working
GTT Spillover	Automatic VRAM → System RAM overflow	✅ Working
Batched GEMM	Multi-head attention support	✅ Working
RT-Sparse GEMM	Ray tracing accelerated sparse operations	✅ Working
Convolution	WMMA-accelerated 2D convolution	✅ Working

---

📦 What's Included

RDNA4_Instinct_Emulation_Library_v2.0/
├── bin/
│   ├── rdna4_hip_wrapper.dll    # Pre-compiled library (Windows x64)
│   └── rdna4_hip_wrapper.lib    # Import library for linking
├── src/
│   ├── rdna4_integration.hpp    # Master header - include this
│   ├── rdna4_wmma_gemm.hpp      # WMMA GEMM kernels
│   ├── rdna4_wmma_kernels.hip.cpp  # Kernel implementations
│   ├── rdna4_fp_types.hpp       # FP8/FP4 data type definitions
│   ├── rdna4_gtt_allocator.hpp  # Smart memory allocator
│   ├── rdna4_batched_gemm.hpp   # Batched operations for transformers
│   ├── rdna4_convolution.hpp    # WMMA convolution kernels
│   ├── rdna4_rt_sparse.hpp      # RT-accelerated sparse operations
│   ├── rdna4_hiprt_bvh.hpp      # BVH construction for sparsity
│   └── CMakeLists.txt           # Build configuration
├── tests/
│   ├── test_multi_precision.cpp # Full precision test suite
│   └── ...
├── docs/
│   └── *.md                     # Integration guides
└── README.md

---

🛠️ Prerequisites

Required Software

Software	Version	Download
Windows	10/11 (64-bit)	-
AMD ROCm SDK	7.1+	AMD Developer Portal
Visual Studio	2022	Download
CMake	3.21+	Download
Ninja	Latest	Download

Supported Hardware

GPU	Architecture	Tensor Cores
Radeon RX 9070 XT	gfx1201	✅ WMMA 16x16x16
Radeon RX 9070	gfx1201	✅ WMMA 16x16x16
Other RDNA4	gfx12xx	✅ Should work

Environment Setup

Ensure these environment variables are set:

$env:HIP_PATH = "C:\Program Files\AMD\ROCm\7.1"
$env:ROCM_PATH = "C:\Program Files\AMD\ROCm\7.1"
$env:PATH += ";C:\Program Files\AMD\ROCm\7.1\bin"

---

🔧 Building from Source

Quick Build

cd RDNA4_Instinct_Emulation_Library_v2.0\src

# Configure
cmake -B build -G Ninja -DROCM_ENABLE_RDNA4=ON

# Build
cmake --build build

# Output: build/rdna4_hip_wrapper.dll

Manual Build with hipcc

hipcc --offload-arch=gfx1201 -std=c++17 -shared ^
  rdna4_wmma_kernels.hip.cpp ^
  rdna4_wmma_exports_minimal.cpp ^
  -o rdna4_hip_wrapper.dll

---

📖 Usage

Quick Start (Using Pre-compiled DLL)

#include <hip/hip_runtime.h>

// Declare the exported functions
extern "C" {
    hipError_t rdna4_gemm_fp16(int M, int N, int K, float alpha,
                               const void* A, int lda,
                               const void* B, int ldb,
                               float beta, void* C, int ldc);

    hipError_t rdna4_gemm_fp8(int M, int N, int K, float alpha,
                              const void* A, int lda, float scale_a,
                              const void* B, int ldb, float scale_b,
                              float beta, void* C, int ldc);
}

int main() {
    // Allocate your matrices...
    __half *d_A, *d_B;
    float *d_C;

    // Call FP16 WMMA GEMM - uses RDNA4 tensor cores
    rdna4_gemm_fp16(1024, 1024, 1024, 1.0f,
                    d_A, 1024, d_B, 1024,
                    0.0f, d_C, 1024);

    return 0;
}

Using the Full Integration Header

#include "rdna4_integration.hpp"

using namespace rocm::rdna4;

int main() {
    // Initialize RDNA4 features
    if (rocm_init_rdna4()) {
        std::cout << "RDNA4 features enabled!" << std::endl;
    }

    // Use smart memory allocation with GTT spillover
    void* ptr = rocm_malloc_smart(1024 * 1024 * 1024); // 1GB

    // Check memory usage
    rocm_print_memory_usage();

    // Clean up
    rocm_free_smart(ptr);

    return 0;
}

---

🧪 Running Tests

# Build the test
hipcc --offload-arch=gfx1201 -std=c++17 ^
  -I src ^
  tests/test_multi_precision.cpp ^
  -L bin -l rdna4_hip_wrapper ^
  -o test_multi_precision.exe ^
  -Xlinker /SUBSYSTEM:CONSOLE

# Run (ensure DLL is in PATH)
$env:PATH = "bin;" + $env:PATH
.\test_multi_precision.exe

Expected Output

HIP Library Path: C:\Windows\SYSTEM32\amdhip64_7.dll
=== RDNA4 Features Detected ===
Architecture: gfx1201
FP8 Native Support: YES
FP4 Support: YES
WMMA (Tensor Cores): YES
GTT Spillover: YES
RT Acceleration: YES
================================
--- Starting RDNA4 Multi-Precision Stress Test ---
[Testing FP16] ... Done.
[Testing FP8]  ... Done.
[Testing FP4]  ... Done.
[Testing FP32] ... Done.
[Testing FP64] ... Done.

=== All 5 RDNA4 Precision Paths Verified ===

---

📊 Performance

Benchmarks on RX 9070 XT (gfx1201) with 16GB VRAM:

Operation	Matrix Size	Precision	TFLOPS	vs. Stock
GEMM	4096x4096	FP16	~45	Baseline
GEMM	4096x4096	FP8	~80	1.8x
GEMM	4096x4096	FP4	~120*	2.7x

*FP4 is experimental and includes dequantization overhead

---

🤝 Contributing

We welcome contributions! Here are some areas where help is needed:

1. Linux Support: Port the build system and test on Linux with ROCm
2. PyTorch Extension: Create Python bindings for easier integration
3. Performance Tuning: Optimize tile sizes for different matrix shapes
4. New Kernels: Add INT8, more sparse formats, Flash Attention
5. Proxy DLL: Implement transparent acceleration for existing apps

Development Setup

git clone https://github.com/YOUR_USERNAME/RDNA4_Instinct_Emulation_Library.git
cd RDNA4_Instinct_Emulation_Library
cmake -B build -G Ninja -DROCM_ENABLE_RDNA4=ON
cmake --build build

---

📄 License

This project is licensed under the MIT License - see the LICENSE file for details.

---

🙏 Acknowledgments

• AMD for ROCm and rocWMMA
• The open-source HIP/ROCm community
• All contributors and testers

---

⚠️ Disclaimer

This is an unofficial, community-driven project. It is not affiliated with or endorsed by AMD. Use at your own risk. The library is provided "as is" without warranty of any kind.

Happy computing! 🎮🚀