[ROCm] Add HIP/ROCm backend support for AMD GPUs

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:
 

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

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 <hip/hip_runtime.h>
+
+// hipRAND replaces cuRAND
+#include <hiprand/hiprand.h>
+#include <hiprand/hiprand_kernel.h>
+
+// 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 <cuda_runtime.h>
+#include <curand.h>
+#include <curand_kernel.h>
+
+#endif  // __HIP_PLATFORM_AMD__ || USE_HIP
+
+#endif  // BRIAN2CUDA_CUDA_TO_HIP_H

brian2cuda/brianlib/cuda_utils.h

@@ -3,7 +3,7 @@
 #include <stdio.h>
 #include <thrust/system_error.h>
 #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 "<unknown>";
+    default:
+      return "<unknown>";
+  }
 }
-#endif
 
 
 inline void _cudaSafeCall(cudaError err, const char *file, const int line, const char *call = "")

brian2cuda/brianlib/curand_buffer.h

@@ -2,8 +2,7 @@
 #define _CURAND_BUFFER_H
 
 #include <stdio.h>
-#include <curand.h>
-#include <cuda.h>
+#include "cuda_to_hip.h"
 
 
 // XXX: for some documentation on random number generation, check out our wiki:

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)

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