Delete network/spike_generation_gpu_cpp directory
This commit is contained in:
David Rotermund
GitHub
parent
ba69190204
commit
8983d63330
4 changed files with 0 additions and 397 deletions
|
|
@ -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)
|
||||
|
||||
|
|
@ -1,19 +0,0 @@
|
|||
|
||||
#include <pybind11/pybind11.h>
|
||||
|
||||
#include "SpikeGenerationGPU.h"
|
||||
|
||||
namespace py = pybind11;
|
||||
|
||||
PYBIND11_MODULE(PySpikeGenerationGPU, m)
|
||||
{
|
||||
m.doc() = "SpikeGenerationGPU Module";
|
||||
py::class_<SpikeGenerationGPU>(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);
|
||||
}
|
||||
|
|
@ -1,275 +0,0 @@
|
|||
#include "SpikeGenerationGPU.h"
|
||||
|
||||
#include <omp.h>
|
||||
#include <stdio.h>
|
||||
#include <string.h>
|
||||
|
||||
#include <algorithm>
|
||||
#include <cassert>
|
||||
#include <iostream>
|
||||
|
||||
|
||||
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<<<grid, block_size >>>(
|
||||
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;
|
||||
};
|
||||
|
|
@ -1,70 +0,0 @@
|
|||
#ifndef SPIKEGENERATIONGPU
|
||||
#define SPIKEGENERATIONGPU
|
||||
|
||||
#include <unistd.h>
|
||||
|
||||
#include <cctype>
|
||||
#include <iostream>
|
||||
|
||||
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 */
|
||||
Loading…
Reference in a new issue