GSoC 2022 Project Idea 4.1: ISPC backends for GeNN (350 h)

GeNN is a C++ library that generates code for efficiently simulating Spiking Neural Networks using GPUs. Currently, GeNN generates CUDA, OpenCL code meaning that it can target a wide range of GPUs . However, for simulating smaller models, focused targeting of SIMD CPUs may be advantageous. For this project you will develop a new GeNN code-generation backend for ISPC (http://ispc.github.io/), which targets the SIMD units in a wide range of modern CPUs using a CUDA-like programming model.

Skills required: C++, Single Instruction Multiple Thread (SIMT) programming

Mentors: Jamie Knight @jamie (J.C.Knight@sussex.ac.uk) and Thomas Nowotny @tnowotny (t.nowotny@sussex.ac.uk)

Tech keywords: C++, SIMT, GPU

For anyone interested in this project, good first steps would be getting familiar with ISPC via the materials linked to from http://ispc.github.io/ and with using GeNN from C++ via the documentation and tutorials here.

If you feel like digging a little deeper in the following areas might be good preparation for writing your application:

Alternatively, if you fancy having a go at a pre-project task, try porting this GeNN-like CUDA code to ISPCC (the CUDA version can be compiled with nvcc test.cc -o test) - the results can be plotted with this script.

Hello, I am using python 3.6.9 on my linux mint cinnamon on a laptop with no nvidia gpu and have intel i5 with amd radeon.while trying to get pygenn binary 4.7 wheel am getting this error:–
Requirement ‘pygenn-4.7.0-cp36-cp36m-linux_x86_64.whl’ looks like a filename, but the file does not exist. pygenn-4.7.0-cp36-cp36m-linux_x86_64.whl is not a supported wheel on this platform.
Is there a way to use and get aquainted with genn from c++ on my system as i believe cuda is necessary.

Sadly the compatibility of the wheel files isn’t very good so it sounds like you’ll need to install GeNN from source. The wheels are also only for python development - if you want to use GeNN from C++ or modify it, you should download/checkout the source

I’ve followed the steps mentioned in GeNN: Installation for pygenn from source just like @jamie suggested.I’ve successfully built dynamic library version of GeNN, directly into the PyGeNN dir as the 1st step.While following the second step i.e. “Build the Python extension with setup tools” I used python setup.py develop and ran into an error namely
Traceback (most recent call last):
File “setup.py”, line 32, in
wsl = “microsoft” in uname().release
AttributeError: ‘tuple’ object has no attribute ‘release’

Is there some problem in uname() or am I doing some mistake.Please provide correction measures.

What version of Python are you using? Based on the documentation for that function and a quick test on various systems, it seems like you must be using Python < 3.3. GeNN 4.X does technically still support Python 2.7 so, if you have a strong reason to still be using this, the uname_2_7 branch should fix your issue.

@jamie Actually I have 2 versions of python which are 2.7.17 AND 3.6.9
I were using the command “python setup.py develop” where I got previous error.
Now I’ve used “sudo python3 setup.py develop” instead so now I am getting new error i.e.
unable to execute ‘swig’: No such file or directory
error: command ‘swig’ failed with exit status 1
However, there is actually a swig directory present which contains files a few *.i files
What might be the cause

Sorry to disturb but installing swig resolves the problem.

great, that would have been what I suggested :slight_smile:

I have been following the c++ tutorial for creating first simulation(tenHHModel) and I’ve successfully built and ran the simulation. While trying to read , by copying intermediate values from the device and save them into a file. Just as mentioned using this code ;-
std::ofstream os(“tenHH_output.V.dat”);
while (t < 1000.0f) {
stepTime();
pullVPop1FromDevice();
os << t << " ";
for (int j= 0; j < 10; j++) {
os << VPop1[j] << " ";
}
os << std::endl;
}
os.close();

After updating the tenHHModel simulation code while trying to build using “make” I get the following error;-
error: variable ‘std::ofstream os’ has initializer but incomplete type
std::ofstream os(“tenHH_output.V.dat”);

Any solutions please…

Sounds like you’re missing a #include<fstream>

Thank you, it worked and now that I have the tenHH.V.dat file I am wondering how to plot the column one (time) against the subsequent 10 columns .Should I use matpoltlib? If yes, then how? I believe the steps would be something like extracting and plotting like;-
import matplotlib.pyplot as plt
import csv

X = []
Y = []

with open(‘tenHH_output.V.dat’, ‘r’) as datafile:
plotting = csv.reader(datafile, delimiter=’ ')

for ROWS in plotting:
    X.append(int(ROWS[0]))
    Y.append(int(ROWS[1]))

plt.plot(X, Y)
plt.title(‘Line Graph using CSV’)
plt.xlabel(‘X’)
plt.ylabel(‘Y’)
plt.show()

Am I missing something or is this correct ?

It’s totally up to you - you could open it in excel/open office or load it into Matlab or Python. If you’re going to use Python, personally, I would use the numpy.loadtxt function.

While trying to port this “genn like cuda code” to ispcc (Sample GeNN-like CUDA code for simulating Izhikevich neurons · GitHub)
Sample GeNN-like CUDA code for simulating Izhikevich neurons · GitHub
I am having a few doubts like ;–

  1. Does the Ispc code need error checking mechanism like CUDA_CHECK_ERRORS which is useful for checking proper memory allocation on device/host.
  2. I believe global qualifier is for compiler to know to run function on device and __syncthreads() is for barrier synchronization among instances which are probably not needed in ispcc so is there something I am missing.?
  3. I think threadidx in cuda c is equivalent to programIndex in ispcc as cuda uses threads in blocks whereas ispc uses program instances in gangs however I am having trouble porting this statement (const unsigned int id = 32 * blockIdx.x + threadIdx.x;). My thought on this is that 1 block contains 32 threads and each id is for unique thread but in ispcc I would use const unsigned int id = programIndex; Will this be correct??.
  4. Should I use “uniform” keyword for declaring global variables like glbSpkCntNeurons and d_glbSpkNeurons as these will be shared across the gang . Also shared keyword is used in cuda c to to make variables reside in shared memory easing communication between threads in the same block , is some similar mechanism possible in ispcc or should I ignore it ??
  5. atomicAdd like atomic operations used in cuda C is good for operating safely on memory without being affected by other threads. Its equivalents in ispcc are atomic_add_global and atomic_add_local. Which one should be used in porting code ??

Will someone please clarify?.

Hello,
I am Mohamed from Egypt, A fresh grad computer systems engineer.
I have a solid understanding of compilers.
I fixed some issues in the LLVM middle-end, and I added most of OMPD(OpenMP Debugging library) support to GCC, I also fixed issues in SpiderMonkey.

I am very interested in this project but I can’t access GPU.
Is it possible to work on this project?

Thanks

Hi Mohamed,

Thanks for your interest in this project, it sounds like you have the kind of skills this project would require. However, this neurostars is for a 2022 project so currently isn’t confirmed for this year. I would happily propose it again though and when/if it gets confirmed I’ll let you know. No GPU access should be required for this project.

Jamie

Hi Jamie,
Thanks for replying.
Until the confirmation I will build GeNN, and try to fix some issues.
Please suggest some issues that may help in the project.
In case if the project is not available, I will happily work on it,
if someone are available to mentor.

Thanks
Mohamed

I wouldn’t worry too much about fixing GeNN issues right now - a new major version is going to be coming out very soon and I don’t really want to have to merge any more PRs! If you do want to do something, I recommend thinking about porting the code linked in the first post to ISPC. Things I discovered are:

  • The ISPC compiler only compiles ISPC so you’ll have to split the kernels from the boilerplate code
  • You should mark some variables as uniform

As a starting point, I also read the GeNN tutorial and ISPC Documentation.
I noticed what you mentioned

  • Yes ISPC Only compiles ISPC.
  • uniform keyword is used to mark variables as shared.

For the part of porting Cuda to ISPC, does this mean porting it manually?
or with some generator.

I am reading the ISPC manual and trying to convert the program above.
I reached the following Cuda/ISPC conversion:
__shared__ —> uniform for shared variables that need synchronization
__syncthreads —> memory_brrier
__global__ —> kernel that should be in a standalone file
atomicAdd —> atomic_add_* operations

BTW should I read the GeNN paper?