diff --git a/Makefile b/Makefile index 91a20f6..2d7ab8f 100644 --- a/Makefile +++ b/Makefile @@ -161,26 +161,34 @@ 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)" -build-tests: prepare-tests $(TESTS_OBJ) $(BUILDDIR)/test-cnn_matrix_multiplication $(BUILDDIR)/test-cnn_convolution +build-tests: prepare-tests $(TESTS_OBJ) $(BUILDDIR)/test-cnn_matrix_multiplication $(BUILDDIR)/test-cnn_convolution $(BUILDDIR)/test-cuda_memory_management prepare-tests: @rm -f $(BUILDDIR)/test-* -build/test-cnn_%: $(TEST_SRCDIR)/cnn_%.c $(CNN_OBJ) $(BUILDDIR)/colors.o $(BUILDDIR)/mnist.o $(BUILDDIR)/utils.o $(BUILDDIR)/memory_management.o +$(BUILDDIR)/test-cnn_%: $(TEST_SRCDIR)/cnn_%.c $(CNN_OBJ) $(BUILDDIR)/colors.o $(BUILDDIR)/mnist.o $(BUILDDIR)/utils.o $(BUILDDIR)/memory_management.o $(CC) $^ -o $@ $(CFLAGS) $(LD_CFLAGS) # mnist.o est déjà inclus en tant que mnist_mnist.o -build/test-mnist_%: $(TEST_SRCDIR)/mnist_%.c $(MNIST_OBJ) $(BUILDDIR)/colors.o +$(BUILDDIR)/test-mnist_%: $(TEST_SRCDIR)/mnist_%.c $(MNIST_OBJ) $(BUILDDIR)/colors.o $(CC) $^ -o $@ $(CFLAGS) $(LD_CFLAGS) -build/test-memory_management: $(TEST_SRCDIR)/memory_management.c $(BUILDDIR)/colors.o $(BUILDDIR)/mnist.o $(BUILDDIR)/utils.o $(BUILDDIR)/test_memory_management.o +$(BUILDDIR)/test-memory_management: $(TEST_SRCDIR)/memory_management.c $(BUILDDIR)/colors.o $(BUILDDIR)/utils.o $(BUILDDIR)/test_memory_management.o $(CC) $^ -o $@ $(CFLAGS) $(LD_CFLAGS) $(BUILDDIR)/test_memory_management.o: $(SRCDIR)/memory_management.c $(SRCDIR)/include/memory_management.h $(CC) -c $< -o $@ $(CFLAGS) -DTEST_MEMORY_MANAGEMENT +ifdef NVCC_INSTALLED +$(BUILDDIR)/test-cuda_memory_management: $(TEST_SRCDIR)/memory_management.cu $(BUILDDIR)/colors.cuda.o $(BUILDDIR)/cuda_utils.o $(BUILDDIR)/cuda_memory_management.o + $(NVCC) $(LD_NVCCFLAGS) $(NVCCFLAGS) $^ -o $@ +else +$(BUILDDIR)/test-cuda_memory_management: + @echo "$(NVCC) not found, skipping" +endif + ifdef NVCC_INSTALLED $(BUILDDIR)/test-cnn_%: $(TEST_SRCDIR)/cnn_%.cu \ $(BUILDDIR)/cnn_cuda_%.o \ diff --git a/src/include/memory_management.h b/src/include/memory_management.h index e128a33..7023d06 100644 --- a/src/include/memory_management.h +++ b/src/include/memory_management.h @@ -19,12 +19,18 @@ typedef struct Memory { } Memory; // Renvoie le nombre d'allocations totales dans la mémoire +#ifdef __CUDACC__ +extern "C" +#endif int get_memory_distinct_allocations(); // Fonction récursive correspondante int get_distinct_allocations(Memory* mem); // Renvoie le nombre d'éléments dans la liste chaînée représentant la mémoire +#ifdef __CUDACC__ +extern "C" +#endif int get_memory_blocks_number(); // Renvoie la taille d'une liste chaînée diff --git a/test/memory_management.cu b/test/memory_management.cu new file mode 100644 index 0000000..1b8240f --- /dev/null +++ b/test/memory_management.cu @@ -0,0 +1,92 @@ +#include +#include +#include + +#include "../src/include/memory_management.h" +#include "../src/include/colors.h" +#include "../src/include/utils.h" + +#define N 350 + +__global__ void check_access(int* array, int range) { + for (int i=0; i < range; i++) { + array[i]++; + } +} + + + +int main() { + printf("Pollution de la mémoire\n"); + int mem_used; + int blocks_used; + // We pollute a little bit the memory before the tests + int* pointeurs[N]; + for (int i=1; i < N; i++) { + pointeurs[i] = (int*)nalloc(i*sizeof(int)); + for (int j=0; j < i; j++) { + pointeurs[i][j] = i; + } + } + + // 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); + 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); + 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"); + /* On lance des kernels detaille 1 ce qui est itératif synchrone + * Donc un peu contraire à CUDA mais l'objectif est de débugger faiclement */ + dim3 gridSize(1, 1, 1); + dim3 blockSize(1, 1, 1); + + for (int i=1; i < N; i++) { + check_access<<>>(pointeurs[i], i); + gpuErrchk( cudaPeekAtLastError() ); + gpuErrchk( cudaDeviceSynchronize() ); + } + printf(GREEN "OK\n" RESET); + + + 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); + void* ptr2 = nalloc(-1+MEMORY_BLOCK/2); + 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); + + + + 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); + } + gree(pointeurs[i]); + } + + gree(ptr1); + gree(ptr2); + 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); + return 0; +} \ No newline at end of file