From 42e046594e91c1283ddd647d2165ba8b662b3999 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Mon, 8 Jun 2026 21:45:12 +0000 Subject: [PATCH 1/2] [ROCm] Add AMD GPU (HIP) support to the RPUCuda backend This adds an optional USE_HIP build of the RPUCuda CUDA backend so the analog tile simulator runs on AMD GPUs through ROCm, alongside the existing USE_CUDA path. Every HIP change is gated behind USE_HIP, so we have made every effort to leave the CUDA build compiling the same sources it did before. The build stays on the project's standalone scikit-build/CMake CUDA toolchain (PyTorch is used only for headers and linkage), so no torch cpp_extension hipify is involved. Review order: start with src/rpucuda/cuda/cuda_to_hip.h, the single HIP-aware header. It is included once from cuda_util.h, ahead of the CUDA runtime / cuBLAS / cuRAND includes and guarded by USE_HIP, and aliases the cuda*/cublas*/ curand* spellings the backend uses to their hip*/hipBLAS/hipRAND equivalents so the CUDA-spelled device sources compile unchanged. cuBLAS maps to hipBLAS (GEMM/GEMV/GER/scal/copy/nrm2, handle, host and device pointer modes, op and status enums) and cuRAND to hipRAND (host generator plus device per-thread state). In CUDA, curandState, curandState_t and curandStateXORWOW are one XORWOW type; hipRAND declares hiprandState and hiprandStateXORWOW_t as distinct structs, so all three spellings are mapped to one hipRAND type to keep the backend's template instantiations consistent (mapping them apart drops a CudaArray instantiation and breaks the module load). The one semantically delicate piece is bit_line_maker.cu. Its stochastic pulse train is a warp-size-coupled serialized bit format: a 32-lane ballot is packed into 32-bit words and the pulsed-weight-updater kernels read those words back by bit position, so the on-device word layout must stay independent of the physical wavefront width. A 64-wide CDNA wavefront holds two independent 32-lane logical warps, and a native 64-lane ballot would merge them into one 64-bit value and corrupt the format. The compat header therefore redefines the ballot/shuffle sync intrinsics as width-32 logical-warp operations: the ballot takes the full wavefront ballot and selects only the calling lane's own 32-lane subgroup as a 32-bit word (shift 0 on wave32; lanes 32-63 select the high word on wave64), and the shuffles use width 32. The word width and the consumer stay unchanged, so a wave32 and a wave64 device emit the same 32-bit-per-32-lane stream the CUDA build produces. CMake gains an option(USE_HIP) and a dependencies_hip.cmake mirroring the CUDA path: it enables the HIP language, marks the .cu sources LANGUAGE HIP, links hipBLAS/hipRAND/hipCUB, takes the architecture list from CMAKE_HIP_ARCHITECTURES (defaulting to gfx90a only when unset, never hardcoded), keeps RPU_USE_CUDA defined so the device tile code and the CudaAnalogTile binding stay built, and forces interprocedural optimization off (HIP device linking does not finalize LTO, which would yield an empty module init). rpu_cub.h selects hipCUB and its hipcub namespace. The pybind translation units are compiled with the HIP toolchain because they inherit the GPU library's offload-arch usage requirement, and rpu_base_tiles_cuda.cpp uses the lightweight c10 HIP stream header instead of the full ATen CUDA context header, which would drag in hipSOLVER/hipSPARSE headers not part of this build. RPU_CXX_STANDARD is a new cache variable (default 17) so the standard can track newer PyTorch headers that require C++20; that also exposed and fixed an ill-formed template-id constructor name in rpu_linearstep_device.h. The pybind layer hand-rolls an at::cuda::getCurrentCUDAStream shim in rpu_base_tiles_cuda.cpp because this build defines its own USE_HIP and never runs PyTorch source-hipify on that translation unit. Which underlying c10 symbol the shim must call depends on the torch hipify generation, not the ROCm version: hipify v2 keeps c10::cuda::getCurrentCUDAStream as the public stream API (c10::hip::getCurrentHIPStream stays guarded by USE_ROCM, which this build does not define), while hipify v1 removes it so the hip-spelled symbol must be used. CMake detects the hipify version from the build's own torch and defines TORCH_HIPIFY_V2 when it is >= 2.0.0; the shim selects its branch from that define, and detection failure leaves it undefined, matching the safe v1 default. Building with amdclang++ on Windows needs two further fixes: guard the MSVC-only /O2 flag with if(MSVC) (Clang-on-Windows is not MSVC) while keeping the C++ standard settings for all Windows builds, and change the forward declaration of PulsedUpdateMetaParameter from class to struct to match its definition (struct and class mangle identically on ELF but differently on the MSVC ABI used by amdclang++, which otherwise leaves virtual-method symbols unresolved at link). The ROCm/HIP build is documented in the developer install guide alongside the CUDA build. This work was authored with the assistance of Claude, an AI assistant. Test Plan: Built and validated on AMD GPUs across CDNA and RDNA, exercising the analog tiles on device against the CPU tiles (tolerance-based comparisons). Linux, AMD Instinct MI250X (gfx90a, CDNA2, wave64), ROCm 7.2.1: ``` USE_HIP=ON USE_CUDA=0 python setup.py build_ext -j16 --inplace \ -DUSE_HIP=ON -DUSE_CUDA=OFF -DRPU_CXX_STANDARD=20 \ -DCMAKE_HIP_ARCHITECTURES=gfx90a \ -DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ \ -DCMAKE_PREFIX_PATH="$TORCH_CMAKE;/opt/rocm" \ -DRPU_BLAS=OpenBLAS -DRPU_USE_TORCH_BUFFERS=OFF -DCMAKE_BUILD_TYPE=Release HIP_VISIBLE_DEVICES=0 PYTHONPATH=src python -m pytest -q \ tests/test_specific_tiles.py tests/test_simulator_tiles.py \ tests/test_bindings_tiles.py tests/test_torch_tiles.py \ tests/test_inference_tiles.py tests/test_layers_linear.py \ tests/test_layers_convolution.py ``` test_specific_tiles (the bit-line-maker plus pulsed-update path) passes 18/18; the tile/binding/torch/inference/layer suites pass (one stochastic weight-programming test is order-sensitive and passes in isolation). A fat binary for "gfx90a;gfx1100" builds and emits both code objects for bit_line_maker.cu, confirming the ballot rewrite is wave-width agnostic. The same suites also pass on a Linux gfx1100 (RDNA3, wave32) device. Windows, AMD Radeon RX 9070 XT (gfx1201, RDNA4, wave32), ROCm 7.14: ``` cmake -DUSE_HIP=ON -DUSE_CUDA=OFF -DRPU_CXX_STANDARD=20 \ -DCMAKE_HIP_ARCHITECTURES=gfx1201 -DRPU_BLAS=OpenBLAS -DBUILD_TEST=OFF ninja rpu_base ``` test_specific_tiles 18/18; the tile/binding/torch/inference/layer suites pass (over 900 tests), with one order-sensitive stochastic test that passes in isolation. Signed-off-by: Jeff Daily --- CMakeLists.txt | 83 ++++++--- README.md | 2 +- cmake/dependencies_hip.cmake | 47 +++++ docs/source/developer_install.rst | 22 +++ src/aihwkit/simulator/CMakeLists.txt | 23 ++- .../rpu_base_src/rpu_base_tiles_cuda.cpp | 26 +++ src/rpucuda/cuda/cuda_to_hip.h | 174 ++++++++++++++++++ src/rpucuda/cuda/cuda_util.h | 4 + src/rpucuda/cuda/pwu_kernel_parameter_base.h | 2 +- src/rpucuda/cuda/rpu_cub.h | 13 ++ src/rpucuda/rpu_linearstep_device.h | 2 +- 11 files changed, 367 insertions(+), 31 deletions(-) create mode 100644 cmake/dependencies_hip.cmake create mode 100644 src/rpucuda/cuda/cuda_to_hip.h diff --git a/CMakeLists.txt b/CMakeLists.txt index 97254ed0d..6af5eaf0e 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -9,6 +9,7 @@ project(aihwkit C CXX) option(BUILD_TEST "Build C++ test binaries" OFF) option(BUILD_EXTENSION "Build additional C++ tools" OFF) option(USE_CUDA "Build with CUDA support" $ENV{USE_CUDA}) +option(USE_HIP "Build with ROCm/HIP support" $ENV{USE_HIP}) # experimental precision flags option(RPU_USE_FP16 "EXPERIMENTAL: Build FP16 support (only available with CUDA)" OFF) @@ -22,14 +23,24 @@ option(RPU_USE_FASTRAND "Use fastrand" OFF) option(RPU_USE_TORCH_BUFFERS "Use torch buffers for RPUCuda" ON) +# The C++ standard tracks the linked PyTorch headers; newer Torch releases +# require C++20. Default stays 17 to match older Torch. +set(RPU_CXX_STANDARD "17" CACHE STRING "C++ standard used for all RPU targets") + set(RPU_BLAS "OpenBLAS" CACHE STRING "BLAS backend of choice (OpenBLAS, MKL)") set(RPU_CUDA_ARCHITECTURES "75;80;89" CACHE STRING "Target CUDA architectures") +if(NOT DEFINED CMAKE_HIP_ARCHITECTURES) + set(CMAKE_HIP_ARCHITECTURES "gfx90a" CACHE STRING "Target HIP (ROCm) architectures") +endif() # Internal variables. set(CUDA_TARGET_PROPERTIES POSITION_INDEPENDENT_CODE ON CUDA_RESOLVE_DEVICE_SYMBOLS ON CUDA_SEPARABLE_COMPILATION ON - CXX_STANDARD 17) + CXX_STANDARD ${RPU_CXX_STANDARD}) +set(HIP_TARGET_PROPERTIES POSITION_INDEPENDENT_CODE ON + INTERPROCEDURAL_OPTIMIZATION OFF + CXX_STANDARD ${RPU_CXX_STANDARD}) # Append the virtualenv library path to cmake. if(DEFINED ENV{VIRTUAL_ENV}) @@ -41,13 +52,16 @@ endif() # Check for dependencies. include(cmake/dependencies.cmake) include(cmake/dependencies_cuda.cmake) +include(cmake/dependencies_hip.cmake) include(cmake/dependencies_test.cmake) # Set compilation flags. if(WIN32) - set(CMAKE_CXX_FLAGS_RELEASE "${CMAKE_CXX_FLAGS_RELEASE} /O2") - set(CMAKE_CXX_STANDARD 17) + set(CMAKE_CXX_STANDARD ${RPU_CXX_STANDARD}) set(CMAKE_CXX_STANDARD_REQUIRED ON) +endif() +if(MSVC) + set(CMAKE_CXX_FLAGS_RELEASE "${CMAKE_CXX_FLAGS_RELEASE} /O2") else() set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wall -Wno-narrowing -Wno-strict-overflow") set(CMAKE_CXX_FLAGS_RELEASE "${CMAKE_CXX_FLAGS_RELEASE} -O3 -ftree-vectorize") @@ -68,7 +82,7 @@ if(WIN32) target_link_libraries(RPU_CPU c10.lib torch_cpu.lib) endif() -set_target_properties(RPU_CPU PROPERTIES CXX_STANDARD 17 +set_target_properties(RPU_CPU PROPERTIES CXX_STANDARD ${RPU_CXX_STANDARD} POSITION_INDEPENDENT_CODE ON) add_compile_definitions(RPU_USE_WITH_TORCH) @@ -78,9 +92,12 @@ if (RPU_USE_DOUBLE) endif(RPU_USE_DOUBLE) -if(USE_CUDA) +if(USE_CUDA OR USE_HIP) add_subdirectory(src/rpucuda/cuda) include_directories(SYSTEM src/rpucuda/cuda) + if (USE_HIP) + set_source_files_properties(${RPU_GPU_SRCS} PROPERTIES LANGUAGE HIP) + endif() add_library(RPU_GPU ${RPU_GPU_SRCS}) if (RPU_USE_FP16) @@ -99,13 +116,19 @@ if(USE_CUDA) message(STATUS "Use FP16 parameters for CUDA (for all RPU number types).") endif(RPU_PARAM_FP16) - target_link_libraries(RPU_GPU RPU_CPU cublas curand ${RPU_DEPENDENCY_LIBS}) - if(WIN32) - target_link_libraries(RPU_GPU c10_cuda.lib torch_cuda.lib) - endif(WIN32) + if (USE_HIP) + target_link_libraries(RPU_GPU RPU_CPU roc::hipblas hip::hiprand hip::hipcub ${RPU_DEPENDENCY_LIBS}) + set_target_properties(RPU_GPU PROPERTIES ${HIP_TARGET_PROPERTIES}) + set_property(TARGET RPU_GPU PROPERTY HIP_ARCHITECTURES ${CMAKE_HIP_ARCHITECTURES}) + else() + target_link_libraries(RPU_GPU RPU_CPU cublas curand ${RPU_DEPENDENCY_LIBS}) + if(WIN32) + target_link_libraries(RPU_GPU c10_cuda.lib torch_cuda.lib) + endif(WIN32) - set_target_properties(RPU_GPU PROPERTIES ${CUDA_TARGET_PROPERTIES}) - set_property(TARGET RPU_GPU PROPERTY CUDA_ARCHITECTURES ${RPU_CUDA_ARCHITECTURES}) + set_target_properties(RPU_GPU PROPERTIES ${CUDA_TARGET_PROPERTIES}) + set_property(TARGET RPU_GPU PROPERTY CUDA_ARCHITECTURES ${RPU_CUDA_ARCHITECTURES}) + endif() if (RPU_USE_TORCH_BUFFERS) if (BUILD_TEST) @@ -114,18 +137,20 @@ if(USE_CUDA) set(RPU_USE_TORCH_BUFFERS OFF) else (BUILD_TEST) add_compile_definitions(RPU_TORCH_CUDA_BUFFERS) - set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --expt-relaxed-constexpr") - set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Xcudafe --diag_suppress=186") + if (NOT USE_HIP) + set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --expt-relaxed-constexpr") + set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Xcudafe --diag_suppress=186") + endif() endif(BUILD_TEST) endif(RPU_USE_TORCH_BUFFERS) - if(${CUDAToolkit_VERSION_MAJOR} LESS 11) + if(USE_CUDA AND CUDAToolkit_VERSION_MAJOR LESS 11) # The "cub" target only exists if cub was downloaded during build. if(TARGET cub) add_dependencies(RPU_GPU cub) endif() endif() -endif(USE_CUDA) +endif() # Add aihwkit targets. add_subdirectory(src/aihwkit/simulator) @@ -136,7 +161,7 @@ if (BUILD_EXTENSION) add_library(AIHWKIT_EXTENSION_OPS ${AIHWKIT_EXTENSION_OPS_CPU_SRCS}) - set_target_properties(AIHWKIT_EXTENSION_OPS PROPERTIES CXX_STANDARD 17 + set_target_properties(AIHWKIT_EXTENSION_OPS PROPERTIES CXX_STANDARD ${RPU_CXX_STANDARD} POSITION_INDEPENDENT_CODE ON) target_link_libraries(AIHWKIT_EXTENSION_OPS @@ -178,7 +203,7 @@ if (BUILD_EXTENSION) target_link_libraries(${extension_module_name} PRIVATE torch_python) target_include_directories(${extension_module_name} PRIVATE src/aihwkit/extension/extension_src) target_include_directories(${extension_module_name} PRIVATE src/aihwkit/extension/extension_src/ops) - set_target_properties(${extension_module_name} PROPERTIES CXX_STANDARD 17 + set_target_properties(${extension_module_name} PROPERTIES CXX_STANDARD ${RPU_CXX_STANDARD} POSITION_INDEPENDENT_CODE ON) if (USE_CUDA) @@ -211,7 +236,7 @@ if(BUILD_TEST) add_executable(${test_name} ${test_src}) target_link_libraries(${test_name} gtest gmock) target_link_libraries(${test_name} torch_python c10 torch_cpu) - set_target_properties(${test_name} PROPERTIES CXX_STANDARD 17 + set_target_properties(${test_name} PROPERTIES CXX_STANDARD ${RPU_CXX_STANDARD} POSITION_INDEPENDENT_CODE ON) if(WIN32) @@ -222,14 +247,20 @@ if(BUILD_TEST) target_link_libraries(${test_name} RPU_CPU ${RPU_DEPENDENCY_LIBS}) if(${test_src} IN_LIST RPU_GPU_TEST_SRCS) - target_link_libraries(${test_name} torch_cuda c10_cuda cudart) - target_link_libraries(${test_name} RPU_GPU RPU_CPU cublas curand ${RPU_DEPENDENCY_LIBS}) - set_target_properties(${test_name} PROPERTIES ${CUDA_TARGET_PROPERTIES}) - set_property(TARGET ${test_name} PROPERTY CUDA_ARCHITECTURES ${RPU_CUDA_ARCHITECTURES}) - - if(WIN32) - target_link_libraries(${test_name} c10_cuda.lib torch_cuda.lib) - endif(WIN32) + if(USE_HIP) + target_link_libraries(${test_name} RPU_GPU RPU_CPU roc::hipblas hip::hiprand hip::hipcub ${RPU_DEPENDENCY_LIBS}) + set_target_properties(${test_name} PROPERTIES ${HIP_TARGET_PROPERTIES}) + set_property(TARGET ${test_name} PROPERTY HIP_ARCHITECTURES ${CMAKE_HIP_ARCHITECTURES}) + else() + target_link_libraries(${test_name} torch_cuda c10_cuda cudart) + target_link_libraries(${test_name} RPU_GPU RPU_CPU cublas curand ${RPU_DEPENDENCY_LIBS}) + set_target_properties(${test_name} PROPERTIES ${CUDA_TARGET_PROPERTIES}) + set_property(TARGET ${test_name} PROPERTY CUDA_ARCHITECTURES ${RPU_CUDA_ARCHITECTURES}) + + if(WIN32) + target_link_libraries(${test_name} c10_cuda.lib torch_cuda.lib) + endif(WIN32) + endif() endif() diff --git a/README.md b/README.md index 5b4b67a48..5175a6f2a 100644 --- a/README.md +++ b/README.md @@ -37,7 +37,7 @@ A series of primitives and features that allow using the toolkit within ### Analog devices simulator -A high-performant (CUDA-capable) C++ simulator that allows for +A high-performant (CUDA- and ROCm/HIP-capable) C++ simulator that allows for simulating a wide range of analog devices and crossbar configurations by using abstract functional models of material characteristics with adjustable parameters. Features include: diff --git a/cmake/dependencies_hip.cmake b/cmake/dependencies_hip.cmake new file mode 100644 index 000000000..5c9559209 --- /dev/null +++ b/cmake/dependencies_hip.cmake @@ -0,0 +1,47 @@ +# (C) Copyright 2020, 2021, 2022, 2023, 2024 IBM. All Rights Reserved. +# +# Licensed under the MIT license. See LICENSE file in the project root for details. + +# ROCm/HIP support, mirroring dependencies_cuda.cmake. The CUDA-spelled device +# sources compile under HIP through the cuda_to_hip.h shim; RPU_USE_CUDA stays +# defined so the device tile code and CudaAnalogTile pybind exposure are built. +if(USE_HIP) + if(NOT DEFINED ROCM_PATH) + if(DEFINED ENV{ROCM_PATH}) + set(ROCM_PATH $ENV{ROCM_PATH}) + else() + set(ROCM_PATH "/opt/rocm") + endif() + endif() + list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}) + + enable_language(HIP) + + find_package(hip REQUIRED) + find_package(hipblas REQUIRED) + find_package(hiprand REQUIRED) + find_package(hipcub REQUIRED) + + add_compile_definitions(RPU_USE_CUDA) + add_compile_definitions(USE_HIP) + + # Detect the torch hipify generation. hipify v1 renames CUDA spellings to HIP + # (c10::cuda::getCurrentCUDAStream is removed, only c10::hip::getCurrentHIPStream + # exists); hipify v2 masquerades, keeping c10::cuda::getCurrentCUDAStream as the + # public API while c10::hip::getCurrentHIPStream stays guarded by USE_ROCM (which + # this build does not define). rpu_base_tiles_cuda.cpp selects the c10 stream + # symbol from this, so probe the torch the build already uses. + execute_process(COMMAND "${RPU_PYTHON_EXECUTABLE}" -c "from torch.utils.hipify import __version__ as v; print(v)" + RESULT_VARIABLE TORCH_HIPIFY_VERSION_RESULT + OUTPUT_VARIABLE TORCH_HIPIFY_VERSION + OUTPUT_STRIP_TRAILING_WHITESPACE + ERROR_QUIET) + if(TORCH_HIPIFY_VERSION_RESULT EQUAL 0 AND NOT TORCH_HIPIFY_VERSION STREQUAL "") + message(STATUS "torch hipify version: ${TORCH_HIPIFY_VERSION}") + if(TORCH_HIPIFY_VERSION VERSION_GREATER_EQUAL "2.0.0") + set(RPU_TORCH_HIPIFY_V2 ON) + endif() + else() + message(STATUS "torch hipify version: unknown (assuming v1 stream API)") + endif() +endif() diff --git a/docs/source/developer_install.rst b/docs/source/developer_install.rst index 489108561..dba5e676d 100644 --- a/docs/source/developer_install.rst +++ b/docs/source/developer_install.rst @@ -102,6 +102,10 @@ Via python command $ python setup.py build_ext -DCMAKE_BUILD_TYPE=Debug -DCMAKE_EXPORT_COMPILE_COMMANDS=TRUE --inplace -DRPU_BLAS=MKL -j16 -DUSE_CUDA=ON -DRPU_CUDA_ARCHITECTURES="60;70" -DCMAKE_PREFIX_PATH=$CONDA_PREFIX + - GPU with MKL (AMD, ROCm/HIP):: + + $ python setup.py build_ext -DCMAKE_BUILD_TYPE=Debug -DCMAKE_EXPORT_COMPILE_COMMANDS=TRUE --inplace -DRPU_BLAS=MKL -j16 -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a -DCMAKE_PREFIX_PATH=$CONDA_PREFIX + If you want to use ``OpenBLAS`` instead ``MKL``, you need to set ``-DRPU_BLAS=OpenBLAS``. @@ -111,6 +115,11 @@ To identify the ``CUDA_ARCH`` for your GPU using ``nvidia-smi`` in your system:: $ export CUDA_ARCH=$(nvidia-smi --query-gpu=compute_cap --format=csv | sed -n '2 p' | tr -d '.') $ echo $CUDA_ARCH +For an AMD GPU, set ``-DCMAKE_HIP_ARCHITECTURES`` to its ``gfx`` target instead +(for example ``gfx90a`` or ``gfx1100``). To identify it:: + + $ rocminfo | grep -m1 -o 'gfx[0-9a-f]*' + This will produce a shared library under the ``src/aihwkit/simulator`` directory, without installing the package. For how to use the shared library see the :ref:`use-the-library` section below. @@ -222,6 +231,7 @@ compilation process: Flag Description Default ========================== ================================================ ======= ``USE_CUDA`` Build with CUDA support ``OFF`` +``USE_HIP`` Build with ROCm/HIP support (AMD GPUs) ``OFF`` ``BUILD_TEST`` Build the C++ test binaries ``OFF`` ``RPU_BLAS`` BLAS backend of choice (``OpenBLAS`` or ``MKL``) ``OpenBLAS`` ``RPU_USE_FASTMOD`` Use fast mod ``ON`` @@ -238,6 +248,18 @@ or if using ``cmake`` directly:: build$ cmake -DUSE_CUDA=ON -DRPU_CUDA_ARCHITECTURES="60;70" .. +For AMD GPUs, build with ROCm/HIP support instead (``USE_HIP`` and ``USE_CUDA`` +are mutually exclusive):: + + $ python setup.py build_ext --inplace -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a + +or if using ``cmake`` directly:: + + build$ cmake -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx90a .. + +Set ``-DCMAKE_HIP_ARCHITECTURES`` to your GPU's ``gfx`` target (for example +``gfx90a`` or ``gfx1100``). + Passing other ``cmake`` flags """"""""""""""""""""""""""""" diff --git a/src/aihwkit/simulator/CMakeLists.txt b/src/aihwkit/simulator/CMakeLists.txt index bf4963887..5a2ee3dec 100644 --- a/src/aihwkit/simulator/CMakeLists.txt +++ b/src/aihwkit/simulator/CMakeLists.txt @@ -5,16 +5,35 @@ set(python_module_name rpu_base) file(GLOB RPU_BINDINGS_SRCS rpu_base_src/*.cpp) +if (USE_HIP) + # The binding TUs pull in HIP device headers transitively from RPU_GPU and + # inherit its --offload-arch usage requirement, which only the HIP (clang) + # driver understands; compile them with the HIP toolchain. + set_source_files_properties(${RPU_BINDINGS_SRCS} PROPERTIES LANGUAGE HIP) +endif() pybind11_add_module(${python_module_name} MODULE ${RPU_BINDINGS_SRCS}) target_link_libraries(${python_module_name} PRIVATE torch_python) -set_target_properties(${python_module_name} PROPERTIES CXX_STANDARD 17) +set_target_properties(${python_module_name} PROPERTIES CXX_STANDARD ${RPU_CXX_STANDARD}) -if (USE_CUDA) +if (USE_CUDA OR USE_HIP) target_link_libraries(${python_module_name} PRIVATE RPU_GPU) else() target_link_libraries(${python_module_name} PRIVATE RPU_CPU) endif() +if (USE_HIP) + # HIP device-code linking does not finalize LTO, which yields an empty + # PyInit_ and an ImportError; keep IPO off for the pybind module. + set_target_properties(${python_module_name} PROPERTIES INTERPROCEDURAL_OPTIMIZATION OFF) + + # rpu_base_tiles_cuda.cpp picks the c10 current-stream symbol by hipify + # generation (see cmake/dependencies_hip.cmake); v2 keeps the c10::cuda + # spelling, v1 (undefined) uses c10::hip. + if(RPU_TORCH_HIPIFY_V2) + target_compile_definitions(${python_module_name} PRIVATE TORCH_HIPIFY_V2) + endif() +endif() + install(TARGETS ${python_module_name} DESTINATION "src/aihwkit/simulator") add_custom_command(TARGET ${python_module_name} diff --git a/src/aihwkit/simulator/rpu_base_src/rpu_base_tiles_cuda.cpp b/src/aihwkit/simulator/rpu_base_src/rpu_base_tiles_cuda.cpp index 1b1a93faa..38d640623 100644 --- a/src/aihwkit/simulator/rpu_base_src/rpu_base_tiles_cuda.cpp +++ b/src/aihwkit/simulator/rpu_base_src/rpu_base_tiles_cuda.cpp @@ -6,13 +6,39 @@ #ifdef RPU_USE_CUDA #define __RPU_CUDA_HALF_DEFINED +#if !defined(USE_HIP) #include "cuda.h" +#endif #include "cuda_util.h" #include "rpu_base.h" #include "rpucuda.h" #include "rpucuda_pulsed.h" +#if defined(USE_HIP) +// The full ATen HIP context header transitively pulls in hipSOLVER / hipSPARSE +// headers that are not part of this build; only the current-stream accessor is +// needed, so include the lightweight c10 stream header and surface it under the +// at::cuda namespace the call sites use. +#include +namespace at { +namespace cuda { +inline auto getCurrentCUDAStream(c10::DeviceIndex device_index = -1) { +#if defined(TORCH_HIPIFY_V2) + // torch hipify v2 masquerades: c10::cuda::getCurrentCUDAStream is the public + // stream API, while c10::hip::getCurrentHIPStream stays guarded by USE_ROCM, + // which this build does not define. + return c10::cuda::getCurrentCUDAStream(device_index); +#else + // torch hipify v1 renames CUDA spellings to HIP: getCurrentCUDAStream is gone + // from c10::cuda, so use the hip-prefixed equivalent. + return c10::hip::getCurrentHIPStream(device_index); +#endif +} +} +} // namespace at +#else #include +#endif #define CHECK_CUDA(x) \ TORCH_CHECK( \ diff --git a/src/rpucuda/cuda/cuda_to_hip.h b/src/rpucuda/cuda/cuda_to_hip.h new file mode 100644 index 000000000..0ab5ff463 --- /dev/null +++ b/src/rpucuda/cuda/cuda_to_hip.h @@ -0,0 +1,174 @@ +/** + * (C) Copyright 2020, 2021, 2022, 2023, 2024 IBM. All Rights Reserved. + * + * Licensed under the MIT license. See LICENSE file in the project root for details. + */ + +#pragma once + +// Compatibility shim that lets the CUDA-spelled RPUCuda sources compile and +// run on ROCm/HIP. It is included once from cuda_util.h before the CUDA +// runtime / cuBLAS / cuRAND headers. On a CUDA build (USE_HIP undefined) it is +// a no-op, so the CUDA path stays byte-identical. +// +// The host C string / allocation headers are pulled in before +// so that calls to memcpy/memset/abs resolve to the host / +// declarations and not to HIP __device__ overloads. + +#if defined(USE_HIP) + +#include +#include + +#include +#include +#include +#include +#include + +// ---- runtime --------------------------------------------------------------- +#define cudaError_t hipError_t +#define cudaSuccess hipSuccess +#define cudaGetErrorString hipGetErrorString +#define cudaPeekAtLastError hipPeekAtLastError +#define cudaDeviceSynchronize hipDeviceSynchronize +#define cudaDeviceReset hipDeviceReset +#define cudaGetDevice hipGetDevice +#define cudaSetDevice hipSetDevice +#define cudaDeviceProp hipDeviceProp_t +#define cudaGetDeviceProperties hipGetDeviceProperties + +#define cudaMalloc hipMalloc +#define cudaMallocPitch hipMallocPitch +#define cudaMallocHost hipHostMalloc +#define cudaFree hipFree +#define cudaFreeHost hipHostFree +#define cudaMemset hipMemset +#define cudaMemcpy hipMemcpy +#define cudaMemcpyAsync hipMemcpyAsync +#define cudaMemcpy2D hipMemcpy2D +#define cudaMemcpy2DAsync hipMemcpy2DAsync +#define cudaMemset2DAsync hipMemset2DAsync +#define cudaMemcpyHostToDevice hipMemcpyHostToDevice +#define cudaMemcpyDeviceToHost hipMemcpyDeviceToHost +#define cudaMemcpyDeviceToDevice hipMemcpyDeviceToDevice + +#define cudaStream_t hipStream_t +#define cudaStreamCreate hipStreamCreate +#define cudaStreamCreateWithFlags hipStreamCreateWithFlags +#define cudaStreamDestroy hipStreamDestroy +#define cudaStreamSynchronize hipStreamSynchronize +#define cudaStreamWaitEvent hipStreamWaitEvent +#define cudaStreamNonBlocking hipStreamNonBlocking + +#define cudaEvent_t hipEvent_t +#define cudaEventCreate hipEventCreate +#define cudaEventDestroy hipEventDestroy +#define cudaEventRecord hipEventRecord +#define cudaEventSynchronize hipEventSynchronize +#define cudaEventElapsedTime hipEventElapsedTime + +// ---- cuBLAS -> hipBLAS ----------------------------------------------------- +#define cublasHandle_t hipblasHandle_t +#define cublasStatus_t hipblasStatus_t +#define cublasPointerMode_t hipblasPointerMode_t +#define cublasCreate hipblasCreate +#define cublasDestroy hipblasDestroy +#define cublasSetStream hipblasSetStream +#define cublasGetPointerMode hipblasGetPointerMode +#define cublasSetPointerMode hipblasSetPointerMode +#define cublasSetMatrix hipblasSetMatrix +#define cublasGetMatrix hipblasGetMatrix + +#define cublasSgemm hipblasSgemm +#define cublasDgemm hipblasDgemm +#define cublasHgemm hipblasHgemm +#define cublasSgemv hipblasSgemv +#define cublasDgemv hipblasDgemv +#define cublasSger hipblasSger +#define cublasDger hipblasDger +#define cublasSscal hipblasSscal +#define cublasDscal hipblasDscal +#define cublasScopy hipblasScopy +#define cublasDcopy hipblasDcopy +#define cublasSnrm2 hipblasSnrm2 +#define cublasDnrm2 hipblasDnrm2 + +#define CUBLAS_OP_N HIPBLAS_OP_N +#define CUBLAS_OP_T HIPBLAS_OP_T +#define CUBLAS_POINTER_MODE_HOST HIPBLAS_POINTER_MODE_HOST +#define CUBLAS_POINTER_MODE_DEVICE HIPBLAS_POINTER_MODE_DEVICE + +#define CUBLAS_STATUS_SUCCESS HIPBLAS_STATUS_SUCCESS +#define CUBLAS_STATUS_NOT_INITIALIZED HIPBLAS_STATUS_NOT_INITIALIZED +#define CUBLAS_STATUS_INVALID_VALUE HIPBLAS_STATUS_INVALID_VALUE +#define CUBLAS_STATUS_ARCH_MISMATCH HIPBLAS_STATUS_ARCH_MISMATCH +#define CUBLAS_STATUS_MAPPING_ERROR HIPBLAS_STATUS_MAPPING_ERROR +#define CUBLAS_STATUS_EXECUTION_FAILED HIPBLAS_STATUS_EXECUTION_FAILED +#define CUBLAS_STATUS_INTERNAL_ERROR HIPBLAS_STATUS_INTERNAL_ERROR +#define CUBLAS_STATUS_NOT_SUPPORTED HIPBLAS_STATUS_NOT_SUPPORTED +// hipBLAS has no LICENSE_ERROR status; map it to a benign unknown so the +// CUBLAS_CALL switch still compiles. +#define CUBLAS_STATUS_LICENSE_ERROR HIPBLAS_STATUS_UNKNOWN + +// ---- cuRAND -> hipRAND (host generator + device per-thread state) ---------- +#define curandGenerator_t hiprandGenerator_t +#define curandStatus_t hiprandStatus_t +// In CUDA, curandState, curandState_t and curandStateXORWOW are all the same +// XORWOW type. hipRAND declares hiprandState and hiprandStateXORWOW_t as +// distinct struct types, so map all three spellings to a single hipRAND type +// to keep the project's template instantiations consistent. +#define curandState hiprandState_t +#define curandState_t hiprandState_t +#define curandStateXORWOW hiprandState_t +#define curandCreateGenerator hiprandCreateGenerator +#define curandDestroyGenerator hiprandDestroyGenerator +#define curandGenerateNormal hiprandGenerateNormal +#define curandGenerateUniform hiprandGenerateUniform +#define curandSetStream hiprandSetStream +#define curandSetPseudoRandomGeneratorSeed hiprandSetPseudoRandomGeneratorSeed +#define curand_init hiprand_init +#define curand_normal hiprand_normal +#define curand_uniform hiprand_uniform +#define CURAND_RNG_PSEUDO_DEFAULT HIPRAND_RNG_PSEUDO_DEFAULT + +#define CURAND_STATUS_SUCCESS HIPRAND_STATUS_SUCCESS +#define CURAND_STATUS_VERSION_MISMATCH HIPRAND_STATUS_VERSION_MISMATCH +#define CURAND_STATUS_NOT_INITIALIZED HIPRAND_STATUS_NOT_INITIALIZED +#define CURAND_STATUS_ALLOCATION_FAILED HIPRAND_STATUS_ALLOCATION_FAILED +#define CURAND_STATUS_TYPE_ERROR HIPRAND_STATUS_TYPE_ERROR +#define CURAND_STATUS_OUT_OF_RANGE HIPRAND_STATUS_OUT_OF_RANGE +#define CURAND_STATUS_LENGTH_NOT_MULTIPLE HIPRAND_STATUS_LENGTH_NOT_MULTIPLE +#define CURAND_STATUS_DOUBLE_PRECISION_REQUIRED HIPRAND_STATUS_DOUBLE_PRECISION_REQUIRED +#define CURAND_STATUS_LAUNCH_FAILURE HIPRAND_STATUS_LAUNCH_FAILURE +#define CURAND_STATUS_PREEXISTING_FAILURE HIPRAND_STATUS_PREEXISTING_FAILURE +#define CURAND_STATUS_INITIALIZATION_FAILED HIPRAND_STATUS_INITIALIZATION_FAILED +#define CURAND_STATUS_ARCH_MISMATCH HIPRAND_STATUS_ARCH_MISMATCH +#define CURAND_STATUS_INTERNAL_ERROR HIPRAND_STATUS_INTERNAL_ERROR + +// ---- warp-size-independent serialized pulse-train bit format --------------- +// +// bit_line_maker.cu packs the stochastic pulse train as 32-bit words, one word +// per 32-lane logical warp (laneId = threadIdx.x & 0x1f), and the pulsed-update +// kernels read those words back with >>5 / &0x1f / __popc. That on-device bit +// layout must stay byte-identical regardless of the physical wavefront width, +// so the producer's ballot/shuffle must operate at 32-lane logical-warp +// granularity. On a 64-wide CDNA wavefront each wavefront holds two independent +// 32-lane logical warps; the native HIP __ballot/__shfl span all 64 lanes, so +// the device-side wrappers below restrict them to the caller's own 32-lane +// subgroup. On a 32-wide RDNA wavefront they reduce to a plain width-32 op. +__device__ __forceinline__ unsigned int +__rpu_logical_warp32_ballot(int pred) { + const unsigned long long full = __ballot(pred); + // Shift the calling lane's own 32-lane subgroup down to bits 0..31. On + // wave32 __lane_id() is 0..31 so the shift is always 0; on wave64 lanes + // 32..63 select the high word. + const unsigned int shift = (__lane_id() & 0x20u); + return static_cast(full >> shift); +} +#define __ballot_sync(mask, pred) __rpu_logical_warp32_ballot(pred) +#define __shfl_sync(mask, val, src) __shfl(val, src, 32) +#define __shfl_up_sync(mask, val, delta) __shfl_up(val, delta, 32) +#define __shfl_down_sync(mask, val, delta) __shfl_down(val, delta, 32) + +#endif // USE_HIP diff --git a/src/rpucuda/cuda/cuda_util.h b/src/rpucuda/cuda/cuda_util.h index d8468170e..80f41d2c4 100644 --- a/src/rpucuda/cuda/cuda_util.h +++ b/src/rpucuda/cuda/cuda_util.h @@ -15,12 +15,16 @@ #include #include +#if defined(USE_HIP) +#include "cuda_to_hip.h" +#else #include "cublas_v2.h" #include "cuda.h" #include "cuda_runtime.h" #include #include #include +#endif #define RPU_BUFFER_IN 0 #define RPU_BUFFER_OUT 1 diff --git a/src/rpucuda/cuda/pwu_kernel_parameter_base.h b/src/rpucuda/cuda/pwu_kernel_parameter_base.h index 01e9564da..c064914f3 100644 --- a/src/rpucuda/cuda/pwu_kernel_parameter_base.h +++ b/src/rpucuda/cuda/pwu_kernel_parameter_base.h @@ -11,7 +11,7 @@ namespace RPU { template class PulsedRPUDeviceCuda; -template class PulsedUpdateMetaParameter; +template struct PulsedUpdateMetaParameter; template class BitLineMaker; template class ChoppedWeightOutput; diff --git a/src/rpucuda/cuda/rpu_cub.h b/src/rpucuda/cuda/rpu_cub.h index 48a756c44..6a6dd2548 100644 --- a/src/rpucuda/cuda/rpu_cub.h +++ b/src/rpucuda/cuda/rpu_cub.h @@ -6,6 +6,17 @@ #pragma once +#if defined(USE_HIP) + +// hipCUB exposes its algorithms in the fixed `hipcub` namespace and does not +// honor the CUB_NS_PREFIX wrapping CUDA's CUB allows, so qualify directly. +#include +#ifndef RPU_CUB_NS_QUALIFIER +#define RPU_CUB_NS_QUALIFIER hipcub:: +#endif + +#else + #ifndef RPU_CUB_NS_QUALIFIER #ifndef CUB_NS_QUALIFIER #undef CUB_NS_PREFIX @@ -20,3 +31,5 @@ #endif #include + +#endif diff --git a/src/rpucuda/rpu_linearstep_device.h b/src/rpucuda/rpu_linearstep_device.h index 0b5e735d2..08d90555d 100644 --- a/src/rpucuda/rpu_linearstep_device.h +++ b/src/rpucuda/rpu_linearstep_device.h @@ -69,7 +69,7 @@ BUILD_PULSED_DEVICE_META_PARAMETER( template struct SoftBoundsRPUDeviceMetaParameter : LinearStepRPUDeviceMetaParameter { - SoftBoundsRPUDeviceMetaParameter() { + SoftBoundsRPUDeviceMetaParameter() { // ctor, to overwrite default values to softbounds this->ls_decrease_up = (T)1.0; // meaning 0 update at bounds this->ls_decrease_down = (T)1.0; // meaning 0 update at bounds From 70577b578b7faa661611e905c0b9ec1a1d85bf33 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Tue, 23 Jun 2026 22:07:39 +0000 Subject: [PATCH 2/2] [ROCm] Auto-detect the HIP architecture instead of pinning gfx90a The top-level CMakeLists set CMAKE_HIP_ARCHITECTURES to gfx90a when it was unset, and that block runs before cmake/dependencies_hip.cmake reaches enable_language(HIP). Pinning ahead of the enable preempts CMake's own host-GPU detection, so a builder on a non-gfx90a AMD GPU who did not pass -DCMAKE_HIP_ARCHITECTURES would silently get gfx90a code objects that fail to load at runtime on their card. Drop the pin and let enable_language(HIP) resolve the architecture: it honors an explicit -DCMAKE_HIP_ARCHITECTURES, otherwise auto-detects the host GPU, and otherwise errors so the builder must choose. The RPU_GPU and test targets continue to read CMAKE_HIP_ARCHITECTURES for their per-target HIP_ARCHITECTURES property, now from the resolved value. Authored with assistance from Claude. Signed-off-by: Jeff Daily --- CMakeLists.txt | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 6af5eaf0e..2c099d35a 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -29,9 +29,8 @@ set(RPU_CXX_STANDARD "17" CACHE STRING "C++ standard used for all RPU targets") set(RPU_BLAS "OpenBLAS" CACHE STRING "BLAS backend of choice (OpenBLAS, MKL)") set(RPU_CUDA_ARCHITECTURES "75;80;89" CACHE STRING "Target CUDA architectures") -if(NOT DEFINED CMAKE_HIP_ARCHITECTURES) - set(CMAKE_HIP_ARCHITECTURES "gfx90a" CACHE STRING "Target HIP (ROCm) architectures") -endif() +# CMAKE_HIP_ARCHITECTURES is left to enable_language(HIP) to auto-detect the host +# GPU when not set explicitly via -DCMAKE_HIP_ARCHITECTURES. # Internal variables. set(CUDA_TARGET_PROPERTIES POSITION_INDEPENDENT_CODE ON