--- 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_` - 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(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 rocblas_status rocblas__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 rocblas_status rocblas__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__launcher(..., static_cast(w_mem)); } // C API wrapper extern "C" rocblas_status rocblas_(...) { return rocblas__impl>(handle, ...); } ``` ### Pattern B: Template Bool Parameter ```cpp template rocblas_status rocblas__launcher(..., void* workspace = nullptr) { if(full_function) { // Error checking, logging, memory allocation auto w_mem = handle->device_malloc(dev_bytes); return rocblas__launcher(..., static_cast(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 __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 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 void kernel_wrapper(TPtr data, TPtr output, rocblas_int n) { // Only specify NB, T is automatically deduced from TPtr kernel_calc(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 // ❌ Wrong order void bad_kernel(const T* data, T* output) { // Caller must specify: bad_kernel(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 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 static constexpr T negative_one = -1; template 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_( rocblas_handle handle, // operation-specific parameters const * input_ptr, * 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 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); ```