# From CUDA to HIP: A Comprehensive Guide to GPU Code Portability **A practical guide for migrating CUDA projects to support both NVIDIA and AMD GPUs** --- ## Table of Contents 1. [Introduction](#introduction) 2. [Understanding HIP (Heterogeneous-Compute Interface for Portability)](#understanding-hip) 3. [Two Approaches to Hipification](#two-approaches-to-hipification) 4. [Environment Setup and Dependencies](#environment-setup-and-dependencies) 5. [Approach 1: Unified Codebase (Recommended)](#approach-1-unified-codebase-recommended) 6. [Case Study: cuda-battery to hip-battery](#case-study-cuda-battery-to-hip-battery) 7. [Approach 2: Separate Codebases](#approach-2-separate-codebases) 8. [Decision Matrix: Which Approach to Choose?](#decision-matrix-which-approach-to-choose) 9. [Conclusion](#conclusion) --- ## Introduction Graphics Processing Units (GPUs) have become essential for high-performance computing, machine learning, scientific simulations, and many other compute-intensive applications. However, the GPU landscape is dominated by two major vendors with incompatible programming models: - **NVIDIA GPUs**: Programmed using CUDA (Compute Unified Device Architecture) - **AMD GPUs**: Programmed using ROCm (Radeon Open Compute) and HIP If you've invested time developing CUDA applications, you might want to make your code portable to AMD GPUs without maintaining completely separate codebases. This is where **HIP (Heterogeneous-Compute Interface for Portability)** comes in. This guide will walk you through two different strategies for making your CUDA code portable, their trade-offs, and provide step-by-step instructions for implementation. ### Who Is This Guide For? - Developers with existing CUDA codebases who want AMD GPU support - Teams maintaining GPU-accelerated libraries - Researchers wanting to maximize the reach of their scientific computing code - Anyone curious about GPU programming portability ### Prerequisites - Basic understanding of C++ programming - Familiarity with CUDA concepts (kernels, device functions, memory management) - Linux operating system (Ubuntu 20.04+ recommended) - Basic command-line knowledge --- ## Understanding HIP (Heterogeneous-Compute Interface for Portability) ### What is HIP? HIP is a C++ runtime API and kernel language developed by AMD that allows developers to write portable GPU code. The key insight behind HIP is that CUDA and ROCm programming models are very similar, differing mainly in: - API function names (e.g., `cudaMalloc` vs `hipMalloc`) - Header includes (e.g., `cuda_runtime.h` vs `hip_runtime.h`) - Error codes and types HIP provides: 1. **A CUDA-like API** that abstracts both NVIDIA and AMD backends 2. **Automatic code conversion tools** (`hipify-perl`, `hipify-clang`) 3. **Runtime backend selection** to target different GPU vendors ### How HIP Works ``` HIP Code | | +------------+------------+ | | [NVIDIA Backend] [AMD Backend] | | Uses NVCC/CUDA Uses ROCm/HIP on NVIDIA GPUs on AMD GPUs ``` When compiling for NVIDIA GPUs, HIP code is translated to CUDA and compiled with `nvcc`. When compiling for AMD GPUs, the same HIP code is compiled directly with ROCm's compiler. ### Key Terminology - **hipify**: Tools that convert CUDA code to HIP code - **hipcc**: HIP compiler wrapper that detects the target platform - **ROCm**: AMD's open-source GPU computing platform - **Backend**: The underlying GPU platform (NVIDIA CUDA or AMD ROCm) - **Device**: The GPU hardware - **Host**: The CPU --- ## Two Approaches to Hipification When porting CUDA code to support multiple GPU vendors, you have two primary strategies: ### Quick Comparison | Aspect | Unified Codebase (HIP) | Separate Codebases (CUDA + HIP) | |--------|----------------------|--------------------------------| | **Code Maintenance** | Single codebase | Two codebases to maintain | | **Compilation** | One build system | Conditional compilation or separate builds | | **End-User Dependencies** | ROCm headers needed even for NVIDIA builds* | Only platform-specific dependencies | | **Developer Dependencies** | CUDA + ROCm (to test both platforms) | CUDA + ROCm (to test both platforms) | | **Performance** | Identical to native | Optimized per platform | | **Development Effort** | Lower long-term | Higher long-term | | **Testing Complexity** | Simpler (one codebase) | More complex (two codebases) | | **Best For** | Most projects | Performance-critical code with platform-specific optimizations | \* *End-users on NVIDIA hardware need ROCm headers but not ROCm runtime; AMD users need full ROCm* --- ## Environment Setup and Dependencies Before implementing either approach, you'll need to set up your development environment. This section covers all the tools and dependencies you may need. **Dependencies vary based on whether you're developing (testing both platforms) or building for end-users (single platform).** ### Overview of Dependencies The table below shows which dependencies are needed for different scenarios: | Dependency | **Development
(Both Approaches)** | **Building for NVIDIA** | **Building for AMD** | Purpose | |------------|:-------------------------------------:|:-----------------------:|:--------------------:|----------| | | | **Approach 1** / **Approach 2** | **Approach 1** / **Approach 2** | | | **HIP/ROCm** | ✅ Required | ✅ Required* / ❌ Not needed | ✅ Required / ✅ Required | HIP runtime and development libraries | | **HIPIFY** | ✅ Required | ✅ Required / ❌ Not needed** | ✅ Required / ❌ Not needed** | Automatic CUDA to HIP conversion tool | | **CUDA Toolkit** | ✅ Required | ✅ Required / ✅ Required | ❌ Not needed / ❌ Not needed | NVIDIA GPU support | | **CMake** | ✅ Required | ✅ Required / ✅ Required | ✅ Required / ✅ Required | Build system configuration | | **Build Tools** | ✅ Required | ✅ Required / ✅ Required | ✅ Required / ✅ Required | Compilers and development utilities | **Key Differences:** - **Development**: Developers need ALL dependencies for both approaches to verify code works on both NVIDIA and AMD GPUs - **Approach 1 (Unified)**: - \* End-users on NVIDIA hardware need hip-runtime-nvidia and hip-dev packages (lightweight HIP layer over CUDA) instead of full ROCm - End-users on AMD hardware need full ROCm stack - Single codebase means HIP is required on all platforms - **Approach 2 (Separate)**: - \*\* HIPIFY only needed during initial code conversion, not for building - End-users only install dependencies for their specific platform - NVIDIA users: CUDA only (no HIP/ROCm needed) - AMD users: Full ROCm/HIP stack --- ### Step 1: Install Build Essentials **Required for:** All approaches Install essential development tools and compilers: ```bash sudo apt update sudo apt install build-essential g++ gdb make git ``` **Verify installation:** ```bash g++ --version make --version git --version ``` --- ### Step 2: Install CMake (Build System) **Required for:** All approaches HIP projects typically use CMake (version 3.24 or higher). > **📚 Official Documentation:** > - [CMake Download Page](https://cmake.org/download/) > - [CMake Documentation](https://cmake.org/documentation/) > - [CMake Tutorial](https://cmake.org/cmake/help/latest/guide/tutorial/index.html) ```bash # Remove old cmake if installed sudo apt remove cmake # Install latest CMake wget https://github.com/Kitware/CMake/releases/download/v4.2.3/cmake-4.2.3-linux-x86_64.sh chmod +x cmake-4.2.3-linux-x86_64.sh sudo ./cmake-4.2.3-linux-x86_64.sh --prefix=/usr/local --skip-license # Verify cmake --version ``` --- ### Step 3: Install CUDA Toolkit (NVIDIA Platform Support) **Required for:** Approach 1 (NVIDIA builds), Approach 2 (CUDA code path) If you want to build for NVIDIA GPUs, install the CUDA Toolkit. > **📚 Official Documentation:** > - [CUDA Toolkit Download](https://developer.nvidia.com/cuda-downloads) > - [CUDA Installation Guide for Linux](https://docs.nvidia.com/cuda/cuda-installation-guide-linux/) > - [CUDA Quick Start Guide](https://docs.nvidia.com/cuda/cuda-quick-start-guide/) **On Ubuntu:** ```bash # Option 1: Use distribution packages (easier but older version) sudo apt install nvidia-cuda-toolkit # Option 2: Install from NVIDIA (latest version) # Download installer from https://developer.nvidia.com/cuda-downloads wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2204/x86_64/cuda-keyring_1.1-1_all.deb sudo dpkg -i cuda-keyring_1.1-1_all.deb sudo apt update sudo apt install cuda-toolkit-12-3 # Add CUDA to PATH echo 'export PATH=/usr/local/cuda/bin:$PATH' >> ~/.bashrc echo 'export LD_LIBRARY_PATH=/usr/local/cuda/lib64:$LD_LIBRARY_PATH' >> ~/.bashrc source ~/.bashrc ``` **Verify CUDA installation:** ```bash nvcc --version nvidia-smi ``` --- ### Step 4: Install HIP and ROCm HIP is required for developing portable GPU code. The installation method differs depending on whether you're targeting NVIDIA or AMD GPUs. > **📚 Official Documentation:** > - [HIP Installation Guide](https://rocm.docs.amd.com/projects/HIP/en/latest/install/install.html) > - [ROCm Installation Guide](https://rocm.docs.amd.com/projects/install-on-linux/en/latest/) > - [ROCm Quick Start Guide](https://rocm.docs.amd.com/en/latest/deploy/linux/quick_start.html) #### Option A: HIP on NVIDIA Platform **Prerequisites:** CUDA Toolkit must be installed first (see [Step 3](#step-3-install-cuda-toolkit-nvidia-platform-support)) **1. Download and install amdgpu-install package:** Choose the appropriate version for your Ubuntu: ```bash # For Ubuntu 24.04 wget https://repo.radeon.com/amdgpu-install/6.2.4/ubuntu/noble/amdgpu-install_6.2.60204-1_all.deb sudo dpkg -i amdgpu-install_6.2.60204-1_all.deb # For Ubuntu 22.04 wget https://repo.radeon.com/amdgpu-install/6.2.4/ubuntu/jammy/amdgpu-install_6.2.60204-1_all.deb sudo dpkg -i amdgpu-install_6.2.60204-1_all.deb # Update package list sudo apt update ``` **2. Install HIP runtime for NVIDIA and development packages:** ```bash # This installs the CUDA SDK and HIP porting layer sudo apt install hip-runtime-nvidia hip-dev ``` **3. Set HIP_PLATFORM environment variable:** ```bash # Set platform to nvidia echo 'export HIP_PLATFORM=nvidia' >> ~/.bashrc # Add HIP and CUDA to PATH echo 'export PATH=/opt/rocm/bin:/usr/local/cuda/bin:$PATH' >> ~/.bashrc echo 'export LD_LIBRARY_PATH=/opt/rocm/lib:/usr/local/cuda/lib64:$LD_LIBRARY_PATH' >> ~/.bashrc source ~/.bashrc ``` **Default installation paths:** - CUDA SDK: `/usr/local/cuda` - HIP: `/opt/rocm` #### Option B: HIP on AMD Platform (Full ROCm) **For Ubuntu 22.04/24.04:** ```bash # Add ROCm repository wget -q -O - https://repo.radeon.com/rocm/rocm.gpg.key | sudo apt-key add - echo 'deb [arch=amd64] https://repo.radeon.com/rocm/apt/debian/ ubuntu main' | \ sudo tee /etc/apt/sources.list.d/rocm.list # Update package list sudo apt update # Install full ROCm stack with HIP development packages sudo apt install rocm-hip-sdk hip-dev rocm-dev # Add ROCm to your PATH echo 'export PATH=/opt/rocm/bin:$PATH' >> ~/.bashrc echo 'export LD_LIBRARY_PATH=/opt/rocm/lib:$LD_LIBRARY_PATH' >> ~/.bashrc source ~/.bashrc ``` #### Verify Installation **Common verification (both platforms):** ```bash hipcc --version hip-config --version ``` **AMD platform only:** ```bash rocminfo rocm-smi ``` --- ### Step 5: Build HIPIFY (Code Conversion Tool) HIPIFY is essential for converting CUDA code to HIP. ROCm typically does not include HIPIFY by default, so you'll need to build it from source. > **📚 Official Documentation:** > - [Building HIPIFY on Linux](https://rocm.docs.amd.com/projects/HIPIFY/en/latest/building/build-hipify-clang-linux.html) **Step 5.1: Build LLVM and Clang** HIPIFY requires LLVM and Clang. We'll build them from source: ```bash # Create a root directory for building LLVM, Clang and HIPIFY export ROOT_DIR=$(pwd)/hipify-build mkdir -p $ROOT_DIR cd $ROOT_DIR # Clone LLVM project (shallow clone for faster download) # If you need full git history, remove the `--depth 1` option git clone --depth 1 https://github.com/llvm/llvm-project.git # Create build and installation directories mkdir build dist cd build # Configure LLVM build cmake \ -DCMAKE_INSTALL_PREFIX=$ROOT_DIR/dist \ -DLLVM_TARGETS_TO_BUILD="X86" \ -DLLVM_ENABLE_PROJECTS="clang" \ -DLLVM_INCLUDE_TESTS=OFF \ -DCMAKE_BUILD_TYPE=Release \ ../llvm-project/llvm # Build and install (this may take 30-60 minutes) make -j$(nproc) install ``` > **⚠️ Important Notes:** > - If you build LLVM in Debug mode (`-DCMAKE_BUILD_TYPE=Debug`), you must also build HIPIFY in Debug mode > - Only 64-bit build mode is supported > - Building LLVM requires significant disk space (~20GB) and time **Step 5.2: Build HIPIFY** Now that LLVM is built, we can build HIPIFY: ```bash # Return to root directory cd $ROOT_DIR # Clone HIPIFY repository git clone https://github.com/ROCm/HIPIFY.git # Navigate to build directory cd build # Configure HIPIFY build # CMAKE_PREFIX_PATH ensures LLVM is found cmake \ -DCMAKE_INSTALL_PREFIX=$ROOT_DIR/dist \ -DCMAKE_BUILD_TYPE=Release \ -DCMAKE_PREFIX_PATH=$ROOT_DIR/dist \ ../HIPIFY # Build and install (much faster than LLVM) make -j$(nproc) install ``` **Step 5.3: Add HIPIFY to PATH** ```bash # Add hipify-clang to your PATH echo "export PATH=$ROOT_DIR/dist/bin:\$PATH" >> ~/.bashrc source ~/.bashrc ``` **Verify HIPIFY installation:** ```bash hipify-clang --version hipify-perl --version # Check what was installed in the build directory ls $ROOT_DIR/dist/bin # Output: hipify-clang hipify-perl ``` The binaries will be located at: - `$ROOT_DIR/dist/bin/hipify-clang` - Main conversion tool - `$ROOT_DIR/dist/bin/hipify-perl` - Perl-based converter > **💡 Tip:** If you encounter issues, verify that: > - LLVM was built successfully > - CMAKE_PREFIX_PATH points to the correct LLVM installation > - You're using the same build type (Release/Debug) for both LLVM and HIPIFY **Step 5.4: Access HIPIFY Helper Scripts (Optional but Recommended)** The HIPIFY source code includes convenient helper scripts that are **NOT copied to the build/install directory** (`dist/bin`). These scripts remain in the **source directory** and must be accessed separately. They can significantly simplify batch conversion tasks. > **⚠️ Important:** The helper scripts are in the **source directory** (`$ROOT_DIR/HIPIFY/bin`), not in the build directory (`$ROOT_DIR/dist/bin`). Only the compiled binaries (`hipify-clang` and `hipify-perl`) are installed to the build directory. ```bash # List available scripts in HIPIFY source directory (NOT in build directory) ls $ROOT_DIR/HIPIFY/bin # Output shows helper scripts: # findcode.sh - Find CUDA code patterns # findcode_headers.sh - Find CUDA headers # findcode_sources.sh - Find CUDA source files # findcode_custom.sh - Custom pattern finding # finduncodep.sh - Find unconverted code # hipconvertinplace.sh - Convert entire directory with hipify-clang # hipconvertinplace-perl.sh - Convert entire directory with hipify-perl # hipexamine.sh - Examine conversion statistics # hipexamine-perl.sh - Examine with perl variant # hipify-perl - Perl-based converter (also available in dist/bin) ``` **Add helper scripts to PATH:** The helper scripts need to be added to your PATH separately from the main binaries: ```bash # Add HIPIFY source scripts to your PATH for easy access # This is the SOURCE directory, not the build directory echo "export PATH=$ROOT_DIR/HIPIFY/bin:\$PATH" >> ~/.bashrc source ~/.bashrc # Verify script access which hipconvertinplace.sh # Should output: /path/to/hipify-build/HIPIFY/bin/hipconvertinplace.sh ``` > **⚠️ Directory Structure Summary:** > - **Build binaries** → `$ROOT_DIR/dist/bin/` (hipify-clang, hipify-perl) > - **Helper scripts** → `$ROOT_DIR/HIPIFY/bin/` (all .sh scripts) > - Both should be in your PATH for full functionality **Common Script Usage Examples:** 1. **Convert an entire directory in-place:** ```bash # Navigate to your CUDA project cd ~/projects/your-cuda-project # Convert all CUDA files in current directory and subdirectories hipconvertinplace.sh . # Or convert a specific directory hipconvertinplace.sh src/ ``` 2. **Examine what would be converted (dry-run):** ```bash # Show statistics about what will be converted without modifying files hipexamine.sh . # Example output shows: # - Number of CUDA API calls found # - Conversion statistics # - Potential issues ``` 3. **Find all CUDA source files:** ```bash # Find all .cu and .cuh files findcode_sources.sh . # Find all CUDA headers findcode_headers.sh . ``` 4. **Find unconverted CUDA code after conversion:** ```bash # Check if any CUDA references remain finduncodep.sh . # Useful for verifying complete conversion ``` 5. **Custom pattern search:** ```bash # Search for specific CUDA patterns findcode_custom.sh "cudaMalloc\|cudaFree" . ``` > **⚠️ Important Notes:** > - `hipconvertinplace.sh` uses `hipify-clang` (requires it to be in PATH from `$ROOT_DIR/dist/bin`) > - `hipconvertinplace-perl.sh` uses `hipify-perl` (also needs to be in PATH) > - These scripts modify files in-place, so always work on a backed-up copy or version-controlled repository > - The scripts recursively process all matching files in the specified directory > - Make sure both `$ROOT_DIR/dist/bin` AND `$ROOT_DIR/HIPIFY/bin` are in your PATH **Example workflow with helper scripts:** ```bash # Ensure both paths are set # export PATH=$ROOT_DIR/dist/bin:$ROOT_DIR/HIPIFY/bin:$PATH # 1. Examine your CUDA codebase cd ~/projects/my-cuda-project hipexamine.sh . > conversion-report.txt # 2. Create a backup branch git checkout -b cuda-to-hip-conversion # 3. Convert the code hipconvertinplace.sh . # 4. Check for any remaining CUDA code finduncodep.sh . # 5. Review changes git diff # 6. Test and commit git add . git commit -m "Convert CUDA to HIP using hipify-clang" ``` --- ## Approach 1: Unified Codebase (Recommended) **Strategy**: Convert all CUDA code to HIP and maintain a single codebase that compiles for both NVIDIA and AMD GPUs. ### Advantages ✅ **Single Source of Truth**: Only one codebase to maintain, debug, and enhance ✅ **Simplified Testing**: Test once, run on any GPU vendor ✅ **Easier Collaboration**: Developers don't need to know which platform they're targeting ✅ **Reduced Technical Debt**: No divergence between platform-specific implementations ✅ **Future-Proof**: Easy to add support for new GPU vendors if HIP expands ### Disadvantages ❌ **Dependency Overhead for End-Users**: Users on NVIDIA systems need to install ROCm headers (though not ROCm runtime) even if only using NVIDIA GPUs ❌ **Less Platform-Specific Optimization**: Can't easily use vendor-specific features without conditional compilation ❌ **Build System Complexity**: Need to properly configure HIP platform detection > **📝 Note for Developers**: When developing with Approach 1, you'll need both CUDA and ROCm installed to test on both platforms and ensure you haven't broken anything. ### When to Use This Approach - You want maximum portability with minimum maintenance - Your code doesn't rely heavily on vendor-specific features - Your end-users can tolerate installing ROCm headers on NVIDIA systems (ROCm runtime not required) - You're starting a new project or library - Development team bandwidth is limited --- ## Approach 1: Step-by-Step Implementation Guide ### Required Dependencies Before implementing this approach, ensure you have completed the following from the [Environment Setup and Dependencies](#environment-setup-and-dependencies) section: ✅ **[Step 1: Install Build Essentials](#step-1-install-build-essentials)** - Required for compilation ✅ **[Step 2: Install CMake](#step-2-install-cmake-build-system)** - Required for building the project ✅ **[Step 3: Install CUDA Toolkit](#step-3-install-cuda-toolkit-nvidia-platform-support)** - Optional (only if targeting NVIDIA GPUs) ✅ **[Step 4: Install HIP and ROCm](#step-4-install-hip-and-rocm)** - Required for HIP runtime and development ✅ **[Step 5: Build HIPIFY](#step-5-build-hipify-code-conversion-tool)** - Required for converting CUDA code to HIP > **💡 Quick Start:** If you only want to build for AMD GPUs, you can skip the CUDA Toolkit installation. However, for maximum portability (supporting both NVIDIA and AMD), install all dependencies. --- ### Phase 1: Code Conversion #### Step 1.1: Automatic Conversion with HIPIFY AMD provides tools to automatically convert CUDA code to HIP. We'll use `hipify-clang` which we built in the [Environment Setup](#step-5-build-hipify-code-conversion-tool). > **📚 Official Documentation:** > - [HIPIFY Tools Documentation](https://rocm.docs.amd.com/projects/HIPIFY/en/latest/) > - [HIPIFY GitHub Repository](https://github.com/ROCm/HIPIFY) > - [HIP Porting Guide](https://rocm.docs.amd.com/en/latest/how-to/hip-porting-guide.html) **Option 1: Using hipify-clang (Recommended)** `hipify-clang` is more accurate than `hipify-perl` as it uses Clang's AST for parsing. ```bash # Clone your CUDA project cd ~/projects git clone cd your-cuda-project # Create a backup git checkout -b pre-hipification # Run hipify-clang on source files find include/ src/ -name "*.cu" -o -name "*.cuh" -o -name "*.hpp" -o -name "*.cpp" | \ xargs -I {} hipify-clang {} -o {} --print-stats # Review changes git diff ``` **Option 2: Using HIPIFY Helper Scripts (Easiest for Entire Projects)** The most convenient method is to use the helper scripts from the HIPIFY source (see [Environment Setup - Step 5.4](#step-5-build-hipify-code-conversion-tool)). These scripts are located in the source directory, not the build directory. > **⚠️ Prerequisite:** Ensure helper scripts are in your PATH: > ```bash > export PATH=$ROOT_DIR/HIPIFY/bin:$PATH # Helper scripts (source directory) > export PATH=$ROOT_DIR/dist/bin:$PATH # hipify-clang and hipify-perl (build directory) > ``` ```bash # Clone your CUDA project cd ~/projects git clone cd your-cuda-project # Create a backup git checkout -b pre-hipification # Examine what will be converted hipexamine.sh . > conversion-report.txt cat conversion-report.txt # Convert entire project in one command hipconvertinplace.sh . # Verify conversion completeness finduncodep.sh . # Review changes git diff ``` > **✅ Why use helper scripts?** > - Automatically finds all CUDA files recursively > - Handles multiple file types (.cu, .cuh, .h, .cpp, etc.) > - Processes entire directory trees in one command > - Provides conversion statistics and reports > - Less error-prone than manual find commands **Option 3: Using hipify-perl (If available)** If you have `hipify-perl` available (some ROCm installations include it): ```bash # Run hipify-perl on source files find include/ src/ -name "*.cu" -o -name "*.cuh" -o -name "*.hpp" -o -name "*.cpp" | \ xargs /opt/rocm/bin/hipify-perl -inplace # Or use the helper script hipconvertinplace-perl.sh . # Review changes git diff ``` **What HIPIFY tools do:** - Convert `cuda*` function calls to `hip*` equivalents - Change `cudaError_t` to `hipError_t` - Replace `cuda_runtime.h` with `hip_runtime.h` - Convert CUDA macros to HIP macros - `hipify-clang` additionally handles complex template code and preprocessor directives more accurately #### Step 2.2: Manual Code Adjustments Not everything can be automatically converted. Here are common manual changes: **1. Header Includes** ```cpp // Before (CUDA) #ifdef __CUDACC__ #include #endif // After (HIP) #ifdef __CUDACC__ #include #endif ``` **2. Compiler Detection Macros** ```cpp // Before (CUDA-only) #ifdef __NVCC__ // NVIDIA-specific code #endif // After (HIP with backend detection) #ifdef __HIP_PLATFORM_NVIDIA__ // NVIDIA-specific code when compiling HIP for NVIDIA #endif #ifdef __HIP_PLATFORM_AMD__ // AMD-specific code #endif ``` **3. Error Handling** ```cpp // Before (CUDA) #define CUDA_CHECK(call) \ do { \ cudaError_t err = call; \ if (err != cudaSuccess) { \ fprintf(stderr, "CUDA error: %s\n", cudaGetErrorString(err)); \ exit(1); \ } \ } while(0) // After (HIP) #define HIP_CHECK(call) \ do { \ hipError_t err = call; \ if (err != hipSuccess) { \ fprintf(stderr, "HIP error: %s\n", hipGetErrorString(err)); \ exit(1); \ } \ } while(0) ``` **4. Kernel Launch Syntax** Good news: The `<<>>` syntax remains the same! ```cpp // Both CUDA and HIP use the same syntax my_kernel<<>>(args); ``` **5. File Extensions** You can keep `.cu` extensions or rename to `.hip`: ```bash # Optional: rename files find . -name "*.cu" -exec rename 's/\.cu$/.hip/' {} \; find . -name "*.cuh" -exec rename 's/\.cuh$/.hip.hpp/' {} \; ``` --- ### Phase 2: Build System Configuration #### Step 3.1: Create CMakeLists.txt for HIP Here's a complete CMake configuration supporting both NVIDIA and AMD: ```cmake cmake_minimum_required(VERSION 3.24) cmake_policy(SET CMP0135 NEW) # Project configuration project(my-hip-project VERSION 1.0 LANGUAGES CXX HIP DESCRIPTION "GPU-accelerated library with HIP" ) # Options option(GPU "Enable GPU support" ON) option(BUILD_TESTS "Build test suite" ON) # HIP Platform selection set(HIP_PLATFORM "nvidia" CACHE STRING "HIP Platform: nvidia or amd") set_property(CACHE HIP_PLATFORM PROPERTY STRINGS "nvidia" "amd") if(GPU) # Set HIP_PLATFORM environment variable for hipcc set(ENV{HIP_PLATFORM} ${HIP_PLATFORM}) message(STATUS "Building for HIP Platform: ${HIP_PLATFORM}") # Platform-specific paths if(HIP_PLATFORM STREQUAL "nvidia") message(STATUS "Configuring for NVIDIA GPU platform") if(NOT DEFINED ENV{CUDA_PATH}) set(ENV{CUDA_PATH} "/usr/local/cuda") endif() message(STATUS "CUDA_PATH: $ENV{CUDA_PATH}") elseif(HIP_PLATFORM STREQUAL "amd") message(STATUS "Configuring for AMD GPU platform") if(NOT DEFINED ENV{ROCM_PATH}) set(ENV{ROCM_PATH} "/opt/rocm") endif() message(STATUS "ROCM_PATH: $ENV{ROCM_PATH}") else() message(FATAL_ERROR "Invalid HIP_PLATFORM: ${HIP_PLATFORM}") endif() # Enable HIP language enable_language(HIP) # Find HIP package find_package(hip REQUIRED) # Set GPU architectures if(NOT DEFINED CMAKE_HIP_ARCHITECTURES) if(HIP_PLATFORM STREQUAL "nvidia") # NVIDIA GPU architectures (Pascal, Volta, Turing, Ampere, Ada, Hopper) set(CMAKE_HIP_ARCHITECTURES "60;70;75;80;86;89;90" CACHE STRING "NVIDIA GPU architectures") elseif(HIP_PLATFORM STREQUAL "amd") # AMD GPU architectures (MI series and consumer GPUs) set(CMAKE_HIP_ARCHITECTURES "gfx906;gfx908;gfx90a;gfx1030;gfx1100" CACHE STRING "AMD GPU architectures") endif() endif() message(STATUS "Target GPU Architectures: ${CMAKE_HIP_ARCHITECTURES}") # Platform-specific include directories if(HIP_PLATFORM STREQUAL "nvidia") find_path(CUDA_INCLUDE_DIR cuda_runtime.h PATHS /usr/local/cuda*/include ENV CUDA_PATH) find_path(ROCM_INCLUDE_DIR hip/hip_runtime.h PATHS /opt/rocm*/include ENV ROCM_PATH) if(CUDA_INCLUDE_DIR) include_directories(SYSTEM ${CUDA_INCLUDE_DIR}) endif() if(ROCM_INCLUDE_DIR) include_directories(SYSTEM ${ROCM_INCLUDE_DIR}) endif() add_compile_definitions(__HIP_PLATFORM_NVIDIA__) elseif(HIP_PLATFORM STREQUAL "amd") find_path(ROCM_INCLUDE_DIR hip/hip_runtime.h PATHS /opt/rocm*/include ENV ROCM_PATH) if(ROCM_INCLUDE_DIR) include_directories(SYSTEM ${ROCM_INCLUDE_DIR}) endif() add_compile_definitions(__HIP_PLATFORM_AMD__) endif() endif() # Add your library/executable add_library(my_hip_lib INTERFACE) target_include_directories(my_hip_lib INTERFACE $ $ ) # Add tests if enabled if(BUILD_TESTS) enable_testing() add_subdirectory(tests) endif() ``` #### Step 3.2: Create CMakePresets.json CMake presets make it easy to switch between platforms and build configurations: ```json { "version": 6, "cmakeMinimumRequired": { "major": 3, "minor": 24, "patch": 0 }, "configurePresets": [ { "name": "hip-nvidia-debug", "displayName": "HIP NVIDIA Debug", "description": "Debug build for NVIDIA GPUs", "binaryDir": "${sourceDir}/build/hip-nvidia-debug", "cacheVariables": { "CMAKE_BUILD_TYPE": "Debug", "HIP_PLATFORM": "nvidia", "GPU": "ON", "BUILD_TESTS": "ON" } }, { "name": "hip-nvidia-release", "displayName": "HIP NVIDIA Release", "description": "Release build for NVIDIA GPUs", "binaryDir": "${sourceDir}/build/hip-nvidia-release", "cacheVariables": { "CMAKE_BUILD_TYPE": "Release", "HIP_PLATFORM": "nvidia", "GPU": "ON", "BUILD_TESTS": "ON" } }, { "name": "hip-amd-debug", "displayName": "HIP AMD Debug", "description": "Debug build for AMD GPUs", "binaryDir": "${sourceDir}/build/hip-amd-debug", "cacheVariables": { "CMAKE_BUILD_TYPE": "Debug", "HIP_PLATFORM": "amd", "GPU": "ON", "BUILD_TESTS": "ON" } }, { "name": "hip-amd-release", "displayName": "HIP AMD Release", "description": "Release build for AMD GPUs", "binaryDir": "${sourceDir}/build/hip-amd-release", "cacheVariables": { "CMAKE_BUILD_TYPE": "Release", "HIP_PLATFORM": "amd", "GPU": "ON", "BUILD_TESTS": "ON" } } ], "buildPresets": [ { "name": "hip-nvidia-debug", "configurePreset": "hip-nvidia-debug" }, { "name": "hip-nvidia-release", "configurePreset": "hip-nvidia-release" }, { "name": "hip-amd-debug", "configurePreset": "hip-amd-debug" }, { "name": "hip-amd-release", "configurePreset": "hip-amd-release" } ], "testPresets": [ { "name": "hip-nvidia-debug", "configurePreset": "hip-nvidia-debug", "output": {"outputOnFailure": true} }, { "name": "hip-amd-debug", "configurePreset": "hip-amd-debug", "output": {"outputOnFailure": true} } ] } ``` #### Step 3.3: Building Your Project **For NVIDIA GPUs:** ```bash # Configure cmake --preset hip-nvidia-release # Build cmake --build --preset hip-nvidia-release # Run tests ctest --preset hip-nvidia-release --output-on-failure ``` **For AMD GPUs:** ```bash # Configure cmake --preset hip-amd-release # Build cmake --build --preset hip-amd-release # Run tests ctest --preset hip-amd-release --output-on-failure ``` --- ### Phase 3: Testing and Validation #### Step 4.1: Create Basic Test Create `tests/basic_test.cpp`: ```cpp #include #include #include __global__ void vector_add(float* a, float* b, float* c, int n) { int idx = blockIdx.x * blockDim.x + threadIdx.x; if (idx < n) { c[idx] = a[idx] + b[idx]; } } int main() { const int n = 1000; const size_t bytes = n * sizeof(float); // Allocate host memory float *h_a = new float[n]; float *h_b = new float[n]; float *h_c = new float[n]; // Initialize data for (int i = 0; i < n; i++) { h_a[i] = i; h_b[i] = i * 2; } // Allocate device memory float *d_a, *d_b, *d_c; hipMalloc(&d_a, bytes); hipMalloc(&d_b, bytes); hipMalloc(&d_c, bytes); // Copy data to device hipMemcpy(d_a, h_a, bytes, hipMemcpyHostToDevice); hipMemcpy(d_b, h_b, bytes, hipMemcpyHostToDevice); // Launch kernel int blockSize = 256; int gridSize = (n + blockSize - 1) / blockSize; vector_add<<>>(d_a, d_b, d_c, n); // Copy result back hipMemcpy(h_c, d_c, bytes, hipMemcpyDeviceToHost); // Verify results for (int i = 0; i < n; i++) { assert(h_c[i] == h_a[i] + h_b[i]); } std::cout << "Test passed!" << std::endl; // Cleanup delete[] h_a; delete[] h_b; delete[] h_c; hipFree(d_a); hipFree(d_b); hipFree(d_c); return 0; } ``` #### Step 4.2: Add Test to CMake Create `tests/CMakeLists.txt`: ```cmake # Find Google Test (optional but recommended) find_package(GTest) if(GTest_FOUND) include(GoogleTest) enable_testing() endif() # Basic HIP test (no external dependencies) add_executable(basic_test basic_test.cpp) set_target_properties(basic_test PROPERTIES LINKER_LANGUAGE HIP HIP_ARCHITECTURES ${CMAKE_HIP_ARCHITECTURES} ) add_test(NAME BasicTest COMMAND basic_test) # If using Google Test if(GTest_FOUND) add_executable(advanced_test advanced_test.cpp) set_target_properties(advanced_test PROPERTIES LINKER_LANGUAGE HIP HIP_ARCHITECTURES ${CMAKE_HIP_ARCHITECTURES} ) target_link_libraries(advanced_test GTest::GTest GTest::Main) gtest_discover_tests(advanced_test) endif() ``` #### Step 4.3: Validation Checklist ✅ **Build Verification** - [ ] Project builds successfully for NVIDIA backend - [ ] Project builds successfully for AMD backend - [ ] No compiler warnings or errors - [ ] All tests compile ✅ **Runtime Verification** - [ ] Basic smoke tests pass on NVIDIA hardware - [ ] Basic smoke tests pass on AMD hardware - [ ] Memory allocations/deallocations work correctly - [ ] Kernel launches execute successfully - [ ] Error handling works properly ✅ **Performance Verification** - [ ] Performance comparable to original CUDA code on NVIDIA - [ ] Performance acceptable on AMD hardware - [ ] No memory leaks detected ✅ **Code Quality** - [ ] Code follows HIP best practices - [ ] Documentation updated - [ ] Examples updated with HIP code --- ## Case Study: cuda-battery to hip-battery Let's examine how the cuda-battery project (a standard library for CUDA programming) was successfully hipified using Approach 1 (Unified Codebase). > **📦 Source Code:** [cuda-battery HIP implementation](https://github.com/icemc/cuda-battery/tree/hip) on GitHub ### Project Overview **cuda-battery** is a C++ library providing: - GPU-compatible data structures (vector, string, bitset, etc.) - Custom allocators for different memory types - Smart pointers (unique_ptr, shared_ptr) for device code - Utility functions and macros for CUDA development ### Migration Strategy The project chose the **Unified Codebase approach** to: 1. Maintain a single source for all GPU vendors 2. Simplify testing and development 3. Future-proof against new GPU vendors ### Key Changes Made #### 1. Header Updates **Before:** ```cpp #ifdef __CUDACC__ #include #endif ``` **After:** ```cpp #ifdef __CUDACC__ #include #endif ``` #### 2. Error Handling Functions **Before:** ```cpp CUDA NI inline void gpuAssert(cudaError_t code, const char *file, int line) { if (code != cudaSuccess) { printf("%s:%d CUDA error %s\n", file, line, cudaGetErrorString(code)); } } ``` **After:** ```cpp CUDA NI inline void gpuAssert(hipError_t code, const char *file, int line) { if (code != hipSuccess) { printf("%s:%d HIP error %s\n", file, line, hipGetErrorString(code)); } } ``` #### 3. CMake Build System The project created a flexible CMake configuration: ```cmake # Support both NVIDIA and AMD via HIP_PLATFORM set(HIP_PLATFORM "nvidia" CACHE STRING "HIP Platform: nvidia or amd") # Configure based on platform if(HIP_PLATFORM STREQUAL "nvidia") set(CMAKE_HIP_ARCHITECTURES "60;70;75;80;86;89;90") elseif(HIP_PLATFORM STREQUAL "amd") set(CMAKE_HIP_ARCHITECTURES "gfx906;gfx908;gfx90a;gfx1030") endif() ``` #### 4. Testing Strategy The project maintained comprehensive tests that run on both platforms: ```bash # Test on NVIDIA cmake --preset hip-nvidia-debug cmake --build --preset hip-nvidia-debug ctest --preset hip-nvidia-debug # Test on AMD (if hardware available) cmake --preset hip-amd-debug cmake --build --preset hip-amd-debug ctest --preset hip-amd-debug ``` ### Results ✅ **Single codebase** maintained across GPU vendors ✅ **Zero performance regression** on NVIDIA hardware ✅ **Full AMD GPU support** added ✅ **Simplified CI/CD pipeline** (one codebase to test) ✅ **Easy for contributors** to work with ⚠️ **Minor tradeoff**: NVIDIA users need ROCm headers installed (but not ROCm runtime) ### Lessons Learned 1. **Automated conversion works well**: `hipify-perl` or `hipify-clang` handled 95% of the conversion. Using helper scripts like `hipconvertinplace.sh` makes bulk conversion even easier 2. **Test early and often**: Run tests on both platforms frequently 3. **Keep it simple**: Don't prematurely optimize for platform-specific features 4. **Documentation matters**: Update all docs to reflect HIP usage 5. **CMake presets**: Make it easy for users to build for their platform --- ## Approach 2: Separate Codebases **Strategy**: Maintain separate CUDA and HIP implementations, selected at compile-time. ### Advantages ✅ **Minimal Dependencies for End-Users**: Users only install what they need for their platform (CUDA for NVIDIA, ROCm for AMD) ✅ **Platform-Specific Optimizations**: Can use vendor-specific features freely ✅ **Independent Evolution**: CUDA and HIP code can evolve separately ✅ **Gradual Migration**: Can port incrementally > **📝 Note for Developers**: While end-users benefit from minimal dependencies, developers maintaining both code paths will need all dependencies (CUDA + ROCm + HIPIFY) installed to develop and test both implementations. ### Disadvantages ❌ **Double Maintenance**: Bug fixes and features need implementation in both codebases ❌ **Code Drift**: Implementations can diverge over time ❌ **Complex Build System**: Requires sophisticated conditional compilation ❌ **Testing Overhead**: Must test both implementations separately ❌ **Higher Development Cost**: More person-hours required ### When to Use This Approach - You have platform-specific optimizations that can't be unified - Dependencies are a major concern for your users - You have dedicated resources for multi-platform maintenance - Your CUDA code uses many NVIDIA-specific features - You're working with legacy code that's difficult to convert completely --- ## Approach 2: Step-by-Step Implementation Guide ### Required Dependencies Before implementing this approach, ensure you have the necessary tools from the [Environment Setup and Dependencies](#environment-setup-and-dependencies) section: **For CUDA Code Path:** - ✅ **[Step 1: Install Build Essentials](#step-1-install-build-essentials)** - Required for compilation - ✅ **[Step 2: Install CMake](#step-2-install-cmake-build-system)** - Required for building - ✅ **[Step 3: Install CUDA Toolkit](#step-3-install-cuda-toolkit-nvidia-platform-support)** - Required for CUDA implementation **For HIP Code Path:** - ✅ **[Step 1: Install Build Essentials](#step-1-install-build-essentials)** - Required for compilation - ✅ **[Step 2: Install CMake](#step-2-install-cmake-build-system)** - Required for building - ✅ **[Step 4: Install HIP and ROCm](#step-4-install-hip-and-rocm)** - Required for HIP runtime - ✅ **[Step 5: Build HIPIFY](#step-5-build-hipify-code-conversion-tool)** - Needed if converting existing CUDA code to HIP > **💡 Note:** With this approach, users only need to install dependencies for their target platform(s). Developers maintaining both code paths will need all dependencies. --- ### Phase 1: Project Structure Organize your code with clear separation: ``` my-project/ ├── CMakeLists.txt ├── include/ │ ├── my_library.hpp # Common interface │ └── my_library/ │ ├── common.hpp # Shared definitions │ ├── cuda/ # CUDA-specific headers │ │ ├── kernels.cuh │ │ └── internal.cuh │ └── hip/ # HIP-specific headers │ ├── kernels.hip.hpp │ └── internal.hip.hpp ├── src/ │ ├── common.cpp # Platform-independent code │ ├── cuda/ # CUDA implementation │ │ ├── kernels.cu │ │ └── api.cpp │ └── hip/ # HIP implementation │ ├── kernels.hip │ └── api.cpp └── tests/ ├── cuda/ # CUDA-specific tests └── hip/ # HIP-specific tests ``` ### Phase 2: CMake Configuration ```cmake cmake_minimum_required(VERSION 3.24) project(my_project LANGUAGES CXX) # Options option(USE_GPU "Enable GPU support" ON) option(USE_CUDA "Use CUDA backend" OFF) option(USE_HIP "Use HIP backend" OFF) # Ensure only one backend is selected if(USE_CUDA AND USE_HIP) message(FATAL_ERROR "Cannot enable both CUDA and HIP backends") endif() if(USE_GPU) if(USE_CUDA) # CUDA-specific configuration enable_language(CUDA) find_package(CUDAToolkit REQUIRED) set(GPU_SOURCES src/cuda/kernels.cu src/cuda/api.cpp ) set(GPU_INCLUDE_DIRS ${CMAKE_CURRENT_SOURCE_DIR}/include/my_library/cuda ${CUDAToolkit_INCLUDE_DIRS} ) add_compile_definitions(USE_CUDA_BACKEND) elseif(USE_HIP) # HIP-specific configuration set(HIP_PLATFORM "nvidia" CACHE STRING "HIP Platform") set(ENV{HIP_PLATFORM} ${HIP_PLATFORM}) enable_language(HIP) find_package(hip REQUIRED) set(GPU_SOURCES src/hip/kernels.hip src/hip/api.cpp ) set(GPU_INCLUDE_DIRS ${CMAKE_CURRENT_SOURCE_DIR}/include/my_library/hip ) add_compile_definitions(USE_HIP_BACKEND) if(HIP_PLATFORM STREQUAL "nvidia") add_compile_definitions(__HIP_PLATFORM_NVIDIA__) else() add_compile_definitions(__HIP_PLATFORM_AMD__) endif() else() message(FATAL_ERROR "USE_GPU=ON requires either USE_CUDA or USE_HIP") endif() else() # CPU-only build set(GPU_SOURCES "") set(GPU_INCLUDE_DIRS "") add_compile_definitions(CPU_ONLY) endif() # Create library add_library(my_library src/common.cpp ${GPU_SOURCES} ) target_include_directories(my_library PUBLIC $ $ $ ) ``` ### Phase 3: Code Organization #### Common Interface Header (`include/my_library.hpp`) ```cpp #pragma once #include namespace my_library { // Forward declarations class DeviceArray; // Common API (platform-agnostic) class GpuContext { public: GpuContext(); ~GpuContext(); DeviceArray allocate(size_t bytes); void copy_to_device(void* dst, const void* src, size_t bytes); void copy_to_host(void* dst, const void* src, size_t bytes); void synchronize(); private: struct Impl; Impl* impl_; }; } // namespace my_library ``` #### CUDA Implementation (`src/cuda/api.cpp`) ```cpp #ifdef USE_CUDA_BACKEND #include "my_library.hpp" #include namespace my_library { struct GpuContext::Impl { cudaStream_t stream; }; GpuContext::GpuContext() : impl_(new Impl) { cudaStreamCreate(&impl_->stream); } GpuContext::~GpuContext() { cudaStreamDestroy(impl_->stream); delete impl_; } void GpuContext::copy_to_device(void* dst, const void* src, size_t bytes) { cudaMemcpyAsync(dst, src, bytes, cudaMemcpyHostToDevice, impl_->stream); } void GpuContext::synchronize() { cudaStreamSynchronize(impl_->stream); } } // namespace my_library #endif // USE_CUDA_BACKEND ``` #### HIP Implementation (`src/hip/api.cpp`) ```cpp #ifdef USE_HIP_BACKEND #include "my_library.hpp" #include namespace my_library { struct GpuContext::Impl { hipStream_t stream; }; GpuContext::GpuContext() : impl_(new Impl) { hipStreamCreate(&impl_->stream); } GpuContext::~GpuContext() { hipStreamDestroy(impl_->stream); delete impl_; } void GpuContext::copy_to_device(void* dst, const void* src, size_t bytes) { hipMemcpyAsync(dst, src, bytes, hipMemcpyHostToDevice, impl_->stream); } void GpuContext::synchronize() { hipStreamSynchronize(impl_->stream); } } // namespace my_library #endif // USE_HIP_BACKEND ``` ### Phase 4: Build Instructions **For CUDA:** ```bash cmake -S . -B build/cuda \ -DUSE_GPU=ON \ -DUSE_CUDA=ON \ -DCMAKE_BUILD_TYPE=Release cmake --build build/cuda ``` **For HIP (NVIDIA):** ```bash cmake -S . -B build/hip-nvidia \ -DUSE_GPU=ON \ -DUSE_HIP=ON \ -DHIP_PLATFORM=nvidia \ -DCMAKE_BUILD_TYPE=Release cmake --build build/hip-nvidia ``` **For HIP (AMD):** ```bash cmake -S . -B build/hip-amd \ -DUSE_GPU=ON \ -DUSE_HIP=ON \ -DHIP_PLATFORM=amd \ -DCMAKE_BUILD_TYPE=Release cmake --build build/hip-amd ``` --- ## Decision Matrix: Which Approach to Choose? ### Choose Unified Codebase (Approach 1) When: - ✅ Your project is a library/framework meant for wide distribution - ✅ Development team is small or has limited GPU expertise - ✅ Code doesn't heavily rely on NVIDIA-specific features (e.g., Tensor Cores, WMMA) - ✅ End-users can tolerate installing ROCm headers on NVIDIA systems (lightweight, no runtime needed) - ✅ You value simplicity and maintainability over all else - ✅ You're starting fresh or have a moderate codebase (<50k LOC) ### Choose Separate Codebases (Approach 2) When: - ✅ You have performance-critical code with vendor-specific optimizations - ✅ Large existing CUDA codebase with tons of legacy code - ✅ Dedicated team for each platform (developers still need both CUDA and ROCm) - ✅ End-users are very sensitive to any cross-platform dependencies - ✅ You use advanced CUDA features not easily portable to HIP - ✅ You need maximum performance on each platform ### Hybrid Approach You can also combine both: - Use unified HIP codebase for 90% of code - Keep separate implementations for performance-critical kernels - Use conditional compilation to select the optimal implementation ```cpp #if defined(USE_HIP_BACKEND) && defined(__HIP_PLATFORM_AMD__) // AMD-optimized implementation amd_optimized_kernel<<>>(args); #elif defined(USE_CUDA_BACKEND) || defined(__HIP_PLATFORM_NVIDIA__) // NVIDIA-optimized implementation nvidia_optimized_kernel<<>>(args); #else // Generic fallback generic_kernel<<>>(args); #endif ``` --- ## Conclusion ### Summary Hipification opens your CUDA code to a wider GPU ecosystem without massive rewrites. The two approaches presented here each have their place: | Aspect | Unified HIP | Separate CUDA+HIP | |--------|-------------|-------------------| | **Maintenance** | Single codebase, easier | Dual codebases, harder | | **Performance** | Same as native | Potential for better optimization | | **End-User Dependencies** | ROCm headers on NVIDIA, full ROCm on AMD | Only platform-specific | | **Developer Dependencies** | CUDA + ROCm (testing both) | CUDA + ROCm (testing both) | | **Best for** | Most projects | Performance-critical, vendor-specific code | ### Recommendations For **most projects**, we recommend **Approach 1 (Unified HIP Codebase)** because: - Lower maintenance burden - Simpler testing (one codebase to validate) - Easier collaboration - Future-proof portability - HIP performance equals CUDA on NVIDIA hardware - ROCm header dependency on NVIDIA is lightweight (no runtime needed) For **specialized scenarios** (high-performance computing with vendor-specific optimizations), **Approach 2 (Separate Codebases)** may be justified. ### Next Steps 1. **Audit your codebase**: Identify NVIDIA-specific features you use 2. **Choose your approach**: Based on the decision matrix above 3. **Set up environment**: Install ROCm and/or CUDA as needed 4. **Start small**: Convert a single module first 5. **Test thoroughly**: Verify correctness on both platforms 6. **Iterate**: Expand converted code, refine build system 7. **Document**: Update README, build instructions, examples ### Resources #### Official Documentation - [HIP Programming Guide](https://rocm.docs.amd.com/projects/HIP/en/latest/how-to/programming_manual.html) - [HIP API Reference](https://rocm.docs.amd.com/projects/HIP/en/latest/reference/index.html) - [HIP Porting Guide](https://rocm.docs.amd.com/en/latest/how-to/hip-porting-guide.html) - [ROCm Documentation Hub](https://rocm.docs.amd.com/) - [CUDA Documentation](https://docs.nvidia.com/cuda/) #### Tools - [HIPIFY Tools](https://github.com/ROCm/HIPIFY) - Automatic CUDA to HIP converter (perl and clang) - [HIP Examples](https://github.com/ROCm/HIP-Examples) - Sample codes and tutorials - [ROCm CMake Utilities](https://github.com/ROCm/rocm-cmake) - CMake modules for ROCm projects - [GPU Open Tools](https://gpuopen.com/tools/) - AMD's suite of GPU development tools #### Community - [ROCm GitHub Organization](https://github.com/ROCm) - [HIP GitHub Repository](https://github.com/ROCm/HIP) - [AMD GPU Open](https://gpuopen.com/) - [ROCm Discussion Forums](https://github.com/ROCm/ROCm/discussions) - [Stack Overflow - ROCm Tag](https://stackoverflow.com/questions/tagged/rocm) - [Stack Overflow - HIP Tag](https://stackoverflow.com/questions/tagged/hip) --- ## Appendix A: Common Conversion Patterns ### API Conversions | CUDA | HIP | Notes | |------|-----|-------| | `cudaMalloc` | `hipMalloc` | Identical signature | | `cudaFree` | `hipFree` | Identical signature | | `cudaMemcpy` | `hipMemcpy` | Identical signature | | `cudaMemcpyAsync` | `hipMemcpyAsync` | Identical signature | | `cudaDeviceSynchronize` | `hipDeviceSynchronize` | Identical signature | | `cudaGetLastError` | `hipGetLastError` | Identical signature | | `cudaSuccess` | `hipSuccess` | Identical usage | | `cudaError_t` | `hipError_t` | Type replacement | | `cudaStream_t` | `hipStream_t` | Type replacement | | `cudaEvent_t` | `hipEvent_t` | Type replacement | ### Compiler Macros | CUDA | HIP (NVIDIA) | HIP (AMD) | |------|--------------|-----------| | `__CUDACC__` | `__CUDACC__` | `__CUDACC__` (same!) | | `__NVCC__` | `__HIP_PLATFORM_NVIDIA__` | Not defined | | N/A | Not defined | `__HIP_PLATFORM_AMD__` | | `__CUDA_ARCH__` | `__CUDA_ARCH__` | `__gfx*__` | ### Memory Types | Memory Type | CUDA | HIP | |-------------|------|-----| | Global | `cudaMalloc` | `hipMalloc` | | Managed | `cudaMallocManaged` | `hipMallocManaged` | | Pinned | `cudaMallocHost` | `hipHostMalloc` | | Shared | `__shared__` | `__shared__` (same!) | | Constant | `__constant__` | `__constant__` (same!) | --- ## Appendix B: Troubleshooting ### Build Errors **Problem**: `fatal error: hip/hip_runtime.h: No such file or directory` **Solution**: ROCm not installed or not in include path ```bash sudo apt install rocm-hip-sdk # Add to CMakeLists.txt: include_directories(SYSTEM /opt/rocm/include) ``` **Problem**: `CMake Error: Could not find HIP compiler` **Solution**: HIP language not properly enabled ```bash # Make sure CMake can find hipcc export PATH=/opt/rocm/bin:$PATH # In CMakeLists.txt: enable_language(HIP) ``` **Problem**: `undefined reference to '__hipRegisterFatBinary'` **Solution**: Linking issue, make sure to link hip libraries ```cmake target_link_libraries(my_target hip::host) ``` ### Runtime Errors **Problem**: Tests pass on NVIDIA but fail on AMD **Solution**: Check for: - Architecture-specific assumptions (warp size, shared memory size) - Undefined behavior exposed by different compiler - Floating-point precision differences **Problem**: Performance regression after hipification **Solution**: - Verify optimization flags are enabled (`-O3`, `--use_fast_math`) - Check architecture targets in `CMAKE_HIP_ARCHITECTURES` - Profile with appropriate tools (`nvprof` or `rocprof`) ---