diff --git a/network/spike_generation_cpu_cpp/PySpikeGenerationCPU.cpp b/network/spike_generation_cpu_cpp/PySpikeGenerationCPU.cpp index 0897c39..c67b35d 100644 --- a/network/spike_generation_cpu_cpp/PySpikeGenerationCPU.cpp +++ b/network/spike_generation_cpu_cpp/PySpikeGenerationCPU.cpp @@ -5,10 +5,15 @@ namespace py = pybind11; -PYBIND11_MODULE(PySpikeGenerationCPU, m) { +PYBIND11_MODULE(PySpikeGenerationCPU, m) +{ m.doc() = "SpikeGenerationCPU Module"; py::class_(m, "SpikeGenerationCPU") .def(py::init<>()) + .def("gpu_occupancy_export", + &SpikeGenerationCPU::gpu_occupancy_export) + .def("gpu_occupancy_import", + &SpikeGenerationCPU::gpu_occupancy_import) .def("spike_generation", &SpikeGenerationCPU::entrypoint); } diff --git a/network/spike_generation_cpu_cpp/SpikeGenerationCPU.cpp b/network/spike_generation_cpu_cpp/SpikeGenerationCPU.cpp index 104fe94..a149ad3 100644 --- a/network/spike_generation_cpu_cpp/SpikeGenerationCPU.cpp +++ b/network/spike_generation_cpu_cpp/SpikeGenerationCPU.cpp @@ -198,3 +198,23 @@ size_t SpikeGenerationCPU::lower_bound(float* data_ptr, } return start_of_range; }; + +void SpikeGenerationCPU::gpu_occupancy_export( + size_t dim_x, + size_t dim_y, + size_t number_of_pattern, + size_t spike_dim, + int64_t setting_memory_addr, + size_t setting_dim_0, + size_t setting_dim_1) +{ + return; +}; + +void SpikeGenerationCPU::gpu_occupancy_import( + int64_t setting_memory_addr, + size_t setting_dim_0, + size_t setting_dim_1) +{ + return; +}; \ No newline at end of file diff --git a/network/spike_generation_cpu_cpp/SpikeGenerationCPU.h b/network/spike_generation_cpu_cpp/SpikeGenerationCPU.h index f8b731a..58d78d1 100644 --- a/network/spike_generation_cpu_cpp/SpikeGenerationCPU.h +++ b/network/spike_generation_cpu_cpp/SpikeGenerationCPU.h @@ -6,33 +6,69 @@ #include #include -class SpikeGenerationCPU { -public: +class SpikeGenerationCPU +{ + public: SpikeGenerationCPU(); ~SpikeGenerationCPU(); void entrypoint( - int64_t input_pointer_addr, int64_t input_dim_0, int64_t input_dim_1, - int64_t input_dim_2, int64_t input_dim_3, - int64_t random_values_pointer_addr, int64_t random_values_dim_0, - int64_t random_values_dim_1, int64_t random_values_dim_2, - int64_t random_values_dim_3, int64_t output_pointer_addr, - int64_t output_dim_0, int64_t output_dim_1, int64_t output_dim_2, - int64_t output_dim_3, int64_t number_of_cpu_processes); + int64_t input_pointer_addr, + int64_t input_dim_0, + int64_t input_dim_1, + int64_t input_dim_2, + int64_t input_dim_3, + int64_t random_values_pointer_addr, + int64_t random_values_dim_0, + int64_t random_values_dim_1, + int64_t random_values_dim_2, + int64_t random_values_dim_3, + int64_t output_pointer_addr, + int64_t output_dim_0, + int64_t output_dim_1, + int64_t output_dim_2, + int64_t output_dim_3, + int64_t number_of_cpu_processes); -private: - void spike_generation(float* input_pointer, size_t input_dim_c0, - size_t input_dim_c1, size_t input_dim_c2, + void gpu_occupancy_export( + size_t dim_x, + size_t dim_y, + size_t number_of_pattern, + size_t spike_dim, + int64_t setting_memory_addr, + size_t setting_dim_0, + size_t setting_dim_1); + + void gpu_occupancy_import( + int64_t setting_memory_addr, + size_t setting_dim_0, + size_t setting_dim_1); + + private: + void spike_generation( + float* input_pointer, + size_t input_dim_c0, + size_t input_dim_c1, + size_t input_dim_c2, float* random_values_pointer, size_t random_values_dim_c0, size_t random_values_dim_c1, - size_t random_values_dim_c2, int64_t* output_pointer, - size_t output_dim_c0, size_t output_dim_c1, - size_t output_dim_c2, size_t x_dim, size_t y_dim, - size_t spike_dim, size_t h_dim, size_t pattern_id); + size_t random_values_dim_c2, + int64_t* output_pointer, + size_t output_dim_c0, + size_t output_dim_c1, + size_t output_dim_c2, + size_t x_dim, + size_t y_dim, + size_t spike_dim, + size_t h_dim, + size_t pattern_id); - size_t lower_bound(float* data_ptr, size_t data_length, - size_t data_ptr_stride, float compare_to_value); + size_t lower_bound( + float* data_ptr, + size_t data_length, + size_t data_ptr_stride, + float compare_to_value); }; #endif /* SPIKEGENERATIONCPU */ diff --git a/network/spike_generation_gpu_cpp_v1/SpikeGenerationGPU.cu b/network/spike_generation_gpu_cpp_v1/SpikeGenerationGPU.cu index d439207..70f7fb1 100644 --- a/network/spike_generation_gpu_cpp_v1/SpikeGenerationGPU.cu +++ b/network/spike_generation_gpu_cpp_v1/SpikeGenerationGPU.cu @@ -22,19 +22,19 @@ SpikeGenerationGPU::~SpikeGenerationGPU() void SpikeGenerationGPU::entrypoint( int64_t input_pointer_addr, int64_t input_dim_0, - int64_t input_dim_1, - int64_t input_dim_2, + int64_t input_dim_1, + int64_t input_dim_2, int64_t input_dim_3, - int64_t random_values_pointer_addr, + int64_t random_values_pointer_addr, int64_t random_values_dim_0, - int64_t random_values_dim_1, + int64_t random_values_dim_1, int64_t random_values_dim_2, - int64_t random_values_dim_3, + int64_t random_values_dim_3, int64_t output_pointer_addr, - int64_t output_dim_0, - int64_t output_dim_1, + int64_t output_dim_0, + int64_t output_dim_1, int64_t output_dim_2, - int64_t output_dim_3, + int64_t output_dim_3, int64_t number_of_cpu_processes) { @@ -87,7 +87,7 @@ void SpikeGenerationGPU::entrypoint( size_t x_dim = output_dim_2; size_t y_dim = output_dim_2; - assert ((number_of_cpu_processes <= 0)); + assert((number_of_cpu_processes <= 0)); gpu_spike_generation( input_pointer, @@ -208,9 +208,9 @@ void SpikeGenerationGPU::gpu_spike_generation( assert((x_dim < 65535)); assert((y_dim < 65535)); - // ////////////////////////////////////// - // Calculate the distribution on the GPU - // ////////////////////////////////////// + // ////////////////////////////////////// + // Calculate the distribution on the GPU + // ////////////////////////////////////// int min_grid_size; int block_size; @@ -222,6 +222,15 @@ void SpikeGenerationGPU::gpu_spike_generation( status = cudaOccupancyMaxPotentialBlockSize(&min_grid_size, &block_size, (void*)kernel_spike_generation, dynamic_s_mem_size, max_threadable_tasks); + if (status != cudaSuccess) + { + std::cerr << "CUDA Runtime Error at: " + << __FILE__ + << ":" + << __LINE__ + << std::endl; + std::cerr << cudaGetErrorString(status) << std::endl; + } assert((status == cudaSuccess)); grid_size = ((number_of_pattern * spike_dim) + block_size - 1) / block_size; @@ -248,8 +257,17 @@ void SpikeGenerationGPU::gpu_spike_generation( h_dim, (number_of_pattern * spike_dim)); - cudaDeviceSynchronize(); - + status = cudaDeviceSynchronize(); + if (status != cudaSuccess) + { + std::cerr << "CUDA Runtime Error at: " + << __FILE__ + << ":" + << __LINE__ + << std::endl; + std::cerr << cudaGetErrorString(status) << std::endl; + } + assert((status == cudaSuccess)); return; }; diff --git a/network/spike_generation_gpu_cpp_v2/kernel_spike_generation.cu b/network/spike_generation_gpu_cpp_v2/kernel_spike_generation.cu index d319eb7..ae33dd7 100644 --- a/network/spike_generation_gpu_cpp_v2/kernel_spike_generation.cu +++ b/network/spike_generation_gpu_cpp_v2/kernel_spike_generation.cu @@ -98,6 +98,15 @@ void occupancy_kernel_spike_generation( status = cudaOccupancyMaxPotentialBlockSize( &min_grid_size, &thread_block_size, (void*)kernel_spike_generation, 0, max_threadable_tasks); + if (status != cudaSuccess) + { + std::cerr << "CUDA Runtime Error at: " + << __FILE__ + << ":" + << __LINE__ + << std::endl; + std::cerr << cudaGetErrorString(status) << std::endl; + } assert((status == cudaSuccess)); grid_size =