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..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 @@ -33,6 +34,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 +49,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 +767,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 +782,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 +794,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 +1284,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 +1534,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 +1730,176 @@ 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). 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) + if os.path.isdir(os.path.join(candidate, 'bin')): + rocm_libs = candidate + break + + # 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', 'rocrand.dll', + ] + + 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