Compare commits

...

2 Commits

Author SHA1 Message Date
56a573b215 Update test/cnn_convolution.cu 2023-11-10 23:14:10 +01:00
37771f76ed Update tests 2023-11-10 23:09:48 +01:00
12 changed files with 125 additions and 142 deletions

View File

@ -220,7 +220,7 @@ endif
#
run-tests: build-tests
$(foreach file, $(wildcard $(TEST_SRCDIR)/*.sh), $(file);)
@echo "$$(for file in build/test-*; do echo -e \\033[33m#####\\033[0m $$file \\033[33m#####\\033[0m; $$file || echo -e "\\033[1m\\033[31mErreur sur $$file\\033[0m"; done)"
@echo "$$(for file in build/test-*; do $$file || exit 1; done)"
build-tests: prepare-tests $(TESTS_OBJ) $(BUILDDIR)/test-cnn_matrix_multiplication $(BUILDDIR)/test-cnn_convolution $(BUILDDIR)/test-cuda_memory_management

View File

@ -6,6 +6,8 @@
#include <time.h>
#include <time.h>
#include "include/test.h"
#include "../src/common/include/memory_management.h"
#include "../src/cnn/include/convolution.h"
#include "../src/common/include/colors.h"
@ -93,7 +95,7 @@ bool check_matrices_equality(float*** m1, float*** m2, int n, int p, int q, int
return true;
}
void run_convolution_test(int input_width, int output_width, int rows, int columns) {
bool run_convolution_test(int input_width, int output_width, int rows, int columns) {
assert(input_width >= output_width);
int k_size = input_width - output_width +1;
@ -142,10 +144,7 @@ void run_convolution_test(int input_width, int output_width, int rows, int colum
// Vérification de l'égalité des matrices
printf("(%d, %d, %d, %d) Checking equality.\n", rows, columns, input_width, output_width);
if (!check_matrices_equality(output_gpu, output_cpu, kernel->columns, output_width, output_width, kernel->k_size)) {// TODO: change acceptation
exit(1);
}
printf(GREEN "OK\n" RESET);
bool valid = check_matrices_equality(output_gpu, output_cpu, kernel->columns, output_width, output_width, kernel->k_size);
free_matrix(kernel->bias, kernel->columns, output_width);
@ -157,10 +156,14 @@ void run_convolution_test(int input_width, int output_width, int rows, int colum
free_matrix(input, kernel->rows, input_width);
free_matrix(output_cpu, kernel->columns, output_width);
free_matrix(output_gpu, kernel->columns, output_width);
return valid;
}
int main() {
_TEST_PRESENTATION("Cuda Convolution");
printf("Checking CUDA compatibility.\n");
bool cuda_compatible = cuda_setup(true);
if (!cuda_compatible) {
@ -171,9 +174,9 @@ int main() {
srand(clock());
run_convolution_test(20, 15, 30, 40);
run_convolution_test(30, 25, 40, 50);
run_convolution_test(250, 200, 3, 3);
_TEST_ASSERT(run_convolution_test(20, 15, 30, 40), "Small");
_TEST_ASSERT(run_convolution_test(30, 25, 40, 50), "Medium");
_TEST_ASSERT(run_convolution_test(250, 200, 3, 3), "Big");
return 0;
}

View File

@ -2,6 +2,8 @@
#include <assert.h>
#include <stdio.h>
#include "include/test.h"
#include "../src/common/include/memory_management.h"
#include "../src/common/include/colors.h"
#include "../src/common/include/utils.h"
@ -25,7 +27,6 @@ __global__ void local_kernel(funcPtr f, float*** input, int depth, int rows, int
void test1(int activation, bool use_local_kernel) {
printf("Test sur la fonction %d\n", activation);
printf("\tInitialisation OK\n");
// Initialise values
int depth = 10;
int rows = 10;
@ -45,7 +46,6 @@ void test1(int activation, bool use_local_kernel) {
}
}
}
printf("\t" GREEN "OK\n" RESET);
funcPtr func_cpu = get_activation_function(activation);
@ -64,16 +64,13 @@ void test1(int activation, bool use_local_kernel) {
gpuErrchk( cudaPeekAtLastError() );
gpuErrchk( cudaDeviceSynchronize() );
}
printf("\t" GREEN "OK\n" RESET);
printf("\tVérification des résultats\n");
for (int i=0; i < depth; i++) {
for (int j=0; j < rows; j++) {
for (int k=0; k < columns; k++) {
if (fabs((*func_cpu)(input_initial[i][j][k]) - input[i][j][k]) > 1e-6) {
printf_error((char*)"Les résultats ne coincident pas\n");
printf("Différence %e\n", fabs((*func_cpu)(input_initial[i][j][k]) - input[i][j][k]));
exit(1);
_TEST_ASSERT(false, "Coïncidence des résultats");
//printf("Différence %e\n", fabs((*func_cpu)(input_initial[i][j][k]) - input[i][j][k]));
}
}
gree(input[i][j], false);
@ -84,12 +81,11 @@ void test1(int activation, bool use_local_kernel) {
}
gree(input, false);
free(input_initial);
printf("\t" GREEN "OK\n" RESET);
printf(GREEN "OK\n" RESET);
}
int main() {
_TEST_PRESENTATION("Cuda fonctions");
printf("Checking CUDA compatibility.\n");
bool cuda_compatible = cuda_setup(true);
if (!cuda_compatible) {

View File

@ -3,6 +3,8 @@
#include <stdint.h>
#include <inttypes.h>
#include "include/test.h"
#include "../src/common/include/colors.h"
#include "../src/cnn/include/neuron_io.h"
#include "../src/cnn/include/creation.h"
@ -12,35 +14,25 @@
int main() {
printf("Création du réseau\n");
_TEST_PRESENTATION("CNN Lecture/Écriture")
Network* network = create_network_lenet5(0, 0, 3, GLOROT, 32, 1, 2); // Pas besoin d'initialiser toute la backprop
printf(GREEN "OK\n" RESET);
_TEST_ASSERT(true, "Création du réseau");
printf("Écriture du réseau\n");
write_network((char*)".test-cache/cnn_neuron_io.bin", network);
printf(GREEN "OK\n" RESET);
_TEST_ASSERT(true, "Écriture du réseau");
printf("Vérification de l'accès en lecture\n");
Network* network2 = read_network((char*)".test-cache/cnn_neuron_io.bin");
Network* network3 = read_network((char*)".test-cache/cnn_neuron_io.bin");
printf(GREEN "OK\n" RESET);
_TEST_ASSERT(true, "Vérification de l'accès en lecture");
printf("Vérification de l'égalité des réseaux\n");
if (! equals_networks(network, network2)) {
printf_error(RED "Le réseau lu ne contient pas les mêmes données.\n" RESET);
exit(1);
}
if (! equals_networks(network2, network3)) {
printf_error(RED "La lecture du réseau donne des résultats différents.\n" RESET);
exit(1);
}
printf(GREEN "OK\n" RESET);
_TEST_ASSERT(equals_networks(network, network2), "Égalité des réseaux");
_TEST_ASSERT(equals_networks(network2, network3), "Égalité de deux lectures");
printf("Libération de la mémoire\n");
free_network(network);
free_network(network2);
free_network(network3);
printf(GREEN "OK\n" RESET);
_TEST_ASSERT(true, "Libération de la mémoire");
return 0;
}

View File

@ -10,12 +10,14 @@
#include "../src/cnn/include/free.h"
#include "../src/cnn/include/cnn.h"
#include "include/test.h"
int main() {
_TEST_PRESENTATION("Description Architecture LeNet5");
Kernel* kernel;
printf("Création du réseau\n");
Network* network = create_network_lenet5(0, 0, 3, 2, 32, 1, NN_ONLY); // Pas besoin d'initialiser toute la backprop
printf(GREEN "OK\n" RESET);
_TEST_ASSERT(true, "Création du réseau");
printf("Architecture LeNet5:\n");
for (int i=0; i < network->size-1; i++) {
@ -48,11 +50,9 @@ int main() {
printf("activation: %d\n", kernel->activation);
}
}
printf("\n" GREEN "OK\n" RESET);
printf("Libération de la mémoire\n");
free_network(network);
printf(GREEN "OK\n" RESET);
_TEST_ASSERT(true, "Libération de la mémoire");
return 0;
}

View File

@ -1,6 +1,8 @@
#include <stdlib.h>
#include <stdio.h>
#include "include/test.h"
#include "../src/common/include/colors.h"
#include "../src/cnn/include/creation.h"
#include "../src/cnn/include/models.h"
@ -8,32 +10,21 @@
#include "../src/cnn/include/free.h"
int main() {
printf("Création du réseau\n");
_TEST_PRESENTATION("Utilitaires du CNN");
Network* network = create_network_lenet5(0, 0, 3, 2, 32, 1, 0);
Network* network2 = create_network_lenet5(0, 0, 3, 2, 32, 1, 0);
printf(GREEN "OK\n" RESET);
_TEST_ASSERT(true, "Création de réseaux");
printf("Copie du réseau via copy_network\n");
Network* network_cp = copy_network(network);
printf(GREEN "OK\n" RESET);
_TEST_ASSERT(true, "Copie de réseau (copy_network)");
printf("Vérification de l'égalité des réseaux\n");
if (! equals_networks(network, network_cp)) {
printf_error(RED "Les deux réseaux obtenus ne sont pas égaux.\n" RESET);
exit(1);
}
printf(GREEN "OK\n" RESET);
_TEST_ASSERT(equals_networks(network, network_cp), "Égalité du réseau copié (copy_network)");
printf("Copie du réseau via copy_network_parameters\n");
copy_network_parameters(network, network2);
printf(GREEN "OK\n" RESET);
_TEST_ASSERT(true, "Copie de réseau (copy_network_parameters)");
printf("Vérification de l'égalité des réseaux\n");
if (! equals_networks(network, network2)) {
printf_error(RED "Les deux réseaux obtenus ne sont pas égaux.\n" RESET);
exit(1);
}
printf(GREEN "OK\n" RESET);
_TEST_ASSERT(equals_networks(network, network2), "Égalité du réseau copié (copy_network_parameters)");
free_network(network_cp);
free_network(network2);

View File

@ -2,17 +2,18 @@
#include <stdio.h>
#include <inttypes.h>
#include "include/test.h"
#include "../src/common/include/colors.h"
#include "../src/common/include/mnist.h"
void read_test(int nb_images, int width, int height, int*** images, unsigned int* labels) {
printf("\tLecture des labels\n");
for (int i=0; i < nb_images; i++) {
(void)labels[i];
}
printf(GREEN "\tOK\n" RESET);
printf("\tLecture des images\n");
_TEST_ASSERT(true, "Accès en lecture des labels");
for (int i=0; i < nb_images; i++) {
for (int j=0; j < height; j++) {
for (int k=0; k < width; k++) {
@ -20,35 +21,30 @@ void read_test(int nb_images, int width, int height, int*** images, unsigned int
}
}
}
printf(GREEN "\tOK\n" RESET);
_TEST_ASSERT(true, "Accès en lecture des images");
}
int main() {
_TEST_PRESENTATION("Mnist: Accès en lecture");
char* image_file = (char*)"data/mnist/t10k-images-idx3-ubyte";
char* labels_file = (char*)"data/mnist/t10k-labels-idx1-ubyte";
printf("Chargement des paramètres\n");
int* parameters = read_mnist_images_parameters(image_file);
int nb_images = parameters[0];
int height = parameters[1];
int width = parameters[2];
printf(GREEN "OK\n" RESET);
printf("Chargement des images\n");
_TEST_ASSERT(true, "Chargement des paramètres");
int*** images = read_mnist_images(image_file);
printf(GREEN "OK\n" RESET);
printf("Chargement des labels\n");
_TEST_ASSERT(true, "Chargement des images");
unsigned int* labels = read_mnist_labels(labels_file);
printf(GREEN "OK\n" RESET);
printf("Vérification de l'accès en lecture\n");
_TEST_ASSERT(true, "Chargement des labels");
read_test(nb_images, width, height, images, labels);
printf(GREEN "OK\n" RESET);
for (int i=0; i < nb_images; i++) {
for (int j=0; j < height; j++) {
free(images[i][j]);

View File

@ -3,21 +3,24 @@
#include <stdint.h>
#include <inttypes.h>
#include "include/test.h"
#include "../src/dense/include/neural_network.h"
#include "../src/dense/include/neuron_io.h"
#include "../src/common/include/colors.h"
int main() {
printf("Création du réseau\n");
_TEST_PRESENTATION("Dense: Création")
Network* network = (Network*)malloc(sizeof(Network));
int tab[5] = {30, 25, 20, 15, 10};
network_creation(network, tab, 5);
printf(GREEN "OK\n" RESET);
_TEST_ASSERT(true, "Création");
printf("Initialisation du réseau\n");
network_initialisation(network);
printf(GREEN "OK\n" RESET);
_TEST_ASSERT(true, "Initialisation");
deletion_of_network(network);
_TEST_ASSERT(true, "Suppression");
return 0;
}

View File

@ -4,6 +4,8 @@
#include <stdint.h>
#include <inttypes.h>
#include "include/test.h"
#include "../src/dense/include/neural_network.h"
#include "../src/dense/include/neuron_io.h"
#include "../src/common/include/colors.h"
@ -64,22 +66,20 @@ Network* create_network(int nb_layers, int nb_max_neurons, int nb_min_neurons) {
}
int main() {
printf("Création du réseau\n");
_TEST_PRESENTATION("Dense: Lecture/Écriture")
Network* network = create_network(5, 300, 10);
printf(GREEN "OK\n" RESET);
_TEST_ASSERT(true, "Création du réseau");
printf("Écriture du réseau\n");
write_network((char*)".test-cache/neuron_io.bin", network);
printf(GREEN "OK\n" RESET);
printf("Vérification de l'accès en lecture\n");
_TEST_ASSERT(true, "Écriture du réseau");
Network* network2 = read_network((char*)".test-cache/neuron_io.bin");
printf(GREEN "OK\n" RESET);
printf("Suppression des réseaux\n");
_TEST_ASSERT(true, "Accès en lecture");
deletion_of_network(network);
deletion_of_network(network2);
printf(GREEN "OK\n" RESET);
_TEST_ASSERT(true, "Suppression des réseaux");
return 0;
}

23
test/include/test.h Normal file
View File

@ -0,0 +1,23 @@
#ifndef DEF_TEST_H
#define DEF_TEST_H
#include <stdio.h>
#include <stdlib.h>
#include <stdbool.h>
#include "../../src/common/include/colors.h"
#define _TEST_PRESENTATION(description) { printf("\n" BLUE "#### %s:" BOLD "%s" RESET BLUE " #####" RESET "\n", __FILE__, description); }
#define _TEST_ASSERT(condition, description) { \
if (condition) { \
printf("[" GREEN "OK" RESET "] %s:%d: %s\n", __FILE__, __LINE__, description); \
} else { \
printf("[" RED "ERREUR" RESET "] %s:%d: %s\n", __FILE__, __LINE__, description); \
exit(1); \
} \
}
#endif

View File

@ -2,13 +2,16 @@
#include <stdio.h>
#include <assert.h>
#include "include/test.h"
#include "../src/common/include/memory_management.h"
#include "../src/common/include/colors.h"
#define N 350
int main() {
printf("Pollution de la mémoire\n");
_TEST_PRESENTATION("Memory management (C part)")
int mem_used;
int blocks_used;
// We pollute a little bit the memory before the tests
@ -20,37 +23,28 @@ int main() {
}
}
_TEST_ASSERT(true, "Pollution de la mémoire");
// We test in a first place that one simple allocation works as expected
mem_used = get_memory_distinct_allocations();
blocks_used = get_memory_blocks_number();
void* ptr = nalloc(15, 1);
if (! (get_memory_distinct_allocations() <= mem_used+1)) {
printf_error((char*)"Plus d'un élément de mémoire alloué en une seule allocation\n");
exit(1);
}
_TEST_ASSERT((get_memory_distinct_allocations() <= mem_used+1), "Un seul bloc mémoire alloué par allocation (max)");
gree(ptr, false);
if (! (get_memory_blocks_number() == blocks_used)) {
printf_error((char*)"La mémoire n'a pas été libérée correctement\n");
exit(1);
}
printf(GREEN "OK\n" RESET);
_TEST_ASSERT((get_memory_blocks_number() == blocks_used), "Libération partielle de la mémoire");
printf("Allocation de deux demi-blocs\n");
// We test that we do not use too much blocks
blocks_used = get_memory_blocks_number();
void* ptr1 = nalloc(-1+MEMORY_BLOCK/2, 1);
void* ptr2 = nalloc(-1+MEMORY_BLOCK/2, 1);
if (! (get_memory_blocks_number() <= blocks_used +1)) {
printf_error((char*)"Trop de blocs ont été alloués par rapport à la mémoire nécessaire\n");
exit(1);
}
printf(GREEN "OK\n" RESET);
_TEST_ASSERT((get_memory_blocks_number() <= blocks_used +1), "Taille d'allocation de deux demi-blocs");
printf("Libération de la mémoire\n");
for (int i=1; i < N; i++) {
for (int j=0; j < i; j++) {
// We test that the memory does not overlap itself
@ -61,10 +55,7 @@ int main() {
gree(ptr1, false);
gree(ptr2, false);
if (! (get_memory_distinct_allocations() == 0 && get_memory_blocks_number() == 0)) {
printf_error((char*)"La mémoire n'a pas été libérée correctement\n");
exit(1);
}
printf(GREEN "OK\n" RESET);
_TEST_ASSERT((get_memory_distinct_allocations() == 0 && get_memory_blocks_number() == 0), "Libération totale de la mémoire");
return 0;
}

View File

@ -2,6 +2,8 @@
#include <stdio.h>
#include <assert.h>
#include "include/test.h"
#include "../src/common/include/memory_management.h"
#include "../src/common/include/colors.h"
#include "../src/common/include/utils.h"
@ -17,6 +19,8 @@ __global__ void check_access(int* array, int range) {
int main() {
_TEST_PRESENTATION("Memory management (Cuda part)")
printf("Checking CUDA compatibility.\n");
bool cuda_compatible = cuda_setup(true);
if (!cuda_compatible) {
@ -25,7 +29,6 @@ int main() {
}
printf(GREEN "OK\n" RESET);
printf("Pollution de la mémoire\n");
int mem_used;
int blocks_used;
// We pollute a little bit the memory before the tests
@ -37,23 +40,19 @@ int main() {
}
}
_TEST_ASSERT(true, "Pollution de la mémoire");
// We test in a first place that one simple allocation works as expected
mem_used = get_memory_distinct_allocations();
blocks_used = get_memory_blocks_number();
void* ptr = nalloc(15, 1);
if (! (get_memory_distinct_allocations() <= mem_used+1)) {
printf("Plus d'un élément de mémoire alloué en une seule allocation\n");
exit(1);
}
gree(ptr, false);
if (! (get_memory_blocks_number() == blocks_used)) {
printf("La mémoire n'a pas été libérée correctement\n");
exit(1);
}
printf(GREEN "OK\n" RESET);
printf("Vérification de l'accès CUDA\n");
_TEST_ASSERT((get_memory_distinct_allocations() <= mem_used+1), "Un seul bloc mémoire alloué par allocation (max)");
gree(ptr, false);
_TEST_ASSERT((get_memory_blocks_number() == blocks_used), "Libération partielle de la mémoire");
/* On lance des kernels de taille 1 ce qui est à la fois itératif et synchrone
* Donc un peu contraire à CUDA mais l'objectif est de pouvoir débugger facilement */
dim3 gridSize(1, 1, 1);
@ -64,37 +63,26 @@ int main() {
gpuErrchk( cudaPeekAtLastError() );
gpuErrchk( cudaDeviceSynchronize() );
}
printf(GREEN "OK\n" RESET);
printf("Allocation de deux demi-blocs\n");
_TEST_ASSERT(true, "accès CUDA à la mémoire")
// We test that we do not use too much blocks
blocks_used = get_memory_blocks_number();
void* ptr1 = nalloc(-1+MEMORY_BLOCK/2, 1);
void* ptr2 = nalloc(-1+MEMORY_BLOCK/2, 1);
if (! (get_memory_blocks_number() <= blocks_used +1)) {
printf("Trop de blocs ont été alloués par rapport à la mémoire nécessaire\n");
exit(1);
}
printf(GREEN "OK\n" RESET);
_TEST_ASSERT((get_memory_blocks_number() <= blocks_used +1), "Taille d'allocation de deux demi-blocs");
printf("Libération de la mémoire\n");
for (int i=1; i < N; i++) {
for (int j=0; j < i; j++) {
// We test that the memory does not overlap itself
assert(pointeurs[i][j] == i+1);
assert(pointeurs[i][j] == i);
}
gree(pointeurs[i], false);
}
gree(ptr1, false);
gree(ptr2, false);
if (! (get_memory_distinct_allocations() == 0 && get_memory_blocks_number() == 0)) {
printf("La mémoire n'a pas été libérée correctement\n");
exit(1);
}
printf(GREEN "OK\n" RESET);
_TEST_ASSERT((get_memory_distinct_allocations() == 0 && get_memory_blocks_number() == 0), "Libération totale de la mémoire");
return 0;
}