<commit_before>//	CUDA code for SPIKE simulator
//
//
//	Author: <NAME>
//	Date: 9/12/2015

//  Adapted by <NAME> and <NAME>
//	Date: 23/3/2016

// For files/manipulations
#include <iostream>
#include <fstream>
#include <stdio.h>
#include <stdlib.h>
#include <algorithm> // For random shuffle
using namespace std;


#include "CUDAcode.h"
#include "Structs.h"
#include <time.h>
#include "CUDAErrorCheckHelpers.h"
// Silences the printfs
//#define QUIETSTART


// CUDA Spiking Simulator Run
//	INPUT (incomplete):
//			total_number_of_neurons = Total number of neurons in simulation
//			total_number_of_connections = Total number of Synapses
//			presynaptic_neuron_indices = vector- entry corresponds to pre-syn neuron ID
//			postsynaptic_neuron_indices = vector- entry corresponds to post-syn neuron ID
//			delays = vector- entry corresponds to axonal delay of syn
//			weight = vector- entry correponds to synaptic weight
//			spikes = vector- entry corresponds to countdown of spike on syn
//			stdp = vector- entry corresponding to whether or not syn is stdp
//			lastactive = vector- indicating last time synapse emitted current
//			parameters=
//			w_max = stdp parameter
//			a_minus = stdp parameter
//			a_plus = stdp parameter
//			tau_minus = stdp parameter
//			tau_plus = stdp parameter
//			timestep = timestep of the simulation
//			total_time_per_epoch = total time for which the sim should run per epoch
void GPUDeviceComputation (
					Neurons * neurons,
					Connections * connections,

					float total_time_per_epoch,
					int number_of_epochs,
					float timestep,
					bool save_spikes,

					int numStimuli,
					int* numEntries,
					int** genids,
					float** gentimes,
					bool randomPresentation
					){


	// Useful to have local versions of totals as used a lot
	size_t total_number_of_neurons = neurons->total_number_of_neurons;
	size_t total_number_of_connections = connections->total_number_of_connections;


	// Creating the Device Pointers I need
	int* d_genids;
	float* d_gentimes;
	
	// I need a location in which to store spikes
	int* d_tempstorenum;
	int* d_tempstoreID;
	float* d_tempstoretimes;
	// And some on host for copy back
	int* h_tempstoreID;
	float* h_tempstoretimes;
	// This is the location in which they shall occasionally be put
	int* h_tempspikenum;
	int h_spikenum = 0;
	int* h_spikestoreID = NULL;
	float* h_spikestoretimes = NULL;


	// For saving spikes (Make seperate class)
	CudaSafeCall(cudaMalloc((void **)&d_tempstoreID, sizeof(int)*total_number_of_neurons));
	CudaSafeCall(cudaMalloc((void **)&d_tempstoretimes, sizeof(float)*total_number_of_neurons));
	CudaSafeCall(cudaMalloc((void **)&d_tempstorenum, sizeof(int)));

	// Send data to device: data for saving spikes
	CudaSafeCall(cudaMemset(d_tempstoreID, -1, sizeof(int)*total_number_of_neurons));
	CudaSafeCall(cudaMemset(d_tempstoretimes, -1.0f, sizeof(float)*total_number_of_neurons));
	CudaSafeCall(cudaMemset(d_tempstorenum, 0, sizeof(int)));


	neurons->initialise_device_pointers();
	connections->initialise_device_pointers();


	// Allocate host data
	h_tempspikenum = (int*)malloc(sizeof(int));
	h_tempstoreID = (int*)malloc(sizeof(int)*total_number_of_neurons);
	h_tempstoretimes = (float*)malloc(sizeof(float)*total_number_of_neurons);
	h_tempspikenum[0] = 0;



	// THREADS&BLOCKS
	// The number of threads per block I shall keep held at 128
	int threads = 128;
	connections->set_threads_per_block_and_blocks_per_grid(threads);
	neurons->set_threads_per_block_and_blocks_per_grid(threads);


	dim3 threadsPerBlock(threads,1,1);
	// I now have to calculate the number of blocks ....
	int number_of_connection_blocks = (total_number_of_connections + threads) / threads; //Previously number_of_connection_blocks
	int vectorblocknum = (total_number_of_neurons + threads) / threads;

	// The maximum dimension for the grid is 65535
	dim3 number_of_connection_blocks_per_grid(number_of_connection_blocks,1,1); // Previously connblocksPerGrid
	printf("number_of_connection_blocks_per_grid: %d\n", number_of_connection_blocks_per_grid.x);
	dim3 vectorblocksPerGrid(vectorblocknum,1,1);  
	// Temp Values which will be replaced

	int genblocknum = 1;
	dim3 genblocksPerGrid(genblocknum,1,1);


	// RANDOM NUMBERS
	// Create the random state seed trackers
	curandState_t* states;
	cudaMalloc((void**) &states, total_number_of_neurons*sizeof(curandState_t));
	// Initialise the random states
	init<<<threadsPerBlock, vectorblocksPerGrid>>>(42, states, total_number_of_neurons);
	CudaCheckError();
	// Keep space for the random numbers
	float* gpu_randfloats;
	CudaSafeCall(cudaMalloc((void**) &gpu_randfloats, total_number_of_neurons*sizeof(float)));
	// REQUIRED DATA SPACES
	float* currentinjection;
	CudaSafeCall(cudaMalloc((void**)&currentinjection, total_number_of_neurons*sizeof(float)));
	// Variables necessary
	clock_t begin = clock();

	// Poisson number
	int numPoisson = 0;
	for (int i = 0; i < total_number_of_neurons; i++){
		if (neurons->group_parameters[i].rate != 0.0f){
			++numPoisson;<commit_msg>number_of_connection_blocks_per_grid calculation moved from CUDACode<commit_after>//	CUDA code for SPIKE simulator
//
//
//	Author: <NAME>
//	Date: 9/12/2015

//  Adapted by <NAME> and <NAME>
//	Date: 23/3/2016

// For files/manipulations
#include <iostream>
#include <fstream>
#include <stdio.h>
#include <stdlib.h>
#include <algorithm> // For random shuffle
using namespace std;


#include "CUDAcode.h"
#include "Structs.h"
#include <time.h>
#include "CUDAErrorCheckHelpers.h"
// Silences the printfs
//#define QUIETSTART


// CUDA Spiking Simulator Run
//	INPUT (incomplete):
//			total_number_of_neurons = Total number of neurons in simulation
//			total_number_of_connections = Total number of Synapses
//			presynaptic_neuron_indices = vector- entry corresponds to pre-syn neuron ID
//			postsynaptic_neuron_indices = vector- entry corresponds to post-syn neuron ID
//			delays = vector- entry corresponds to axonal delay of syn
//			weight = vector- entry correponds to synaptic weight
//			spikes = vector- entry corresponds to countdown of spike on syn
//			stdp = vector- entry corresponding to whether or not syn is stdp
//			lastactive = vector- indicating last time synapse emitted current
//			parameters=
//			w_max = stdp parameter
//			a_minus = stdp parameter
//			a_plus = stdp parameter
//			tau_minus = stdp parameter
//			tau_plus = stdp parameter
//			timestep = timestep of the simulation
//			total_time_per_epoch = total time for which the sim should run per epoch
void GPUDeviceComputation (
					Neurons * neurons,
					Connections * connections,

					float total_time_per_epoch,
					int number_of_epochs,
					float timestep,
					bool save_spikes,

					int numStimuli,
					int* numEntries,
					int** genids,
					float** gentimes,
					bool randomPresentation
					){


	// Useful to have local versions of totals as used a lot
	size_t total_number_of_neurons = neurons->total_number_of_neurons;
	size_t total_number_of_connections = connections->total_number_of_connections;


	// Creating the Device Pointers I need
	int* d_genids;
	float* d_gentimes;
	
	// I need a location in which to store spikes
	int* d_tempstorenum;
	int* d_tempstoreID;
	float* d_tempstoretimes;
	// And some on host for copy back
	int* h_tempstoreID;
	float* h_tempstoretimes;
	// This is the location in which they shall occasionally be put
	int* h_tempspikenum;
	int h_spikenum = 0;
	int* h_spikestoreID = NULL;
	float* h_spikestoretimes = NULL;


	// For saving spikes (Make seperate class)
	CudaSafeCall(cudaMalloc((void **)&d_tempstoreID, sizeof(int)*total_number_of_neurons));
	CudaSafeCall(cudaMalloc((void **)&d_tempstoretimes, sizeof(float)*total_number_of_neurons));
	CudaSafeCall(cudaMalloc((void **)&d_tempstorenum, sizeof(int)));

	// Send data to device: data for saving spikes
	CudaSafeCall(cudaMemset(d_tempstoreID, -1, sizeof(int)*total_number_of_neurons));
	CudaSafeCall(cudaMemset(d_tempstoretimes, -1.0f, sizeof(float)*total_number_of_neurons));
	CudaSafeCall(cudaMemset(d_tempstorenum, 0, sizeof(int)));


	neurons->initialise_device_pointers();
	connections->initialise_device_pointers();


	// Allocate host data
	h_tempspikenum = (int*)malloc(sizeof(int));
	h_tempstoreID = (int*)malloc(sizeof(int)*total_number_of_neurons);
	h_tempstoretimes = (float*)malloc(sizeof(float)*total_number_of_neurons);
	h_tempspikenum[0] = 0;



	// THREADS&BLOCKS
	// The number of threads per block I shall keep held at 128
	int threads = 128;
	connections->set_threads_per_block_and_blocks_per_grid(threads);
	neurons->set_threads_per_block_and_blocks_per_grid(threads);


	dim3 threadsPerBlock(threads,1,1);
	// I now have to calculate the number of blocks ....
	int number_of_connection_blocks = (total_number_of_connections + threads) / threads; //Previously number_of_connection_blocks
	int vectorblocknum = (total_number_of_neurons + threads) / threads;

	// The maximum dimension for the grid is 65535
	dim3 vectorblocksPerGrid(vectorblocknum,1,1);  
	// Temp Values which will be replaced

	int genblocknum = 1;
	dim3 genblocksPerGrid(genblocknum,1,1);


	// RANDOM NUMBERS
	// Create the random state seed trackers
	curandState_t* states;
	cudaMalloc((void**) &states, total_number_of_neurons*sizeof(curandState_t));
	// Initialise the random states
	init<<<threadsPerBlock, vectorblocksPerGrid>>>(42, states, total_number_of_neurons);
	CudaCheckError();
	// Keep space for the random numbers
	float* gpu_randfloats;
	CudaSafeCall(cudaMalloc((void**) &gpu_randfloats, total_number_of_neurons*sizeof(float)));
	// REQUIRED DATA SPACES
	float* currentinjection;
	CudaSafeCall(cudaMalloc((void**)&currentinjection, total_number_of_neurons*sizeof(float)));
	// Variables necessary
	clock_t begin = clock();

	// Poisson number
	int numPoisson = 0;
	for (int i = 0; i < total_number_of_neurons; i++){
		if (neurons->group_parameters[i].rate != 0.0f){
			++numPoisson;
{'fraction_non_alphanumeric': 0.10041293181589284, 'fraction_numerical': 0.00997079262765636, 'mean_word_length': 4.7498552403011, 'pattern_counts': {'":': 0, '<': 29, '<?xml version=': 0, '>': 41, 'https://': 0, 'lorem ipsum': 0, 'www.': 0, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 0, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '1714882', 'n_tokens_mistral': 3642, 'n_tokens_neox': 3321, 'n_words': 862}