TheRock/rocm-libraries/projects/rocblas/.cursor/rules/cpp-style.mdc

---
globs: ["*.cpp", "*.hpp", "*.h", "*.cxx", "*.cc"]
---

# C++ Code Style

> **Note to AI Agents:** If you find code patterns that contradict these style guidelines, suggest rule updates to match actual practices.

## Naming Conventions

- Use snake_case for functions and variables (e.g., `rocblas_gemm`, `matrix_size`)
- Use SCREAMING_SNAKE_CASE for macros and constants (e.g., `ROCBLAS_API_CALL`)
- Use PascalCase for class and struct names when appropriate
- Prefix internal/private functions with underscore or place in detail namespace
- Follow BLAS naming conventions for API functions:
  - Format: `rocblas_<precision><operation>`
  - Examples: `rocblas_sgemm`, `rocblas_daxpy`, `rocblas_cgemv`

## File Headers

- Add copyright header to all source files, updating present year in text:

```cpp
/* ************************************************************************
 * Copyright (C) 2016-[CURRENT YEAR] Advanced Micro Devices, Inc. All rights reserved.
 *
 * Permission is hereby granted, free of charge, to any person obtaining a copy
 * of this software and associated documentation files (the "Software"), to deal
 * in the Software without restriction, including without limitation the rights
 * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell cop-
 * ies of the Software, and to permit persons to whom the Software is furnished
 * to do so, subject to the following conditions:
 *
 * The above copyright notice and this permission notice shall be included in all
 * copies or substantial portions of the Software.
 *
 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IM-
 * PLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS
 * FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR
 * COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER
 * IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNE-
 * CTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
 *
 * ************************************************************************ */
```

- For header files (.h, .hpp), use `#pragma once` for include guards

## Code Practices

- Use explicit type conversions with `static_cast<>` where needed
- Prefer const correctness for function parameters
- **Never use hipMalloc/hipFree** - use `handle->device_malloc()` for device memory (see Device Memory section)
- Use other HIP functions as needed (hipMemcpy, hipStreamSynchronize, etc.)
- Always check return codes from rocBLAS API calls in client code
- Use braces for all control flow statements, even single-line bodies
- Follow const-correctness: use `const` for read-only parameters and member functions
- Use RAII patterns for resource management (memory, pointer mode changes, etc.)

## Device Memory Allocation

**Critical Rule:** rocBLAS code must NOT call `hipMalloc()` or `hipFree()` directly - they are synchronizing APIs.

### Using rocBLAS Device Memory Manager

```cpp
// Allocate device memory using RAII wrapper
auto w_mem = handle->device_malloc(dev_bytes);
if(!w_mem)
    return rocblas_status_memory_error;

// Use the memory
void* workspace = static_cast<void*>(w_mem);
// Memory automatically freed when w_mem goes out of scope
```

### Key Requirements

1. **Allocate upfront:** All device memory must be allocated at the start of the function
2. **RAII pattern:** Use `rocblas_device_malloc` for automatic cleanup
3. **Pass down:** Lower-level kernels receive pre-allocated memory from higher-level routines
4. **Naming:** Use `w_mem`, `workspace`, or `w_` prefix for workspace variables
5. **No nested allocation:** Don't allocate/deallocate at different kernel call levels

### Multiple Memory Blocks

```cpp
// Allocate multiple blocks
auto mem = handle->device_malloc(size1, size2, size3);
if(!mem)
    return rocblas_status_memory_error;

// Access individual blocks
void* ptr1 = mem[0];
void* ptr2 = mem[1];
void* ptr3 = mem[2];
```

## Implementation Pattern (_impl/_launcher)

rocBLAS uses a two-tier pattern for functions needing device memory or error checking:

### Pattern A: Separate Templates (Preferred)

```cpp
// _launcher: Fast computation, no error checking or allocation
template <typename API_INT, typename T>
rocblas_status rocblas_<operation>_launcher(..., void* workspace)
{
    // Launch device kernels
    // No argument checking
    // No logging
    // No device memory allocation
    // Can be called by other computational functions
    return rocblas_status_success;
}

// _impl: Error checking, logging, memory allocation
template <typename API_INT, typename T>
rocblas_status rocblas_<operation>_impl(rocblas_handle handle, ...)
{
    // Argument error checking
    if(!handle) return rocblas_status_invalid_handle;
    
    // Logging
    log_trace(handle, ...);
    
    // Allocate device memory
    auto w_mem = handle->device_malloc(dev_bytes);
    if(!w_mem)
        return rocblas_status_memory_error;
    
    // Call launcher with pre-allocated memory
    return rocblas_<operation>_launcher<API_INT, T>(..., static_cast<void*>(w_mem));
}

// C API wrapper
extern "C" rocblas_status rocblas_<precision><operation>(...)
{
    return rocblas_<operation>_impl<rocblas_int, <type>>(handle, ...);
}
```

### Pattern B: Template Bool Parameter

```cpp
template <bool full_function, typename API_INT, typename T>
rocblas_status rocblas_<operation>_launcher(..., void* workspace = nullptr)
{
    if(full_function)
    {
        // Error checking, logging, memory allocation
        auto w_mem = handle->device_malloc(dev_bytes);
        return rocblas_<operation>_launcher<false, API_INT, T>(..., static_cast<void*>(w_mem));
    }
    // Fast computation path
}
```

Note: Pattern A is preferred for clarity. Pattern B is useful for large existing functions.

## Pointer Mode Patterns

rocBLAS supports scalar parameters in both host and device memory.

### Switching Pointer Mode

```cpp
// Temporarily switch to host pointer mode (RAII)
auto saved_pointer_mode = handle->push_pointer_mode(rocblas_pointer_mode_host);

// Get scalar value regardless of original mode
T alpha_h;
if(saved_pointer_mode == rocblas_pointer_mode_host)
    alpha_h = *alpha;  // Already on host
else
    RETURN_IF_HIP_ERROR(hipMemcpy(&alpha_h, alpha, sizeof(T), hipMemcpyDeviceToHost));

// Use host constants in kernel calls
static constexpr T one = 1;
static constexpr T zero = 0;
rocblas_internal_gemm_template(handle, ..., &one, ..., &zero, ...);

// Pointer mode automatically restored when saved_pointer_mode goes out of scope
```

### load_scalar() for Kernels

```cpp
// HIP kernel that works with both host and device scalars
template <typename T, typename U>
__global__ void axpy_kernel(rocblas_int n, U alpha_device_host, const T* x, T* y)
{
    auto alpha = load_scalar(alpha_device_host);  // Handles both pointer and value
    ptrdiff_t tid = blockIdx.x * blockDim.x + threadIdx.x;
    if(tid < n)
        y[tid] += alpha * x[tid];
}

// Launch with device pointer or host value
if(handle->pointer_mode == rocblas_pointer_mode_device)
    ROCBLAS_LAUNCH_KERNEL(axpy_kernel, ..., alpha, ...);  // Pass pointer
else if(*alpha)
    ROCBLAS_LAUNCH_KERNEL(axpy_kernel, ..., *alpha, ...);  // Pass value
```

## Template Ordering

**Rule:** Place non-type template parameters before type parameters.

```cpp
// Correct: Non-type (NB) before type (T)
template <rocblas_int NB, typename T>
void kernel_calc(const T* data, T* output, rocblas_int n)
{
    // Kernel implementation - T is part of function signature
}

// Caller can omit T - it's deduced from the pointer types
template <rocblas_int NB, typename TPtr>
void kernel_wrapper(TPtr data, TPtr output, rocblas_int n)
{
    // Only specify NB, T is automatically deduced from TPtr
    kernel_calc<NB>(data, output, n);  // T deduced from pointer arguments
}

// Example usage
float* d_data;
float* d_output;
kernel_wrapper<256>(d_data, d_output, n);  // Only NB specified, TPtr=float* deduced

// Incorrect: Type before non-type
template <typename T, rocblas_int NB>  // ❌ Wrong order
void bad_kernel(const T* data, T* output)
{
    // Caller must specify: bad_kernel<float, 256>(data, output)  // Both required!
}
```

**Rationale:** Type parameters can be automatically deduced from function arguments, but non-type parameters cannot. Placing non-type parameters first allows calling templates with only explicit non-type arguments while types are auto-deduced from function parameters.

## RAII Patterns

Use RAII for all resource management:

```cpp
// Memory allocation (automatic cleanup)
auto w_mem = handle->device_malloc(bytes);

// Pointer mode (automatic restore)
auto saved_mode = handle->push_pointer_mode(rocblas_pointer_mode_host);

// One-time initialization
void my_func()
{
    static int dummy = (init_once_function(), 0);
    // init_once_function() called exactly once, thread-safely
}

// Singleton pattern for static variables
static auto& get_table()
{
    static std::unordered_map<std::string, size_t> table;
    return table;
}
```

Never use explicit `new`/`delete`, `malloc`/`free`, or `hipMalloc`/`hipFree`.

## Constants and Macros

### Constants: Use static constexpr

**Use `static constexpr` for constants. Never use macros for constants.**

```cpp
// Template variable for constants
template <typename T>
static constexpr T negative_one = -1;

template <typename T>
static constexpr T zero = 0;

// Inside functions, use static constexpr (not just constexpr)
void my_function()
{
    static constexpr int BLOCK_SIZE = 256;
}
```

### Macros: Only for Control Flow

**Only use macros for control flow constructs where necessary.** Macros with control flow must use `do-while(0)` wrapper.

```cpp
// Macro with control flow - wraps early return
#define RETURN_ZERO_IF_QUERIED(h) \
    do                            \
    {                             \
        if((h)->is_query())       \
            return rocblas_status_success; \
    } while(0)

// Usage
void some_function(rocblas_handle handle)
{
    RETURN_ZERO_IF_QUERIED(handle);  // No semicolon needed with do-while
    // ... rest of function
}
```

**Why do-while(0)?** Ensures the macro behaves like a single statement in all contexts (if/else, etc.).

## API Function Signature Pattern

```cpp
// Standard BLAS API pattern
rocblas_status rocblas_<precision><operation>(
    rocblas_handle handle,
    // operation-specific parameters
    const <type>* input_ptr,
    <type>* output_ptr);

// Example: GEMM
rocblas_status rocblas_sgemm(
    rocblas_handle handle,
    rocblas_operation transA,
    rocblas_operation transB,
    rocblas_int m,
    rocblas_int n,
    rocblas_int k,
    const float* alpha,
    const float* A,
    rocblas_int lda,
    const float* B,
    rocblas_int ldb,
    const float* beta,
    float* C,
    rocblas_int ldc);
```

## Build System

- Use CMake for all build configuration
- Dependencies managed via `install.sh` (Linux) or `rdeps.py` (Windows)
- Tensile integration handled automatically during build
- For new dependencies, update `deps/CMakeLists.txt` and dependency scripts

## Example Code Structure

```cpp
/* ************************************************************************
 * Copyright (C) 2016-2024 Advanced Micro Devices, Inc. All rights reserved.
 * ... (full copyright header)
 * ************************************************************************ */

#pragma once

#include "rocblas.h"

// API function following BLAS naming conventions
rocblas_status rocblas_sgemm(rocblas_handle handle,
                             rocblas_operation transA,
                             rocblas_operation transB,
                             rocblas_int m,
                             rocblas_int n,
                             rocblas_int k,
                             const float* alpha,
                             const float* A,
                             rocblas_int lda,
                             const float* B,
                             rocblas_int ldb,
                             const float* beta,
                             float* C,
                             rocblas_int ldc);

// Internal helper function
namespace detail
{
    template<typename T>
    rocblas_status gemm_template(rocblas_handle handle,
                                  rocblas_operation transA,
                                  rocblas_operation transB,
                                  rocblas_int m,
                                  rocblas_int n,
                                  rocblas_int k,
                                  const T* alpha,
                                  const T* A,
                                  rocblas_int lda,
                                  const T* B,
                                  rocblas_int ldb,
                                  const T* beta,
                                  T* C,
                                  rocblas_int ldc);
}
```

## HIP/GPU Code Patterns

```cpp
// Device function
__device__ void kernel_helper(float* data, int idx)
{
    data[idx] = data[idx] * 2.0f;
}

// Kernel launch
__global__ void my_kernel(float* data, int n)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if(idx < n)
    {
        kernel_helper(data, idx);
    }
}

// Host code launching kernel
hipLaunchKernelGGL(my_kernel,
                   dim3(blocks),
                   dim3(threads),
                   0,  // shared memory
                   stream,
                   data,
                   n);
```