Optimisation d'un réseau neuronal
Bonsoir,
Je me re-penche actuellement sur la création d'un Spiking Neural Network ; tout fonctionne plutôt bien, ce qui m'amène donc à vouloir bosser au maximum l'optimisation de mes kernels CUDA.
Veuillez m'excuser par avance si je fais beaucoup de franglais, je n'ai littéralement rien trouvé en français (même pas de page Wikipédia!) sur ce type de NN, du coup, les termes sont ceux que j'ai trouvés dans les documentations anglophones.
J'ai choisi (ou plutôt, c'est la solution qui m'est venu naturellement) de séparer mes data en trois structures principales:
- t_neuron: qui est la représentation d'un neurone, elle contient pas mal de trucs, donc l'action potential, le threshold, le weight, et autre trucs dans le genre.
- t_synapse: qui contient un ID in vers une case de mon tableau de neurone, et un ID out qui fait la même ; ainsi que son axonal delay. -- mes synapses sont donc à sens unique.
- t_spike: qui symbolise une spike, se calquant sur une t_synapse, qui contient donc un ID synapse, étant une des case de mon tableau de synapse, un ID out, comme celui de t_synapse, pour ne pas avoir à passer le tableau de synapse à un endroit inutile, le timestep de début et de fin (début + axonal delay), sa value, et si elle est active ou non.
Pour prendre un cas concret, j'aurais un tableau de t_neuron de taille 20000, un tableau de t_synapse de taille 400000, et un tableau de t_spike de taille 400000 * 4. Oui, ça fait beaucoup :s. J'ai une sorte de "multiple buffer circulaire" où j'autorise chaque synapse à conduire 4 spikes en même temps, ce qui n'arrive jamais avec une configuration normale. Ca semble plutôt overkill côté mémoire, mais je pense que ça me permet de garder la mémoire 'coalesced', au moins à cet endroit, vu que syn[0] accède à spike[0-3], syn[1] à spike[4-7] etc... si j'ai bien compris l'idée.
Du coup, mon code est splitté en 3 kernels (4 théoriquement), le premier prend en parallèle tous les neurones de type INPUT, et regarde s'ils doivent faire feu ou non (selon le timestep actuel). Le deuxième prend toutes les synapses, et regarde si un spike doit être crée (AP >= threshold), si oui, elle s'exécute en cherchant un spike libre dans son mini-espace de 4 places (le pseudo buffer circulaire). Et le troisième et dernier prend tous les spikes, et si elle approche de la fin, update l'AP de la liaison post-synaptique, puis s'auto-détruit.
Pour être plus clair, voici un pseudo-code (qui est presque un vrai code, au final...) exprimant mieux mon point de vue.
Code:
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
|
| idx = blockIdx.x * blockDim.x + threadIdx.x
| = First Kernel =
| INPUT
| n -> le tableau contenant toutes mes neurones
| timestep -> le timestep actuel
| CODE
| if timestep >= n[idx].next_time
| n[idx].action_potential += n[idx].input_value
| n[idx].next_time += n[idx].input_time // next_time étant donc une valeur absolue
| = Second Kernel =
| INPUT
| n -> le tableau contenant toutes mes neurones
| s -> le tableau contenant toutes mes synapses
| sp -> le tableau contenant toutes mes spikes
| timestep -> le timestep actuel
| CODE
| if n[s[idx].id_in].action-potential < n[s[idx].id_in].threshold)
| return
| n[s[idx].id_in].carry = 1 // le flag sert à dire qu'on doit reset la variable, ce qui est fait dans le kernel 2.5
| for i = idx * 4 to i < idx * 4 + 4 step 1 // mon pseudo buffer circulaire où chaque thread/synapse a 4 slots
| if sp[i].active
| continue ;
| // INIT d'une nouvelle spike en sp[i] //
| break ;
|
| = Second Kernel Bis = (prend tous les neurones en parallèle)
| INPUT
| n -> le tableau contenant toutes mes neurones
| CODE
| if n[idx].carry
| n[idx].action_potential = 0.0f
| n[idx].carry = 0
| = Third Kernel =
| INPUT
| n -> le tableau contenant toutes mes neurones
| sp -> le tableau contenant toutes mes spikes
| timestep -> le timestep actuel
| CODE
| for i = idx * 4 to i < idx * 4 + 4 step 1 // de nouveau mon pseudo buffer circulaire où chaque thread/synapse a 4 slots
| if not sp[i].active
| continue ;
| if timestep >= sp[i].end_time
| sp[i].active = false
| n[sp[i].id_out].action_potential += (sp[i].value * n[sp][i].id_out].weight);
| if n[sp[i].id_out].action_potential < 0.0f
| n[sp[i].id_out].action_potential = 0.0f |
Je m'aperçois en me relisant qu'à chaque fois que j'essaie de mette du contexte, ça devient absolument obscure.
Enfin bon, si l'un de vous s'est déjà amusé à faire ça, où remarque à vu d'oeil une mauvais façon d'utiliser CUDA de façon context-free, je vous remercie d'avance pour votre aide précieuse! :)
Bonne soirée.
