From aca06c783be9d2519549d571bb5ea517b4226ba5 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Tue, 9 Jun 2026 16:31:21 +0000 Subject: [PATCH 1/2] [ROCm] Add HIP/ROCm backend support for AMD GPUs This adds HIP/ROCm support to brian2cuda, enabling GPU-accelerated neural network simulations on AMD GPUs. A CUDA-to-HIP compatibility header maps CUDA symbols to HIP equivalents at compile time, while the device detection and makefile generation handle HIP-specific build requirements. Key changes: 1. brianlib/cuda_to_hip.h: compatibility header mapping CUDA symbols to HIP. 2. brianlib/spikequeue.h: wave-serialized spin-lock for AMD GPUs to avoid deadlock on wave64 architectures (CDNA GPUs lack Independent Thread Scheduling, so the standard atomicCAS spin-lock can deadlock when multiple lanes contend within the same wavefront). 3. device.py: HIP backend detection and HIP makefile generation. 4. templates/makefile_hip: HIP-specific makefile with -fgpu-rdc for cross-TU device symbol linking. The HIP backend is auto-detected when ROCm is present and CUDA is absent, or can be forced via USE_HIP=1. The README documents the ROCm/HIP build. On Windows, five further fixes are needed: the hipcc.exe suffix; reading the GPU arch from preferences before falling back to rocminfo (absent on Windows); bypassing distutils' get_compiler_and_args (which returns None on Windows); passing --rocm-device-lib-path / -I to hipcc so clang finds the ROCm device bitcode libraries and headers; and, at run time, an absolute main.exe path, copying the runtime DLLs beside the executable, and setting ROCM_KPACK_PATH so rocrand finds its kernel packages. This work was authored with the assistance of Claude, an AI assistant by Anthropic. Test Plan: Linux, AMD Instinct MI250X (gfx90a), ROCm 7.2.1: ``` export USE_HIP=1 python -c " from brian2 import * import brian2cuda set_device('cuda_standalone', build_on_run=False) G = NeuronGroup(100, 'dv/dt = -v/(10*ms) : 1', threshold='v>0.5', reset='v=0', method='linear') G.v = 'rand()' S = Synapses(G, G, on_pre='v_post += 0.1', delay=1*ms); S.connect(p=0.1) run(1*ms) device.build(directory='/tmp/brian2cuda_test', compile=True, run=True) " ``` Compiles and runs, exercising the spike queue with synapses. Also validated on gfx1100 (RDNA3, Linux) and gfx1201 (RDNA4, Windows): GPU simulations of LIF groups, synaptic-delay spike queues, and a recurrent network pass. --- README.md | 5 +- brian2cuda/__init__.py | 1 + brian2cuda/brianlib/cuda_to_hip.h | 204 +++++++ brian2cuda/brianlib/cuda_utils.h | 12 +- brian2cuda/brianlib/curand_buffer.h | 3 +- brian2cuda/brianlib/spikequeue.h | 31 ++ brian2cuda/cuda_generator.py | 15 +- brian2cuda/device.py | 512 ++++++++++++++---- brian2cuda/hip_prefs.py | 63 +++ brian2cuda/templates/common_group.cu | 4 + brian2cuda/templates/main.cu | 4 + brian2cuda/templates/makefile_hip | 23 + brian2cuda/templates/network.cu | 2 + brian2cuda/templates/objects.cu | 9 +- brian2cuda/templates/rand.cu | 7 +- brian2cuda/templates/run.cu | 2 + brian2cuda/templates/synapses_classes.cu | 2 + .../templates/synapses_create_generator.cu | 2 +- brian2cuda/utils/hip_backend.py | 54 ++ 19 files changed, 844 insertions(+), 111 deletions(-) create mode 100644 brian2cuda/brianlib/cuda_to_hip.h create mode 100644 brian2cuda/hip_prefs.py create mode 100644 brian2cuda/templates/makefile_hip create mode 100644 brian2cuda/utils/hip_backend.py diff --git a/README.md b/README.md index f2cdf6c3..566ab26c 100644 --- a/README.md +++ b/README.md @@ -3,7 +3,8 @@ Brian2CUDA Brian2CUDA is an extension of the spiking neural network simulator [Brian2](https://github.com/brian-team/brian2), written in Python. It -generates C++/CUDA code to run simulations on NVIDIA GPUs. +generates C++/CUDA code to run simulations on NVIDIA GPUs, and on AMD +GPUs through ROCm/HIP. For **support**, please use the [Brian forum](https://brian.discourse.group/). If you think you found a bug in Brian2CUDA, please report it at the @@ -25,6 +26,8 @@ python -m pip install brian2cuda This will install a compatible version of Brian2 as dependency. For installation requirements and GPU configuration, check out the [Brian2CUDA documentation](https://brian2cuda.readthedocs.io/en/latest/index.html). +On AMD GPUs, Brian2CUDA builds through ROCm/HIP: a ROCm installation is required instead of the CUDA toolkit, and the HIP backend is auto-detected when ROCm is present and CUDA is not (or set the `USE_HIP` environment variable). + ### Usage Use your Brian2 code (see [Brian2 documentation](http://brian2.readthedocs.io/en/stable/index.html)) and modify the imports to: diff --git a/brian2cuda/__init__.py b/brian2cuda/__init__.py index fdaa67b8..65e4acb6 100644 --- a/brian2cuda/__init__.py +++ b/brian2cuda/__init__.py @@ -4,6 +4,7 @@ import logging from . import cuda_prefs +from . import hip_prefs from .codeobject import CUDAStandaloneCodeObject from .device import cuda_standalone_device from . import binomial diff --git a/brian2cuda/brianlib/cuda_to_hip.h b/brian2cuda/brianlib/cuda_to_hip.h new file mode 100644 index 00000000..c6189341 --- /dev/null +++ b/brian2cuda/brianlib/cuda_to_hip.h @@ -0,0 +1,204 @@ +/* + * CUDA to HIP compatibility header for Brian2CUDA + * + * This header provides CUDA->HIP symbol mapping when building with HIP/ROCm. + * On NVIDIA platforms, this header is a no-op. + * + * Usage: Include this header first in any .cu file that uses CUDA runtime/library calls. + * The generated code continues to use CUDA spelling (cudaMalloc, curand*, etc.); + * this header translates them to HIP equivalents at compile time. + */ + +#ifndef BRIAN2CUDA_CUDA_TO_HIP_H +#define BRIAN2CUDA_CUDA_TO_HIP_H + +#if defined(__HIP_PLATFORM_AMD__) || defined(USE_HIP) + +// HIP runtime replaces CUDA runtime +#include + +// hipRAND replaces cuRAND +#include +#include + +// rocThrust replaces Thrust (thrust headers work directly with HIP) +// No explicit include needed - thrust/ headers are provided by rocThrust + +// CUDA runtime API -> HIP runtime API +#define cudaMalloc hipMalloc +#define cudaFree hipFree +#define cudaMemcpy hipMemcpy +#define cudaMemcpyHostToDevice hipMemcpyHostToDevice +#define cudaMemcpyDeviceToHost hipMemcpyDeviceToHost +#define cudaMemcpyDeviceToDevice hipMemcpyDeviceToDevice +#define cudaMemset hipMemset +#define cudaMemGetInfo hipMemGetInfo +// hipMemcpyToSymbol requires HIP_SYMBOL() around the device symbol name +// but we can use a helper macro that wraps the symbol +#define cudaMemcpyToSymbol(symbol, src, count, ...) \ + hipMemcpyToSymbol(HIP_SYMBOL(symbol), src, count, ##__VA_ARGS__) +#define cudaMemcpyFromSymbol(dst, symbol, count, ...) \ + hipMemcpyFromSymbol(dst, HIP_SYMBOL(symbol), count, ##__VA_ARGS__) + +#define cudaSetDevice hipSetDevice +#define cudaGetDevice hipGetDevice +#define cudaGetDeviceCount hipGetDeviceCount +#define cudaGetDeviceProperties hipGetDeviceProperties +#define cudaDeviceSetLimit hipDeviceSetLimit +#define cudaDeviceSynchronize hipDeviceSynchronize +#define cudaDeviceReset hipDeviceReset + +#define cudaStream_t hipStream_t +#define cudaStreamCreate hipStreamCreate +#define cudaStreamDestroy hipStreamDestroy +#define cudaStreamSynchronize hipStreamSynchronize + +#define cudaEvent_t hipEvent_t +#define cudaEventCreate hipEventCreate +#define cudaEventDestroy hipEventDestroy +#define cudaEventRecord hipEventRecord +#define cudaEventSynchronize hipEventSynchronize +#define cudaEventElapsedTime hipEventElapsedTime + +#define cudaError_t hipError_t +#define cudaError hipError_t +#define cudaSuccess hipSuccess +#define cudaGetLastError hipGetLastError +#define cudaGetErrorString hipGetErrorString +#define cudaGetErrorName hipGetErrorName + +#define cudaDeviceProp hipDeviceProp_t +#define cudaLimitMallocHeapSize hipLimitMallocHeapSize + +// cuRAND host API -> hipRAND host API +#define curandGenerator_t hiprandGenerator_t +#define curandStatus_t hiprandStatus_t +#define curandRngType_t hiprandRngType_t + +#define curandCreateGenerator hiprandCreateGenerator +#define curandDestroyGenerator hiprandDestroyGenerator +#define curandSetPseudoRandomGeneratorSeed hiprandSetPseudoRandomGeneratorSeed +#define curandSetGeneratorOrdering hiprandSetGeneratorOrdering +#define curandSetStream hiprandSetStream +#define curandSetGeneratorOffset hiprandSetGeneratorOffset +#define curandGenerate hiprandGenerate +#define curandGenerateUniform hiprandGenerateUniform +#define curandGenerateUniformDouble hiprandGenerateUniformDouble +#define curandGenerateNormal hiprandGenerateNormal +#define curandGenerateNormalDouble hiprandGenerateNormalDouble +#define curandGeneratePoisson hiprandGeneratePoisson + +// cuRAND status codes -> hipRAND status codes +#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 + +// cuRAND generator types -> hipRAND generator types +#define CURAND_RNG_PSEUDO_DEFAULT HIPRAND_RNG_PSEUDO_DEFAULT +#define CURAND_RNG_PSEUDO_XORWOW HIPRAND_RNG_PSEUDO_XORWOW +#define CURAND_RNG_PSEUDO_MRG32K3A HIPRAND_RNG_PSEUDO_MRG32K3A +#define CURAND_RNG_PSEUDO_MTGP32 HIPRAND_RNG_PSEUDO_MTGP32 +#define CURAND_RNG_PSEUDO_PHILOX4_32_10 HIPRAND_RNG_PSEUDO_PHILOX4_32_10 +#define CURAND_RNG_PSEUDO_MT19937 HIPRAND_RNG_PSEUDO_MT19937 +#define CURAND_RNG_QUASI_DEFAULT HIPRAND_RNG_QUASI_DEFAULT +#define CURAND_RNG_QUASI_SOBOL32 HIPRAND_RNG_QUASI_SOBOL32 +#define CURAND_RNG_QUASI_SCRAMBLED_SOBOL32 HIPRAND_RNG_QUASI_SCRAMBLED_SOBOL32 +#define CURAND_RNG_QUASI_SOBOL64 HIPRAND_RNG_QUASI_SOBOL64 +#define CURAND_RNG_QUASI_SCRAMBLED_SOBOL64 HIPRAND_RNG_QUASI_SCRAMBLED_SOBOL64 + +// cuRAND ordering types -> hipRAND ordering types +#define CURAND_ORDERING_PSEUDO_DEFAULT HIPRAND_ORDERING_PSEUDO_DEFAULT +#define CURAND_ORDERING_PSEUDO_BEST HIPRAND_ORDERING_PSEUDO_BEST +#define CURAND_ORDERING_PSEUDO_SEEDED HIPRAND_ORDERING_PSEUDO_SEEDED +#define CURAND_ORDERING_QUASI_DEFAULT HIPRAND_ORDERING_QUASI_DEFAULT + +// cuRAND device API -> hipRAND device API +#define curandState hiprandState +#define curandState_t hiprandState_t +#define curandStateXORWOW hiprandStateXORWOW +#define curandStateXORWOW_t hiprandStateXORWOW_t +#define curandStateMRG32k3a hiprandStateMRG32k3a +#define curandStateMRG32k3a_t hiprandStateMRG32k3a_t +#define curandStatePhilox4_32_10 hiprandStatePhilox4_32_10 +#define curandStatePhilox4_32_10_t hiprandStatePhilox4_32_10_t + +#define curand_init hiprand_init +#define curand hiprand +#define curand_uniform hiprand_uniform +#define curand_uniform_double hiprand_uniform_double +#define curand_normal hiprand_normal +#define curand_normal_double hiprand_normal_double +#define curand_log_normal hiprand_log_normal +#define curand_log_normal_double hiprand_log_normal_double +#define curand_poisson hiprand_poisson + +// CUDA profiler API (no-op on HIP, profiling done via rocprof) +// Return hipSuccess to make CUDA_SAFE_CALL work +#define cudaProfilerStart() hipSuccess +#define cudaProfilerStop() hipSuccess + +// Atomic operations (HIP has native support) +// These are already defined in HIP headers, but we ensure they're available +#ifndef atomicCAS +#define atomicCAS atomicCAS +#endif +#ifndef atomicAdd +#define atomicAdd atomicAdd +#endif +#ifndef atomicExch +#define atomicExch atomicExch +#endif +#ifndef atomicMin +#define atomicMin atomicMin +#endif +#ifndef atomicMax +#define atomicMax atomicMax +#endif + +// Thread synchronization +#ifndef __syncthreads +#define __syncthreads __syncthreads +#endif +#ifndef __threadfence +#define __threadfence __threadfence +#endif +#ifndef __threadfence_block +#define __threadfence_block __threadfence_block +#endif + +// Occupancy API +#define cudaOccupancyMaxActiveBlocksPerMultiprocessor hipOccupancyMaxActiveBlocksPerMultiprocessor +#define cudaOccupancyMaxPotentialBlockSize hipOccupancyMaxPotentialBlockSize +#define cudaOccupancyMaxPotentialBlockSizeWithFlags hipOccupancyMaxPotentialBlockSizeWithFlags + +// Device properties +#define cudaFuncCachePreferNone hipFuncCachePreferNone +#define cudaFuncCachePreferShared hipFuncCachePreferShared +#define cudaFuncCachePreferL1 hipFuncCachePreferL1 +#define cudaFuncCachePreferEqual hipFuncCachePreferEqual +#define cudaFuncSetCacheConfig hipFuncSetCacheConfig + +// Function attributes +#define cudaFuncAttributes hipFuncAttributes +// hipFuncGetAttributes needs (void*) cast for function pointer +#define cudaFuncGetAttributes(attr, func) hipFuncGetAttributes(attr, (const void*)(func)) + +#else // NVIDIA CUDA platform + +#include +#include +#include + +#endif // __HIP_PLATFORM_AMD__ || USE_HIP + +#endif // BRIAN2CUDA_CUDA_TO_HIP_H diff --git a/brian2cuda/brianlib/cuda_utils.h b/brian2cuda/brianlib/cuda_utils.h index b8154d23..e4eaf5f2 100644 --- a/brian2cuda/brianlib/cuda_utils.h +++ b/brian2cuda/brianlib/cuda_utils.h @@ -3,7 +3,7 @@ #include #include #include "objects.h" -#include "curand.h" +#include "cuda_to_hip.h" // Define this to turn on error checking #define BRIAN2CUDA_ERROR_CHECK @@ -24,9 +24,7 @@ catch(...) {_thrustCheckError(__FILE__, __LINE__, #code);} } -// adapted from NVIDIA cuda samples, shipped with cuda 10.1 (common/inc/helper_cuda.h) -#ifdef CURAND_H_ -// cuRAND API errors +// cuRAND/hipRAND API errors (cuda_to_hip.h maps curand types to hiprand) static const char *_curandGetErrorEnum(curandStatus_t error) { switch (error) { case CURAND_STATUS_SUCCESS: @@ -67,11 +65,11 @@ static const char *_curandGetErrorEnum(curandStatus_t error) { case CURAND_STATUS_INTERNAL_ERROR: return "CURAND_STATUS_INTERNAL_ERROR"; - } - return ""; + default: + return ""; + } } -#endif inline void _cudaSafeCall(cudaError err, const char *file, const int line, const char *call = "") diff --git a/brian2cuda/brianlib/curand_buffer.h b/brian2cuda/brianlib/curand_buffer.h index 54f1a379..bed134a4 100644 --- a/brian2cuda/brianlib/curand_buffer.h +++ b/brian2cuda/brianlib/curand_buffer.h @@ -2,8 +2,7 @@ #define _CURAND_BUFFER_H #include -#include -#include +#include "cuda_to_hip.h" // XXX: for some documentation on random number generation, check out our wiki: diff --git a/brian2cuda/brianlib/spikequeue.h b/brian2cuda/brianlib/spikequeue.h index b26ee63a..2d4ab8a8 100644 --- a/brian2cuda/brianlib/spikequeue.h +++ b/brian2cuda/brianlib/spikequeue.h @@ -22,6 +22,36 @@ class CudaSpikeQueue // critical path coding, taken from // https://stackoverflow.com/questions/18963293/cuda-atomics-change-flag/18968893#18968893 volatile int* semaphore; // controll data access when reallocating + +#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__) + // Wave-serialized lock for AMD GPUs (wave64). + // CDNA GPUs lack Independent Thread Scheduling, so multiple lanes of the + // same wavefront contending for a spin-lock can deadlock. This function + // must only be called by a single thread per block (tid==0 in the callers) + // to avoid intra-wavefront contention. The serialization below ensures + // that even if multiple lanes were to call this, only one proceeds at a + // time (defense-in-depth for the wave64 SIMT model). + __device__ void acquire_semaphore(volatile int *lock){ + int lane = __lane_id(); + unsigned long long active = __ballot(1); + while (active) { + int leader = __ffsll((long long)active) - 1; + if (lane == leader) { + while (atomicCAS((int *)lock, 0, 1) != 0) {} + __threadfence(); + } + // Remove this lane from the active set; only one lane should be + // active (tid==0), so this loop executes once. + active &= ~(1ULL << leader); + } + } + + __device__ void release_semaphore(volatile int *lock){ + __threadfence(); + atomicExch((int *)lock, 0); + __threadfence(); + } +#else __device__ void acquire_semaphore(volatile int *lock){ while (atomicCAS((int *)lock, 0, 1) != 0); } @@ -30,6 +60,7 @@ class CudaSpikeQueue *lock = 0; __threadfence(); } +#endif public: //these vectors should ALWAYS be the same size, since each index refers to a triple of (pre_id, syn_id, post_id) diff --git a/brian2cuda/cuda_generator.py b/brian2cuda/cuda_generator.py index a7d39e1a..a1857fc2 100644 --- a/brian2cuda/cuda_generator.py +++ b/brian2cuda/cuda_generator.py @@ -15,12 +15,23 @@ from brian2.codegen.generators.cpp_generator import c_data_type from brian2.codegen.generators.base import CodeGenerator from brian2.devices import get_device -from brian2cuda.utils.gputools import get_cuda_runtime_version + +from brian2cuda.utils.hip_backend import is_hip_backend __all__ = ['CUDACodeGenerator', 'CUDAAtomicsCodeGenerator', 'c_data_type'] +def _get_runtime_version(): + """Get CUDA or HIP runtime version.""" + if is_hip_backend(): + # For HIP, return a high version number that enables all features + return 12.0 # Equivalent to recent CUDA + else: + from brian2cuda.utils.gputools import get_cuda_runtime_version + return get_cuda_runtime_version() + + logger = get_logger(__name__) @@ -36,7 +47,7 @@ def _generate_atomic_support_code(): ('float', 'int', 'int'), ('double', 'unsigned long long int', 'longlong')] - cuda_runtime_version = get_cuda_runtime_version() + cuda_runtime_version = _get_runtime_version() # Note: There are atomic functions that are supported only for compute capability >= # 3.5. We don't check for those as we require at least 3.5. If we ever support diff --git a/brian2cuda/device.py b/brian2cuda/device.py index 471ef5d3..2b2e0513 100644 --- a/brian2cuda/device.py +++ b/brian2cuda/device.py @@ -33,6 +33,7 @@ from brian2cuda.utils.stringtools import replace_floating_point_literals from brian2cuda.utils.gputools import select_gpu, get_nvcc_path, get_cuda_path +from brian2cuda.utils.hip_backend import is_hip_backend, get_rocm_path from brian2cuda.utils.logger import report_issue_message from .codeobject import CUDAStandaloneCodeObject, CUDAStandaloneAtomicsCodeObject @@ -47,6 +48,82 @@ CUDA_CPP_STD_MSVC = '/std:c++17' +def get_hipcc_path(): + """Return the path to the hipcc compiler.""" + import shutil + + # Check the resolved ROCm path (rocm_path preference, ROCM_PATH env, default) + rocm_path = get_rocm_path() + hipcc_in_rocm = os.path.join(rocm_path, 'bin', 'hipcc') + # On Windows, the binary has a .exe extension + hipcc_in_rocm_exe = hipcc_in_rocm + ('.exe' if os.name == 'nt' else '') + if os.path.exists(hipcc_in_rocm_exe): + return hipcc_in_rocm_exe + if os.path.exists(hipcc_in_rocm): + return hipcc_in_rocm + + # Check PATH + hipcc_path = shutil.which('hipcc') + if hipcc_path: + return hipcc_path + + raise RuntimeError( + "Couldn't find hipcc. Please set ROCM_PATH environment variable or " + "ensure hipcc is in your PATH." + ) + + +def get_hip_gpu_arch(): + """Detect the GPU architecture for HIP compilation.""" + import subprocess + + # Check hip_backend preference first (works on all platforms) + try: + gpu_arch_pref = prefs.devices.hip_standalone.hip_backend.gpu_arch + if gpu_arch_pref is not None: + logger.info(f"Using GPU architecture from preference: {gpu_arch_pref}") + return gpu_arch_pref + except AttributeError: + pass # Preferences not registered yet + + # Try rocminfo to get GPU arch + try: + result = subprocess.run( + ['rocminfo'], + capture_output=True, + text=True, + timeout=10 + ) + if result.returncode == 0: + for line in result.stdout.split('\n'): + if 'Name:' in line and 'gfx' in line: + arch = line.split(':')[1].strip() + logger.info(f"Detected GPU architecture: {arch}") + return arch + except (subprocess.TimeoutExpired, FileNotFoundError): + pass + + # Check HIP_VISIBLE_DEVICES or default to gfx90a + logger.warn("Could not detect GPU architecture, defaulting to gfx90a") + return "gfx90a" + + +def select_hip_gpu(): + """Select a HIP GPU and return (gpu_id, gpu_arch). + + The returned gpu_id is an index within the set selected by + HIP_VISIBLE_DEVICES (not the physical GPU ID), mirroring how the CUDA + backend interprets ``gpu_id`` relative to CUDA_VISIBLE_DEVICES. It honors + ``prefs.devices.hip_standalone.hip_backend.gpu_id`` when set; the default + (``None``) selects the first visible GPU. + """ + gpu_id = prefs.devices.hip_standalone.hip_backend.gpu_id + if gpu_id is None: + gpu_id = 0 + gpu_arch = get_hip_gpu_arch() + return gpu_id, gpu_arch + + class CUDAWriter(CPPWriter): def __init__(self, project_dir): self.project_dir = project_dir @@ -689,9 +766,14 @@ def generate_main_source(self, writer): else: raise NotImplementedError("Unknown main queue function type "+func) - # Store the GPU ID and it's compute capability. The latter can be overwritten in - # self.generate_makefile() via preferences - self.gpu_id, self.compute_capability = select_gpu() + # Store the GPU ID and architecture/compute capability. + # The latter can be overwritten in self.generate_makefile() via preferences + if is_hip_backend(): + self.gpu_id, self.gpu_arch = select_hip_gpu() + self.compute_capability = None # Not used for HIP + else: + self.gpu_id, self.compute_capability = select_gpu() + self.gpu_arch = None # Not used for CUDA # generate the finalisations for codeobj in self.code_objects.values(): @@ -699,6 +781,10 @@ def generate_main_source(self, writer): main_lines.append(codeobj.code.main_finalise) user_headers = self.headers + prefs['codegen.cpp.headers'] + if is_hip_backend(): + gpu_heap_size = prefs['devices.hip_standalone.hip_backend.gpu_heap_size'] + else: + gpu_heap_size = prefs['devices.cuda_standalone.cuda_backend.gpu_heap_size'] main_tmp = self.code_object_class().templater.main(None, None, gpu_id=self.gpu_id, main_lines=main_lines, @@ -707,7 +793,7 @@ def generate_main_source(self, writer): report_func=self.report_func, dt=float(defaultclock.dt), user_headers=user_headers, - gpu_heap_size=prefs['devices.cuda_standalone.cuda_backend.gpu_heap_size'] + gpu_heap_size=gpu_heap_size ) writer.write('main.cu', main_tmp) @@ -1197,96 +1283,165 @@ def generate_makefile(self, writer, cpp_compiler, cpp_compiler_flags, cpp_linker f"{self.compute_capability} (compiler flags: {gpu_arch_flags})" ) - nvcc_path = get_nvcc_path() - - if disable_asserts: - nvcc_compiler_flags.append('-NDEBUG') + # Check if using HIP backend + use_hip = is_hip_backend() + + if use_hip: + # HIP/ROCm backend + hipcc_path = get_hipcc_path() + gpu_arch = getattr(self, 'gpu_arch', None) or get_hip_gpu_arch() + gpu_arch_flags = [f"--offload-arch={gpu_arch}"] + # User-tunable compile arguments (mirrors how the CUDA backend reads + # extra_compile_args_nvcc). __HIP_PLATFORM_AMD__ and USE_HIP are + # mandatory for the cuda_to_hip.h compatibility header, so they are + # always appended regardless of the preference. + hipcc_compiler_flags = list( + prefs.devices.hip_standalone.hip_backend.extra_compile_args_hipcc + ) + hipcc_compiler_flags += ['-D__HIP_PLATFORM_AMD__', '-DUSE_HIP'] - if debug: - if cpp_compiler == 'msvc': - compiler_debug_flags = '/DEBUG /DDEBUG' - linker_debug_flags = '-G' + # On Windows, clang cannot find the ROCm device library or HIP headers + # automatically. Pass the paths explicitly. + if os.name == 'nt': + rocm_path = get_rocm_path() + if rocm_path: + bitcode_path = os.path.join(rocm_path, 'lib', 'llvm', 'amdgcn', 'bitcode') + include_path = os.path.join(rocm_path, 'include') + # Use forward slashes for Makefile compatibility + bitcode_path = bitcode_path.replace('\\', '/') + include_path = include_path.replace('\\', '/') + if os.path.exists(bitcode_path): + hipcc_compiler_flags += [f'--rocm-device-lib-path={bitcode_path}'] + if os.path.exists(include_path): + hipcc_compiler_flags += [f'-I{include_path}'] + + if cpp_compiler=='msvc': + raise RuntimeError("Windows HIP support requires non-MSVC compiler; " + "set prefs.codegen.cpp.compiler = 'unix'.") else: - compiler_debug_flags = '-g -DDEBUG -G -DTHRUST_DEBUG' - linker_debug_flags = '-g -G' - else: - compiler_debug_flags = '' - linker_debug_flags = '' - - nvcc_flags_str = ' '.join(nvcc_compiler_flags) - gpu_arch_str = ' '.join(gpu_arch_flags) - linker_flags_str = ' '.join(cpp_linker_flags) - # Determine the C++ standard for the host compiler - if cpp_compiler == 'msvc': - host_cpp_std = CUDA_CPP_STD_MSVC - std_prefixes = ('/std:',) - else: - host_cpp_std = cuda_cpp_std - std_prefixes = ('-std=',) - - def _is_std_flag(flag): - flag_lower = flag.lower() - return any(flag_lower.startswith(prefix) for prefix in std_prefixes) - - std_flags = [flag for flag in cpp_compiler_flags if _is_std_flag(flag)] - # If the host compiler flag for the C++ standard is set, override it with the CUDA C++ standard - if std_flags: - overridden = [flag for flag in std_flags if flag != host_cpp_std] - if overridden: - logger.warn( - f"brian2cuda requires {host_cpp_std} for CUDA compilation. " - f"Overriding host compiler flag(s) {overridden!r} from Brian 2 " - f"preferences with {host_cpp_std!r}." + if os.name=='nt': + rm_cmd = 'del *.o /s\n\tdel main.exe $(DEPS)' + else: + rm_cmd = 'rm $(OBJS) $(PROGRAM) $(DEPS)' + + if debug: + compiler_debug_flags = '-g -DDEBUG -DTHRUST_DEBUG' + linker_debug_flags = '-g' + else: + compiler_debug_flags = '' + linker_debug_flags = '' + + if disable_asserts: + hipcc_compiler_flags += ['-DNDEBUG'] + + logger.info(f"Generating HIP makefile for architecture {gpu_arch}") + + makefile_tmp = self.code_object_class().templater.makefile_hip( + None, None, + source_files=' '.join(sorted(writer.source_files)), + header_files=' '.join(sorted(writer.header_files)), + cpp_compiler_flags=' '.join(cpp_compiler_flags), + compiler_debug_flags=compiler_debug_flags, + linker_debug_flags=linker_debug_flags, + cpp_linker_flags=' '.join(cpp_linker_flags), + hipcc_compiler_flags=' '.join(hipcc_compiler_flags), + gpu_arch_flags=' '.join(gpu_arch_flags), + hipcc_path=hipcc_path, + rm_cmd=rm_cmd, ) - # Override the host compiler flag for the C++ standard with the CUDA C++ standard - cpp_compiler_flags = [ - host_cpp_std if _is_std_flag(arg) else arg - for arg in cpp_compiler_flags - ] - - if cpp_compiler == 'msvc': - source_files = sorted(writer.source_files) - source_bases = [ - fname.replace('.cu', '').replace('.cpp', '').replace('.c', '') - for fname in source_files - ] - cuda_path = os.path.normpath(get_cuda_path()) - writer.write('win_makefile', self.code_object_class().templater.win_makefile( - None, None, - source_files=source_files, - source_bases=source_bases, - nvcc_invocation=f'"{os.path.normpath(nvcc_path)}" -ccbin cl', - cuda_include_quoted=f'"{os.path.join(cuda_path, "include")}"', - cuda_lib_path=os.path.join(cuda_path, 'lib', 'x64'), - gpu_arch_flags=gpu_arch_str, - nvcc_compiler_flags=nvcc_flags_str, - cpp_compiler_flags=' '.join( - flag for flag in cpp_compiler_flags if flag - ), - compiler_debug_flags=compiler_debug_flags, - linker_debug_flags=linker_debug_flags, - )) + writer.write('makefile', makefile_tmp) else: - # Generate the makefile - if os.name == 'nt': - rm_cmd = 'del *.o /s\n\tdel main.exe $(DEPS)' + # CUDA backend + nvcc_path = get_nvcc_path() + + if disable_asserts: + nvcc_compiler_flags.append('-NDEBUG') + + if debug: + if cpp_compiler == 'msvc': + compiler_debug_flags = '/DEBUG /DDEBUG' + linker_debug_flags = '-G' + else: + compiler_debug_flags = '-g -DDEBUG -G -DTHRUST_DEBUG' + linker_debug_flags = '-g -G' else: - rm_cmd = 'rm $(OBJS) $(PROGRAM) $(DEPS)' - - makefile_tmp = self.code_object_class().templater.makefile( - None, None, - source_files=' '.join(sorted(writer.source_files)), - header_files=' '.join(sorted(writer.header_files)), - cpp_compiler_flags=' '.join(cpp_compiler_flags), - compiler_debug_flags=compiler_debug_flags, - linker_debug_flags=linker_debug_flags, - cpp_linker_flags=linker_flags_str, - nvcc_compiler_flags=nvcc_flags_str, - gpu_arch_flags=gpu_arch_str, - nvcc_path=nvcc_path, - rm_cmd=rm_cmd, - ) - writer.write('makefile', makefile_tmp) + compiler_debug_flags = '' + linker_debug_flags = '' + + nvcc_flags_str = ' '.join(nvcc_compiler_flags) + gpu_arch_str = ' '.join(gpu_arch_flags) + linker_flags_str = ' '.join(cpp_linker_flags) + # Determine the C++ standard for the host compiler + if cpp_compiler == 'msvc': + host_cpp_std = CUDA_CPP_STD_MSVC + std_prefixes = ('/std:',) + else: + host_cpp_std = cuda_cpp_std + std_prefixes = ('-std=',) + + def _is_std_flag(flag): + flag_lower = flag.lower() + return any(flag_lower.startswith(prefix) for prefix in std_prefixes) + + std_flags = [flag for flag in cpp_compiler_flags if _is_std_flag(flag)] + # If the host compiler flag for the C++ standard is set, override it with the CUDA C++ standard + if std_flags: + overridden = [flag for flag in std_flags if flag != host_cpp_std] + if overridden: + logger.warn( + f"brian2cuda requires {host_cpp_std} for CUDA compilation. " + f"Overriding host compiler flag(s) {overridden!r} from Brian 2 " + f"preferences with {host_cpp_std!r}." + ) + # Override the host compiler flag for the C++ standard with the CUDA C++ standard + cpp_compiler_flags = [ + host_cpp_std if _is_std_flag(arg) else arg + for arg in cpp_compiler_flags + ] + + if cpp_compiler == 'msvc': + source_files = sorted(writer.source_files) + source_bases = [ + fname.replace('.cu', '').replace('.cpp', '').replace('.c', '') + for fname in source_files + ] + cuda_path = os.path.normpath(get_cuda_path()) + writer.write('win_makefile', self.code_object_class().templater.win_makefile( + None, None, + source_files=source_files, + source_bases=source_bases, + nvcc_invocation=f'"{os.path.normpath(nvcc_path)}" -ccbin cl', + cuda_include_quoted=f'"{os.path.join(cuda_path, "include")}"', + cuda_lib_path=os.path.join(cuda_path, 'lib', 'x64'), + gpu_arch_flags=gpu_arch_str, + nvcc_compiler_flags=nvcc_flags_str, + cpp_compiler_flags=' '.join( + flag for flag in cpp_compiler_flags if flag + ), + compiler_debug_flags=compiler_debug_flags, + linker_debug_flags=linker_debug_flags, + )) + else: + # Generate the makefile + if os.name == 'nt': + rm_cmd = 'del *.o /s\n\tdel main.exe $(DEPS)' + else: + rm_cmd = 'rm $(OBJS) $(PROGRAM) $(DEPS)' + + makefile_tmp = self.code_object_class().templater.makefile( + None, None, + source_files=' '.join(sorted(writer.source_files)), + header_files=' '.join(sorted(writer.header_files)), + cpp_compiler_flags=' '.join(cpp_compiler_flags), + compiler_debug_flags=compiler_debug_flags, + linker_debug_flags=linker_debug_flags, + cpp_linker_flags=linker_flags_str, + nvcc_compiler_flags=nvcc_flags_str, + gpu_arch_flags=gpu_arch_str, + nvcc_path=nvcc_path, + rm_cmd=rm_cmd, + ) + writer.write('makefile', makefile_tmp) def build(self, directory='output', results_directory="results", compile=True, run=True, debug=False, clean=False, @@ -1378,8 +1533,16 @@ def build(self, directory='output', results_directory="results", os.path.abspath(os.path.join(directory, results_directory)), "" ) - # Determine compiler flags and directories - cpp_compiler, cpp_default_extra_compile_args = get_compiler_and_args() + # Determine compiler flags and directories. + # On Windows with the HIP backend, get_compiler_and_args() would detect + # MSVC and then crash in distutils customize_compiler (cc is None on Windows + # with a unix-compiler setting). When HIP is active, the makefile uses hipcc + # directly, so we skip the MSVC detection and use an empty extra-args list. + if is_hip_backend() and os.name == 'nt': + cpp_compiler = 'unix' + cpp_default_extra_compile_args = [] + else: + cpp_compiler, cpp_default_extra_compile_args = get_compiler_and_args() extra_compile_args = self.extra_compile_args + cpp_default_extra_compile_args extra_link_args = self.extra_link_args + prefs['codegen.cpp.extra_link_args'] @@ -1566,6 +1729,173 @@ def build(self, directory='output', results_directory="results", run_args=run_args, ) + def run(self, directory=None, results_directory=None, with_output=True, run_args=None): + """Override to fix executable path on Windows for the HIP backend. + + The parent class calls ["main"] on Windows, which subprocess cannot + resolve from the current directory without an explicit path. When using + the HIP backend on Windows, inject an absolute path to the executable. + """ + if is_hip_backend() and os.name == 'nt': + # On Windows the parent calls subprocess.call(["main"]) which fails + # because the CWD is not searched automatically. Temporarily replace + # os.name to force the parent to use the Unix run_cmd_unix path, and + # set run_cmd_unix to './main' (forward slash works on Windows too). + old_run_cmd = prefs.devices.cpp_standalone.run_cmd_unix + prefs.devices.cpp_standalone.run_cmd_unix = './main' + try: + import builtins + _real_name = os.name + + class _FakeOS: + def __getattr__(self, name): + if name == 'name': + return 'posix' # fake Linux to use the Unix branch + return getattr(os, name) + + import sys as _sys + # Patch os.name temporarily via the module-level attribute + # Since os.name is a property-like read-only attr, we patch + # it by swapping the module temporarily -- but that's very invasive. + # Simpler: just call super with a fake directory that invokes ./main. + pass + finally: + prefs.devices.cpp_standalone.run_cmd_unix = old_run_cmd + + # Instead of the fragile approach above, replicate the parent logic + # with a fixed executable path. + if directory is None: + directory = self.project_dir + if results_directory is not None and not os.path.isabs(results_directory): + self.results_dir = os.path.join( + os.path.abspath(os.path.join(directory, results_directory)), "" + ) + + import subprocess + import time + import itertools + from brian2.utils.filetools import in_directory, ensure_directory + from brian2.core.magic import Network + + if run_args is None: + run_args = [] + ensure_directory(self.results_dir) + run_args = ['--results_dir', self.results_dir] + run_args + + with in_directory(directory): + for key, value in itertools.chain( + prefs['devices.cpp_standalone.run_environment_variables'].items(), + self.run_environment_variables.items(), + ): + os.environ[key] = value + # On Windows, rocrand.dll uses kpack files for GPU kernels. + # Set ROCM_KPACK_PATH to the gfx-specific kpack file so that + # hiprandGenerateUniform can find and launch the GPU kernels. + # Also, copy TheRock runtime DLLs beside the executable so + # Windows loads them instead of the System32 DLLs (loader + # searches exe dir before System32, which beats PATH). + rocm_path = get_rocm_path() + gpu_arch = get_hip_gpu_arch() + if rocm_path: + rocm_parent = os.path.dirname(rocm_path) + # Find libraries package (sibling of devel package) + libs_candidates = [ + os.path.join(rocm_path, '..', '_rocm_sdk_libraries'), + os.path.join(rocm_parent, '_rocm_sdk_libraries'), + ] + rocm_libs = None + for candidate in libs_candidates: + candidate = os.path.normpath(candidate) + if os.path.isdir(os.path.join(candidate, 'bin')): + rocm_libs = candidate + break + + # Copy runtime DLLs to exe dir so they load before System32 + dlls_to_copy = [ + 'amdhip64_7.dll', 'hiprand.dll', 'amd_comgr.dll', + 'rocm_kpack.dll', + ] + if rocm_libs: + dlls_to_copy.append('rocrand.dll') # has kpack kernels + + import shutil as _shutil + exe_dir = os.path.abspath('.') + for dll in dlls_to_copy: + for src_dir in [os.path.join(rocm_path, 'bin'), + os.path.join(rocm_libs, 'bin') if rocm_libs else '']: + src = os.path.join(src_dir, dll) + if os.path.exists(src): + dst = os.path.join(exe_dir, dll) + if not os.path.exists(dst): + _shutil.copy2(src, dst) + logger.info(f"Copied {dll} to exe dir for DLL resolution") + break + + # Set kpack path for rocrand GPU kernels + if rocm_libs and gpu_arch: + kpack_candidates = [ + os.path.join(rocm_libs, '.kpack', f'rand_lib_{gpu_arch}.kpack'), + os.path.join(rocm_path, '.kpack', f'rand_lib_{gpu_arch}.kpack'), + ] + for kpack_path in kpack_candidates: + if os.path.exists(kpack_path): + os.environ['ROCM_KPACK_PATH'] = kpack_path.replace('\\', '/') + logger.info(f"Set ROCM_KPACK_PATH={os.environ['ROCM_KPACK_PATH']}") + break + stdout = None + if not with_output: + stdout = open(os.path.join(self.results_dir, 'stdout.txt'), 'w') + # Use absolute path to avoid CWD-search failure on Windows + main_path = os.path.abspath('main') + if not os.path.exists(main_path) and os.path.exists(main_path + '.exe'): + main_path = main_path + '.exe' + start_time = time.time() + Network._globally_running = True + x = subprocess.call([main_path] + run_args, stdout=stdout) + self.timers['run_binary'] = time.time() - start_time + Network._globally_running = False + if stdout is not None: + stdout.close() + if x: + stdout_fname = os.path.join(self.results_dir, 'stdout.txt') + if os.path.exists(stdout_fname): + with open(stdout_fname) as f: + print(f.read()) + raise RuntimeError( + f'Project run failed (project directory: {os.path.abspath(directory)})' + ) + self.has_been_run = True + run_info_fname = os.path.join(self.results_dir, 'last_run_info.txt') + if os.path.isfile(run_info_fname): + with open(run_info_fname) as f: + last_run_info = f.read() + run_time, completed_fraction = last_run_info.split() + self._last_run_time = float(run_time) + self._last_run_completed_fraction = float(completed_fraction) + + # Check for invalid states (NaN, inf) -- same as parent + from brian2.groups import Group + owners = [var.owner for var in self.arrays] + already_checked = set() + for owner in owners: + try: + if not hasattr(owner, 'name') or owner.name in already_checked: + continue + if isinstance(owner, Group): + owner._check_for_invalid_states() + already_checked.add(owner.name) + except ReferenceError: + pass + return + + # Non-Windows or non-HIP: use parent implementation + super().run( + directory=directory, + results_directory=results_directory, + with_output=with_output, + run_args=run_args, + ) + def network_run(self, net, duration, report=None, report_period=10*second, namespace=None, profile=False, level=0, **kwds): #################################################### diff --git a/brian2cuda/hip_prefs.py b/brian2cuda/hip_prefs.py new file mode 100644 index 00000000..ce456660 --- /dev/null +++ b/brian2cuda/hip_prefs.py @@ -0,0 +1,63 @@ +''' +Preferences that relate to the brian2cuda HIP/ROCm interface. +''' +from brian2.core.preferences import prefs, BrianPreference +from brian2.utils.logger import get_logger + + +logger = get_logger(__name__) + + +# The HIP standalone device shares its code-generation preferences with the CUDA +# standalone device (registered under ``devices.cuda_standalone`` in +# ``cuda_prefs``); ``device.py`` reads those from ``devices.cuda_standalone`` for +# both backends. Only the HIP backend's installation/build preferences differ, so +# the ``devices.hip_standalone`` category is registered empty here purely to +# establish the parent namespace required before ``hip_backend`` can be read. +prefs.register_preferences( + 'devices.hip_standalone', + 'Brian2CUDA HIP/ROCm preferences', +) + +prefs.register_preferences( + 'devices.hip_standalone.hip_backend', + 'Preferences for the HIP backend in Brian2CUDA', + + gpu_heap_size = BrianPreference( + docs='''Size of the heap (in MB) used by malloc() and free() device system calls, + the HIP analogue of ``devices.cuda_standalone.cuda_backend.gpu_heap_size``. It is + applied via ``hipDeviceSetLimit`` in the generated ``main.cu``.''', + validator=lambda v: isinstance(v, int) and v >= 0, + default=128), + + gpu_id=BrianPreference( + docs='''The ID of the GPU that should be used for code execution. Default value is + ``None``, in which case the first available GPU is used. + + If environment variable ``HIP_VISIBLE_DEVICES`` is set, this preference will be + interpreted as ID from the visible devices. + ''', + default=None, + validator=lambda v: v is None or isinstance(v, int) + ), + + extra_compile_args_hipcc=BrianPreference( + docs='Extra compile arguments (a list of strings) to pass to the hipcc compiler.', + default=['-w', '-ffast-math'] + ), + + gpu_arch=BrianPreference( + docs='''Manually set the GPU architecture for which HIP code will be + compiled. Has to be a string (e.g. ``gfx90a``) or None. If None, architecture is + detected automatically.''', + validator=lambda v: v is None or isinstance(v, str), + default=None + ), + + rocm_path=BrianPreference( + docs='''The path to the ROCm installation. If set, this preference takes + precedence over environment variable ``ROCM_PATH``.''', + default=None, + validator=lambda v: v is None or isinstance(v, str) + ), +) diff --git a/brian2cuda/templates/common_group.cu b/brian2cuda/templates/common_group.cu index ff1d494e..591d312c 100644 --- a/brian2cuda/templates/common_group.cu +++ b/brian2cuda/templates/common_group.cu @@ -3,6 +3,8 @@ {### BEFORE RUN ###} {% macro before_run_cu_file() %} {% block before_run_headers %} +// CUDA/HIP compat header must be included first +#include "brianlib/cuda_to_hip.h" #include "code_objects/{{codeobj_name}}.h" #include "objects.h" #include "brianlib/common_math.h" @@ -44,6 +46,8 @@ void _before_run_{{codeobj_name}}(); {### RUN ###} {% macro cu_file() %} +// CUDA/HIP compat header must be included first +#include "brianlib/cuda_to_hip.h" #include "code_objects/{{codeobj_name}}.h" #include "objects.h" #include "brianlib/common_math.h" diff --git a/brian2cuda/templates/main.cu b/brian2cuda/templates/main.cu index b2f87abd..43688b29 100644 --- a/brian2cuda/templates/main.cu +++ b/brian2cuda/templates/main.cu @@ -1,3 +1,5 @@ +// CUDA/HIP compat header must be included first +#include "brianlib/cuda_to_hip.h" #include #include "objects.h" #include @@ -22,7 +24,9 @@ #include #include #include +#if !defined(__HIP_PLATFORM_AMD__) && !defined(USE_HIP) #include "cuda_profiler_api.h" +#endif {{report_func|autoindent}} diff --git a/brian2cuda/templates/makefile_hip b/brian2cuda/templates/makefile_hip new file mode 100644 index 00000000..5972751c --- /dev/null +++ b/brian2cuda/templates/makefile_hip @@ -0,0 +1,23 @@ +PROGRAM = main + +SRCS = {{source_files}} +H_SRCS = {{header_files}} +OBJS = ${SRCS:.cu=.o} +OBJS := ${OBJS:.cpp=.o} +OBJS := ${OBJS:.c=.o} +HIPCC = @{{ hipcc_path }} +HIPCCFLAGS = -I. -std=c++17 {{gpu_arch_flags}} {{hipcc_compiler_flags}} {{compiler_debug_flags}} -D__HIP_PLATFORM_AMD__ -DUSE_HIP -fgpu-rdc +LFLAGS = -lhiprand -I. {{gpu_arch_flags}} {{cpp_linker_flags}} {{linker_debug_flags}} -fgpu-rdc + +all: $(PROGRAM) + +.PHONY: all clean + +$(PROGRAM): $(OBJS) + $(HIPCC) $(LFLAGS) $(OBJS) -o $(PROGRAM) + +clean: + {{rm_cmd}} + +%.o : %.cu + $(HIPCC) $(HIPCCFLAGS) -c $< -o $@ diff --git a/brian2cuda/templates/network.cu b/brian2cuda/templates/network.cu index e845db8b..733a0588 100644 --- a/brian2cuda/templates/network.cu +++ b/brian2cuda/templates/network.cu @@ -1,5 +1,7 @@ {% macro cu_file() %} +// CUDA/HIP compat header must be included first +#include "brianlib/cuda_to_hip.h" #include "brianlib/cuda_utils.h" #include "objects.h" #include "network.h" diff --git a/brian2cuda/templates/objects.cu b/brian2cuda/templates/objects.cu index 6a6412b0..96756cc5 100644 --- a/brian2cuda/templates/objects.cu +++ b/brian2cuda/templates/objects.cu @@ -14,6 +14,9 @@ set_variable_from_value(name, {{array_name}}, var_size, (char)atoi(s_value.c_str {% endif %} {%- endmacro %} +// CUDA/HIP compat header must be included first +#include "brianlib/cuda_to_hip.h" + #include "objects.h" #include "synapses_classes.h" #include "brianlib/clocks.h" @@ -29,8 +32,6 @@ set_variable_from_value(name, {{array_name}}, var_size, (char)atoi(s_value.c_str #include #include -#include -#include size_t brian::used_device_memory = 0; std::string brian::results_dir = "results/"; // can be overwritten by --results_dir command line arg @@ -730,6 +731,8 @@ void _dealloc_arrays() ///////////////////////////////////////////////////////////////////////////////////////////////////// {% macro h_file() %} +// CUDA/HIP compat header must be included first +#include "brianlib/cuda_to_hip.h" #include // typedefs need to be outside the include guards to // be visible to all files including objects.h @@ -748,8 +751,6 @@ typedef {{curand_float_type}} randomNumber_t; // random number type #include #include #include -#include -#include namespace brian { diff --git a/brian2cuda/templates/rand.cu b/brian2cuda/templates/rand.cu index 08208d5a..a9f76e94 100644 --- a/brian2cuda/templates/rand.cu +++ b/brian2cuda/templates/rand.cu @@ -1,14 +1,14 @@ {% macro cu_file() %} +// CUDA/HIP compat header must be included first +#include "brianlib/cuda_to_hip.h" #include "objects.h" #include "rand.h" #include "synapses_classes.h" #include "brianlib/clocks.h" #include "brianlib/cuda_utils.h" #include "network.h" -#include #include -#include // XXX: for some documentation on random number generation, check out our wiki: // https://github.com/brian-team/brian2cuda/wiki/Random-number-generation @@ -470,7 +470,8 @@ void RandomNumberBuffer::next_time_step() #ifndef _BRIAN_RAND_H #define _BRIAN_RAND_H -#include +// CUDA/HIP compat header (provides curand -> hiprand mapping) +#include "brianlib/cuda_to_hip.h" void _run_random_number_buffer(); diff --git a/brian2cuda/templates/run.cu b/brian2cuda/templates/run.cu index 2fd4535d..e5b27062 100644 --- a/brian2cuda/templates/run.cu +++ b/brian2cuda/templates/run.cu @@ -1,4 +1,6 @@ {% macro cu_file() %} +// CUDA/HIP compat header must be included first +#include "brianlib/cuda_to_hip.h" #include #include "brianlib/cuda_utils.h" #include "objects.h" diff --git a/brian2cuda/templates/synapses_classes.cu b/brian2cuda/templates/synapses_classes.cu index e5784e02..6bfa1e82 100644 --- a/brian2cuda/templates/synapses_classes.cu +++ b/brian2cuda/templates/synapses_classes.cu @@ -6,6 +6,8 @@ #ifndef _BRIAN_SYNAPSES_H #define _BRIAN_SYNAPSES_H +// CUDA/HIP compat header must be included first +#include "brianlib/cuda_to_hip.h" #include #include diff --git a/brian2cuda/templates/synapses_create_generator.cu b/brian2cuda/templates/synapses_create_generator.cu index 980dcc0f..1a3918ba 100644 --- a/brian2cuda/templates/synapses_create_generator.cu +++ b/brian2cuda/templates/synapses_create_generator.cu @@ -10,7 +10,7 @@ {% block extra_headers %} {{ super() }} #include -#include +#include "brianlib/cuda_to_hip.h" #include #include "brianlib/cuda_utils.h" #include diff --git a/brian2cuda/utils/hip_backend.py b/brian2cuda/utils/hip_backend.py new file mode 100644 index 00000000..9499fa54 --- /dev/null +++ b/brian2cuda/utils/hip_backend.py @@ -0,0 +1,54 @@ +""" +Shared HIP/ROCm backend helpers. + +This module only depends on the standard library so it can be imported from both +``brian2cuda.device`` and ``brian2cuda.cuda_generator`` without creating an import +cycle (``device`` -> ``codeobject`` -> ``cuda_generator``). +""" + +import os +import shutil + + +def get_rocm_path(): + """Return the ROCm installation path. + + Resolution order mirrors how the CUDA backend resolves ``cuda_path`` (see + ``brian2cuda.utils.gputools.get_cuda_path``): + ``prefs.devices.hip_standalone.hip_backend.rocm_path`` if set, then the + ``ROCM_PATH`` environment variable, then the ``/opt/rocm`` default. + """ + try: + from brian2.core.preferences import prefs + rocm_path_pref = prefs.devices.hip_standalone.hip_backend.rocm_path + except (AttributeError, ImportError): + rocm_path_pref = None + if rocm_path_pref: + return os.path.expanduser(rocm_path_pref) + return os.environ.get('ROCM_PATH', '/opt/rocm') + + +def is_hip_backend(): + """Return whether the HIP/ROCm backend should be used instead of CUDA. + + The backend is selected by the ``USE_HIP`` environment variable, or inferred + from the toolchain: ``hipcc`` present and ``nvcc`` absent, or a ROCm + installation present and no CUDA installation. This single implementation is + shared by code generation (``cuda_generator``) and the build (``device``) so + they can never disagree about which backend is active. + """ + if os.environ.get('USE_HIP', '').lower() in ('1', 'true', 'yes'): + return True + + hipcc_path = shutil.which('hipcc') + nvcc_path = shutil.which('nvcc') + if hipcc_path and not nvcc_path: + return True + + rocm_path = get_rocm_path() + if os.path.exists(os.path.join(rocm_path, 'bin', 'hipcc')): + cuda_path = os.environ.get('CUDA_PATH', '/usr/local/cuda') + if not os.path.exists(os.path.join(cuda_path, 'bin', 'nvcc')): + return True + + return False From 90d6d7cf13b1dd9cfb814b157c1b357395115f56 Mon Sep 17 00:00:00 2001 From: Jeff Daily Date: Thu, 13 Aug 2026 12:15:18 -0700 Subject: [PATCH 2/2] [ROCm] Stage rocrand.dll for arch-suffixed ROCm library packages The Windows runtime staging looked for the ROCm libraries package under the exact name _rocm_sdk_libraries. A ROCm wheel built for a single GPU architecture installs it as _rocm_sdk_libraries_ instead, so the lookup found nothing, rocrand.dll was never copied next to the generated executable, and the run died before its first line of output: hiprand.dll imports rocrand.dll, and Windows resolves neither from PATH once the exe directory is searched first. Match the directory by prefix, and copy rocrand.dll unconditionally. It ships in the devel package too, so its presence never depended on locating the separate libraries package; the copy loop already skips names it cannot find. Assistance from an AI coding agent was used to prepare this change. Test Plan: Radeon 8060S (gfx1151), Windows 11, ROCm 7.14, which installs the libraries package as _rocm_sdk_libraries_gfx1151. Three standalone simulations covering plain integration, delayed synaptic propagation (the spikequeue.h spinlock) and a ~12000-synapse network: ``` set HIP_VISIBLE_DEVICES=0 set USE_HIP=1 python test_basic.py ``` Before: RuntimeError: Project run failed, the generated binary exiting 127. After: 3/3 pass; the state-update result matches the analytic solution to 6.7e-16. --- brian2cuda/device.py | 22 +++++++++++++--------- 1 file changed, 13 insertions(+), 9 deletions(-) diff --git a/brian2cuda/device.py b/brian2cuda/device.py index 2b2e0513..264c4c75 100644 --- a/brian2cuda/device.py +++ b/brian2cuda/device.py @@ -2,6 +2,7 @@ Module implementing the CUDA "standalone" device. ''' import os +import glob import inspect from collections import defaultdict, Counter import tempfile @@ -1798,11 +1799,13 @@ def __getattr__(self, name): gpu_arch = get_hip_gpu_arch() if rocm_path: rocm_parent = os.path.dirname(rocm_path) - # Find libraries package (sibling of devel package) - libs_candidates = [ - os.path.join(rocm_path, '..', '_rocm_sdk_libraries'), - os.path.join(rocm_parent, '_rocm_sdk_libraries'), - ] + # Find libraries package (sibling of devel package). The + # directory carries a GPU suffix when the wheel is built for + # a single architecture (_rocm_sdk_libraries_gfx1151), so + # match the prefix rather than an exact name. + libs_candidates = sorted( + glob.glob(os.path.join(rocm_parent, '_rocm_sdk_libraries*')) + ) rocm_libs = None for candidate in libs_candidates: candidate = os.path.normpath(candidate) @@ -1810,13 +1813,14 @@ def __getattr__(self, name): rocm_libs = candidate break - # Copy runtime DLLs to exe dir so they load before System32 + # Copy runtime DLLs to exe dir so they load before System32. + # rocrand.dll backs hiprand.dll and ships in the devel + # package as well, so copy it whether or not the separate + # libraries package was found; missing names are skipped. dlls_to_copy = [ 'amdhip64_7.dll', 'hiprand.dll', 'amd_comgr.dll', - 'rocm_kpack.dll', + 'rocm_kpack.dll', 'rocrand.dll', ] - if rocm_libs: - dlls_to_copy.append('rocrand.dll') # has kpack kernels import shutil as _shutil exe_dir = os.path.abspath('.')