diff --git a/brian2cuda/device.py b/brian2cuda/device.py index b4610ee3..0f134042 100644 --- a/brian2cuda/device.py +++ b/brian2cuda/device.py @@ -49,6 +49,17 @@ ''', ), + gpu_heap_size = BrianPreference( + docs=''' + Size of the heap (in MB) used by malloc() and free() device system calls, which + are used in the `cudaVector` implementation. `cudaVectors` are used to + dynamically allocate device memory for `Spikemonitors` and the synapse + queues in the `CudaSpikeQueue` implementation for networks with + heterogeneously distributed delays. + ''', + validator=lambda v: isinstance(v, int) and v >= 0, + default=128), + curand_float_type=BrianPreference( docs=''' Floating point type of generated random numbers (float/double). @@ -373,7 +384,8 @@ def generate_main_source(self, writer, main_includes): code_objects=self.code_objects.values(), report_func=self.report_func, dt=float(defaultclock.dt), - additional_headers=main_includes + additional_headers=main_includes, + gpu_heap_size=prefs['devices.cuda_standalone.gpu_heap_size'] ) writer.write('main.cu', main_tmp) diff --git a/brian2cuda/templates/main.cu b/brian2cuda/templates/main.cu index af7cedf0..6c56044a 100644 --- a/brian2cuda/templates/main.cu +++ b/brian2cuda/templates/main.cu @@ -28,7 +28,7 @@ int main(int argc, char **argv) cudaDeviceProp props; cudaGetDeviceProperties(&props, 0); - size_t limit = 128 * 1024 * 1024; + size_t limit = {{gpu_heap_size}} * 1024 * 1024; cudaDeviceSetLimit(cudaLimitMallocHeapSize, limit); cudaDeviceSynchronize();