rocBLAS: SWMMAC full precision family + MXFP4/FP8/BF8/UE8M0 types

.cursor/rules/rocblas-architecture.mdc

@@ -0,0 +1,74 @@
+---
+alwaysApply: true
+---
+
+# rocBLAS Project Architecture
+
+> **Note to AI Agents:** If you notice these rules don't match the actual codebase or could be improved, please suggest updates. See `.cursorrules` for details.
+
+rocBLAS is the AMD ROCm Basic Linear Algebra Subprograms (BLAS) library, implemented in HIP and optimized for AMD GPUs.
+
+## Key Documentation
+
+- **[README](./README.md)** - Project overview and requirements
+- **[Linux Install Guide](./docs/install/Linux_Install_Guide.rst)** - Linux build and installation
+- **[Windows Install Guide](./docs/install/Windows_Install_Guide.rst)** - Windows build and installation
+- **[Programmer's Guide](./docs/how-to/Programmers_Guide.rst)** - API usage and programming guide
+- **[Design Notes](./docs/conceptual/rocblas-design-notes.rst)** - Architecture and design decisions
+
+## Component Architecture
+
+| Component | Location | Purpose |
+|-----------|----------|---------|
+| **Library** (`library/`) | Core implementation | BLAS operations (Level 1, 2, 3, Extensions) |
+| **Clients** (`clients/`) | Testing & benchmarking | Test suite (gtest), benchmarks, samples |
+| **Tensile** (`library/src/blas3/`) | GEMM kernels | Optimized matrix multiplication kernels |
+| **Dependencies** (`deps/`) | External deps | GTest, LAPACK dependencies |
+| **Scripts** (`scripts/`) | Utilities | Performance testing, YAML generation |
+
+## BLAS Operation Levels
+
+### Level 1: Vector-Vector Operations
+- Location: `library/src/blas1/`
+- Examples: `axpy`, `dot`, `scal`, `nrm2`
+- Characteristics: O(n) complexity, vector inputs/outputs
+
+### Level 2: Matrix-Vector Operations
+- Location: `library/src/blas2/`
+- Examples: `gemv`, `ger`, `trmv`, `symv`
+- Characteristics: O(n²) complexity, matrix-vector operations
+
+### Level 3: Matrix-Matrix Operations
+- Location: `library/src/blas3/`
+- Examples: `gemm`, `trmm`, `symm`, `syrk`
+- Characteristics: O(n³) complexity, highly optimized via Tensile
+- **Tensile Integration:** Generates optimized GPU kernels for GEMM operations
+
+### BLAS Extensions
+- Location: `library/src/blas_ex/`
+- Examples: `gemm_ex`, `gemm_strided_batched_ex`
+- Purpose: Extended precision, batched operations, mixed precision
+
+## Tensile Integration
+
+Tensile is a kernel generator that creates highly optimized GEMM (General Matrix Multiply) kernels:
+
+- **Configuration:** YAML files in `library/src/blas3/Tensile/Logic/`
+- **Build-Time Generation:** Kernels generated during build process
+- **Architecture-Specific:** Optimized for specific GPU architectures (gfx90a, gfx942, etc.)
+- **Performance Critical:** GEMM is the most performance-critical BLAS operation
+
+## Data Type Support
+
+rocBLAS supports multiple data types with naming conventions:
+
+- **s** - Single precision float (float)
+- **d** - Double precision float (double)
+- **c** - Complex single precision (rocblas_float_complex)
+- **z** - Complex double precision (rocblas_double_complex)
+- **h** - Half precision (rocblas_half)
+- **bf16** - Brain float 16 (rocblas_bfloat16)
+- **i8** - 8-bit integer (int8_t)
+- **i32** - 32-bit integer (int32_t)
+
+Example: `rocblas_sgemm` = single precision GEMM, `rocblas_daxpy` = double precision AXPY

.cursor/rules/testing.mdc

@@ -0,0 +1,297 @@
+---
+globs: ["*.cpp", "*.hpp", "*.h", "clients/gtest/**", "clients/benchmarks/**"]
+---
+
+# Testing Guidelines
+
+> **Note to AI Agents:** If testing patterns have evolved or examples are misleading, suggest updates to these guidelines.
+
+## Framework
+
+- Use Google Test (gtest) framework for all C/C++ tests
+- Test files located in `clients/gtest/`
+- YAML test configurations in `clients/gtest/*.yaml`
+- Use `TEST()` or `TEST_P()` macros for parameterized tests
+- Never generate a main() function in test files - gtest provides its own
+
+## Test Organization
+
+### Test File Structure
+
+**Note on API Usage:** Tests use `rocblas_gemm<T>` which is a **test infrastructure wrapper** that internally dispatches to precision-specific functions (rocblas_sgemm, rocblas_dgemm, etc.). End users of the rocBLAS library should call the precision-specific functions directly (e.g., `rocblas_sgemm` for float, `rocblas_dgemm` for double).
+
+```cpp
+/* ************************************************************************
+ * Copyright (C) 2016-2024 Advanced Micro Devices, Inc. All rights reserved.
+ * ... (full copyright header)
+ * ************************************************************************ */
+
+#include "rocblas_test.hpp"
+#include <gtest/gtest.h>
+
+// Template test function using test infrastructure
+template <typename T>
+void testing_gemm(const Arguments& arg)
+{
+    // Test infrastructure provides a templated wrapper around precision-specific APIs
+    // rocblas_gemm<T> internally calls rocblas_sgemm, rocblas_dgemm, etc.
+    auto rocblas_gemm_fn = arg.api & c_API_FORTRAN ? rocblas_gemm<T, true> : rocblas_gemm<T, false>;
+
+    rocblas_local_handle handle{arg};
+    rocblas_int m = arg.M;
+    rocblas_int n = arg.N;
+    rocblas_int k = arg.K;
+    rocblas_int lda = arg.lda;
+    rocblas_int ldb = arg.ldb;
+    rocblas_int ldc = arg.ldc;
+
+    T h_alpha = arg.get_alpha<T>();
+    T h_beta = arg.get_beta<T>();
+
+    rocblas_operation transA = rocblas_operation_none;
+    rocblas_operation transB = rocblas_operation_none;
+
+    // Allocate host and device memory using test infrastructure
+    HOST_MEMCHECK(host_matrix<T>, hA, (m, k, lda));
+    HOST_MEMCHECK(host_matrix<T>, hB, (k, n, ldb));
+    HOST_MEMCHECK(host_matrix<T>, hC, (m, n, ldc));
+
+    DEVICE_MEMCHECK(device_matrix<T>, dA, (m, k, lda));
+    DEVICE_MEMCHECK(device_matrix<T>, dB, (k, n, ldb));
+    DEVICE_MEMCHECK(device_matrix<T>, dC, (m, n, ldc));
+
+    // Initialize data on host
+    rocblas_init_matrix(hA, arg, rocblas_client_general_matrix);
+    rocblas_init_matrix(hB, arg, rocblas_client_general_matrix);
+    rocblas_init_matrix(hC, arg, rocblas_client_general_matrix);
+
+    // Copy data from CPU to device
+    CHECK_HIP_ERROR(dA.transfer_from(hA));
+    CHECK_HIP_ERROR(dB.transfer_from(hB));
+    CHECK_HIP_ERROR(dC.transfer_from(hC));
+
+    // Execute operation via test wrapper
+    CHECK_ROCBLAS_ERROR(rocblas_set_pointer_mode(handle, rocblas_pointer_mode_host));
+    DAPI_CHECK(rocblas_gemm_fn, 
+               (handle, transA, transB, m, n, k, &h_alpha, dA, lda, dB, ldb, &h_beta, dC, ldc));
+
+    // Verify results (unit_check, norm_check, etc.)
+    if(arg.unit_check || arg.norm_check)
+    {
+        // Validation code here
+    }
+}
+
+// Test instantiation via YAML or direct arguments
+TEST(gemm_gtest, float)
+{
+    Arguments arg;
+    arg.M = 128;
+    arg.N = 128;
+    arg.K = 128;
+    arg.lda = 128;
+    arg.ldb = 128;
+    arg.ldc = 128;
+    testing_gemm<float>(arg);
+}
+```
+
+## Test Naming Conventions
+
+### Test Suite Naming
+
+Format: `<operation>_<variant>_gtest`
+
+Examples:
+- `gemm_gtest` - Basic GEMM tests
+- `gemm_strided_batched_gtest` - Strided batched GEMM tests
+- `axpy_gtest` - AXPY tests
+
+### Test Case Naming
+
+Format: `<datatype>_<variant>`
+
+Examples:
+- `float` - Single precision test
+- `double` - Double precision test
+- `complex_float` - Complex single precision
+- `half_precision` - Half precision (fp16)
+
+### Test Categories (via gtest filters)
+
+- `*quick*` - Fast tests for rapid iteration (< 1 minute)
+- `*pre_checkin*` - Tests for PR validation (< 1 hour)
+- `*nightly*` - Comprehensive regression tests (1-2 hours)
+- `*known_bug*` - Tests for known issues (excluded from runs)
+
+## YAML Test Configuration
+
+YAML files define test parameters for comprehensive testing:
+
+```yaml
+---
+include: rocblas_common.yaml
+---
+
+Tests:
+- name: gemm_small
+  category: quick
+  function: gemm
+  precision: *single_double_precisions
+  transA: N
+  transB: N
+  M: [128, 256]
+  N: [128, 256]
+  K: [128, 256]
+  alpha: 1.0
+  beta: 0.0
+  lda: [128, 256]
+  ldb: [128, 256]
+  ldc: [128, 256]
+```
+
+## Test Execution Patterns
+
+### Using rtest.py
+
+```bash
+# Quick smoke tests
+python3 rtest.py -t smoke
+
+# Pre-submit tests (PR validation)
+python3 rtest.py -t psdb
+
+# Nightly regression
+python3 rtest.py -t osdb
+
+# Complete quality engineering
+python3 rtest.py -t cqe
+```
+
+### Direct Test Execution
+
+```bash
+# Run all tests
+./build/release/clients/staging/rocblas-test
+
+# Run specific operation
+./build/release/clients/staging/rocblas-test --gtest_filter=*gemm*
+
+# Run quick tests only
+./build/release/clients/staging/rocblas-test --gtest_filter=*quick*
+
+# Run with YAML
+./build/release/clients/staging/rocblas-test --yaml clients/gtest/rocblas_smoke.yaml
+
+# Exclude known bugs
+./build/release/clients/staging/rocblas-test --gtest_filter=*quick*-*known_bug*
+```
+
+## Verification Patterns
+
+### Numerical Verification
+
+```cpp
+// Helper to get epsilon for real and complex types (from clients/include/unit.hpp)
+template <typename T, std::enable_if_t<(!rocblas_is_complex<T>), int> = 0>
+constexpr double get_epsilon()
+{
+    return std::numeric_limits<T>::epsilon();
+}
+
+template <typename T, std::enable_if_t<(rocblas_is_complex<T>), int> = 0>
+constexpr auto get_epsilon()
+{
+    // For complex types, get epsilon of the underlying real type
+    return get_epsilon<decltype(std::real(T{}))>();
+}
+
+// Compare against reference implementation (works for real and complex types)
+template <typename T>
+void verify_gemm_result(const T* computed,
+                        const T* reference,
+                        rocblas_int m,
+                        rocblas_int n,
+                        rocblas_int ldc)
+{
+    double tolerance = get_epsilon<T>() * 10;
+
+    for(rocblas_int i = 0; i < m; i++)
+    {
+        for(rocblas_int j = 0; j < n; j++)
+        {
+            auto idx = i + j * ldc;
+            // std::abs returns magnitude for both real and complex types
+            auto diff = std::abs(computed[idx] - reference[idx]);
+            EXPECT_NEAR(diff, 0.0, tolerance);
+        }
+    }
+}
+```
+
+### Error Handling Tests
+
+```cpp
+TEST(gemm_gtest, invalid_handle)
+{
+    rocblas_status status = rocblas_sgemm(nullptr, ...);
+    EXPECT_EQ(status, rocblas_status_invalid_handle);
+}
+
+TEST(gemm_gtest, invalid_size)
+{
+    rocblas_handle handle;
+    rocblas_create_handle(&handle);
+
+    rocblas_status status = rocblas_sgemm(handle,
+                                          rocblas_operation_none,
+                                          rocblas_operation_none,
+                                          -1,  // invalid m
+                                          128,
+                                          128,
+                                          ...);
+    EXPECT_EQ(status, rocblas_status_invalid_size);
+
+    rocblas_destroy_handle(handle);
+}
+```
+
+## Benchmarking
+
+### Using rocblas-bench
+
+```bash
+# GEMM benchmark
+./build/release/clients/staging/rocblas-bench -f gemm -r f32_r -m 4096 -n 4096 -k 4096
+
+# GEMV benchmark
+./build/release/clients/staging/rocblas-bench -f gemv -r f32_r -m 8192 -n 8192
+
+# Batched operations
+./build/release/clients/staging/rocblas-bench -f gemm_strided_batched -r f32_r -m 256 -n 256 -k 256 --batch_count 100
+
+# Load from YAML
+./build/release/clients/staging/rocblas-bench --yaml scripts/performance/gemm_nn.yaml
+```
+
+## Environment Variables for Testing
+
+```bash
+# Enable verbose logging
+export ROCBLAS_LAYER=1
+
+# Enable numerical checks (1-4, higher = more checks)
+export ROCBLAS_CHECK_NUMERICS=4
+
+# Enable profiling
+export ROCBLAS_LAYER=4
+```
+
+## Best Practices
+
+1. **Test Coverage:** Cover normal cases, edge cases, and error conditions
+2. **Data Types:** Test all supported precisions (s, d, c, z, h, bf16)
+3. **Sizes:** Test small, medium, large, and edge case sizes
+4. **Batching:** Test both single and batched operations where applicable
+5. **Reference Implementation:** Use AOCL-BLAS or OpenBLAS for verification
+6. **Performance:** Use `rocblas-bench` for performance regression testing