diff --git a/network/spike_generation_gpu_cpp/Makefile b/network/spike_generation_gpu_cpp/Makefile deleted file mode 100644 index 5817c60..0000000 --- a/network/spike_generation_gpu_cpp/Makefile +++ /dev/null @@ -1,33 +0,0 @@ -include ../.env -export - -name = SpikeGeneration -type = GPU - -PYPOSTFIX := $(shell $(PYBIN)python3-config --extension-suffix) -PYBIND11INCLUDE := $(shell $(PYBIN)python3 -m pybind11 --includes) -PARAMETERS_O = $(PARAMETERS_O_GPU) $(PYBIND11INCLUDE) -PARAMETERS_Linker = $(PARAMETERS_Linker_GPU) - -so_file = Py$(name)$(type)$(PYPOSTFIX) -pyi_file = Py$(name)$(type).pyi -all: ../$(so_file) - -$(O_DIRS)$(name)$(type).o: $(name)$(type).h $(name)$(type).cu - mkdir -p $(O_DIRS) - $(NVCC) $(PARAMETERS_O) -c $(name)$(type).cu -o $(O_DIRS)$(name)$(type).o - -$(O_DIRS)Py$(name)$(type).o: $(name)$(type).h Py$(name)$(type).cpp - mkdir -p $(O_DIRS) - $(NVCC) $(PARAMETERS_O) -c Py$(name)$(type).cpp -o $(O_DIRS)Py$(name)$(type).o - -../$(so_file): $(O_DIRS)$(name)$(type).o $(O_DIRS)Py$(name)$(type).o - $(NVCC) $(PARAMETERS_Linker) -o ../$(so_file) $(O_DIRS)$(name)$(type).o $(O_DIRS)Py$(name)$(type).o - - -####################### -clean: - rm -rf $(O_DIRS) - rm -f ../$(so_file) - rm -f ../$(pyi_file) - diff --git a/network/spike_generation_gpu_cpp/PySpikeGenerationGPU.cpp b/network/spike_generation_gpu_cpp/PySpikeGenerationGPU.cpp deleted file mode 100644 index 70f4581..0000000 --- a/network/spike_generation_gpu_cpp/PySpikeGenerationGPU.cpp +++ /dev/null @@ -1,19 +0,0 @@ - -#include - -#include "SpikeGenerationGPU.h" - -namespace py = pybind11; - -PYBIND11_MODULE(PySpikeGenerationGPU, m) -{ - m.doc() = "SpikeGenerationGPU Module"; - py::class_(m, "SpikeGenerationGPU") - .def(py::init<>()) - .def("gpu_occupancy_export", - &SpikeGenerationGPU::gpu_occupancy_export) - .def("gpu_occupancy_import", - &SpikeGenerationGPU::gpu_occupancy_import) - .def("spike_generation", - &SpikeGenerationGPU::entrypoint); -} diff --git a/network/spike_generation_gpu_cpp/SpikeGenerationGPU.cu b/network/spike_generation_gpu_cpp/SpikeGenerationGPU.cu deleted file mode 100644 index d439207..0000000 --- a/network/spike_generation_gpu_cpp/SpikeGenerationGPU.cu +++ /dev/null @@ -1,275 +0,0 @@ -#include "SpikeGenerationGPU.h" - -#include -#include -#include - -#include -#include -#include - - -SpikeGenerationGPU::SpikeGenerationGPU() -{ - -}; - -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_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) -{ - - float* input_pointer = (float*)input_pointer_addr; - float* random_values_pointer = (float*)random_values_pointer_addr; - int64_t* output_pointer = (int64_t*)output_pointer_addr; - - // Input - assert((input_pointer != nullptr)); - assert((input_dim_0 > 0)); - assert((input_dim_1 > 0)); - assert((input_dim_2 > 0)); - assert((input_dim_3 > 0)); - - // Random - assert((random_values_pointer != nullptr)); - assert((random_values_dim_0 > 0)); - assert((random_values_dim_1 > 0)); - assert((random_values_dim_2 > 0)); - assert((random_values_dim_3 > 0)); - - // Output - assert((output_pointer != nullptr)); - assert((output_dim_0 > 0)); - assert((output_dim_1 > 0)); - assert((output_dim_2 > 0)); - assert((output_dim_3 > 0)); - - // Input - size_t input_dim_c0 = input_dim_1 * input_dim_2 * input_dim_3; - size_t input_dim_c1 = input_dim_2 * input_dim_3; - size_t input_dim_c2 = input_dim_3; - - // Random - size_t random_values_dim_c0 = - random_values_dim_1 * random_values_dim_2 * random_values_dim_3; - size_t random_values_dim_c1 = - random_values_dim_2 * random_values_dim_3; - size_t random_values_dim_c2 = random_values_dim_3; - - // Output - size_t output_dim_c0 = - output_dim_1 * output_dim_2 * output_dim_3; - size_t output_dim_c1 = output_dim_2 * output_dim_3; - size_t output_dim_c2 = output_dim_3; - - size_t number_of_pattern = input_dim_0; - size_t h_dim = input_dim_1; - size_t spike_dim = output_dim_1; - size_t x_dim = output_dim_2; - size_t y_dim = output_dim_2; - - assert ((number_of_cpu_processes <= 0)); - - gpu_spike_generation( - input_pointer, - input_dim_c0, - input_dim_c1, - input_dim_c2, - random_values_pointer, - random_values_dim_c0, - random_values_dim_c1, - random_values_dim_c2, - output_pointer, - output_dim_c0, - output_dim_c1, - output_dim_c2, - x_dim, - y_dim, - spike_dim, - h_dim, - number_of_pattern); - - return; -}; - -__device__ size_t gpu_lower_bound(float* __restrict__ data_ptr, - size_t data_length, - size_t data_ptr_stride, - float compare_to_value) -{ - - size_t start_of_range = 0; - size_t length_of_range = data_length; - - while (length_of_range != 0) - { - size_t half_length = length_of_range >> 1; - size_t actual_position = start_of_range + half_length; - - if (data_ptr[actual_position * data_ptr_stride] < compare_to_value) - { - start_of_range = ++actual_position; - length_of_range -= half_length + 1; - } - else - length_of_range = half_length; - } - return start_of_range; -}; - -__global__ void kernel_spike_generation( - float* __restrict__ input_pointer, - size_t input_dim_c0, - size_t input_dim_c1, - size_t input_dim_c2, - float* __restrict__ random_values_pointer, - size_t random_values_dim_c0, - size_t random_values_dim_c1, - size_t random_values_dim_c2, - int64_t* __restrict__ 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 max_threadable_tasks) -{ - int idx = threadIdx.x + blockIdx.x * blockDim.x; - - if (idx < max_threadable_tasks) - { - - size_t pattern_id = idx / spike_dim; - size_t position_spike = idx - (pattern_id * spike_dim); - - size_t position_x = blockIdx.y; - size_t position_y = blockIdx.z; - - float* p_ptr = input_pointer + pattern_id * input_dim_c0 + - position_x * input_dim_c2 + position_y; - - int64_t* out_ptr = output_pointer + pattern_id * output_dim_c0 + - position_x * output_dim_c2 + position_y - + position_spike * output_dim_c1; - - float* rand_ptr = random_values_pointer + - pattern_id * random_values_dim_c0 + - position_x * random_values_dim_c2 + position_y - + position_spike * random_values_dim_c1; - - *out_ptr = gpu_lower_bound(p_ptr, - h_dim, - input_dim_c1, - *rand_ptr); - } -}; - -void SpikeGenerationGPU::gpu_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 number_of_pattern) -{ - cudaError_t status; - assert((x_dim < 65535)); - assert((y_dim < 65535)); - - // ////////////////////////////////////// - // Calculate the distribution on the GPU - // ////////////////////////////////////// - - int min_grid_size; - int block_size; - int grid_size; - - size_t dynamic_s_mem_size = 0; - size_t max_threadable_tasks = number_of_pattern * spike_dim * x_dim * y_dim; - - status = cudaOccupancyMaxPotentialBlockSize(&min_grid_size, &block_size, - (void*)kernel_spike_generation, - dynamic_s_mem_size, max_threadable_tasks); - assert((status == cudaSuccess)); - - grid_size = ((number_of_pattern * spike_dim) + block_size - 1) / block_size; - - dim3 grid(grid_size, x_dim, y_dim); - - - kernel_spike_generation<<>>( - input_pointer, - input_dim_c0, - input_dim_c1, - input_dim_c2, - random_values_pointer, - random_values_dim_c0, - random_values_dim_c1, - random_values_dim_c2, - output_pointer, - output_dim_c0, - output_dim_c1, - output_dim_c2, - x_dim, - y_dim, - spike_dim, - h_dim, - (number_of_pattern * spike_dim)); - - cudaDeviceSynchronize(); - - return; -}; - - -void SpikeGenerationGPU::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 SpikeGenerationGPU::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_gpu_cpp/SpikeGenerationGPU.h b/network/spike_generation_gpu_cpp/SpikeGenerationGPU.h deleted file mode 100644 index 7059b4b..0000000 --- a/network/spike_generation_gpu_cpp/SpikeGenerationGPU.h +++ /dev/null @@ -1,70 +0,0 @@ -#ifndef SPIKEGENERATIONGPU -#define SPIKEGENERATIONGPU - -#include - -#include -#include - -class SpikeGenerationGPU -{ -public: - SpikeGenerationGPU(); - ~SpikeGenerationGPU(); - - 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); - - 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 gpu_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 number_of_pattern); - -}; - -#endif /* SPIKEGENERATIONGPU */