diff --git a/TPs/TP0/CODE/common/helper_cuda.h b/TPs/TP0/CODE/common/helper_cuda.h
index 3dd446dced85de9d10eb9f1597c56099fa358bdb..77ca0bf8a3091a41d462238e6b8cda540b6868c3 100644
--- a/TPs/TP0/CODE/common/helper_cuda.h
+++ b/TPs/TP0/CODE/common/helper_cuda.h
@@ -46,17 +46,16 @@
#define MIN(a, b) (a < b ? a : b)
#endif
-#define checkCudaErrors(val) check((val), #val, __FILE__, __LINE__)
#define getLastCudaError(msg) __getLastCudaError(msg, __FILE__, __LINE__)
+#define checkCudaErrors(ans) { gpuAssert((ans), __FILE__, __LINE__); }
-void check(cudaError_t result, char const *const func, const char *const file,
- int const line)
+inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=false)
{
- if (result)
- {
- fprintf(stderr, "CUDA error at %s:%d code=%d (%s) \"%s\" \n", file, line, (int)result, cudaGetErrorName(result), func);
- exit(EXIT_FAILURE);
- }
+ if (code != cudaSuccess)
+ {
+ fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
+ if (abort) exit(code);
+ }
}
inline void __getLastCudaError(const char *errorMessage, const char *file,
@@ -180,4 +179,4 @@ inline const char *_ConvertSMVer2ArchName(int major, int minor)
// end of CUDA Helper Functions
-#endif // COMMON_HELPER_CUDA_H_
\ No newline at end of file
+#endif // COMMON_HELPER_CUDA_H_
diff --git a/TPs/TP0/CODE/compute-sanitizer/Makefile b/TPs/TP0/CODE/compute-sanitizer/Makefile
index 2283121ddd828e6e4f8750965ef592866b4acec1..098d87b12f2989aa389f02fddd4ca65df28bd525 100644
--- a/TPs/TP0/CODE/compute-sanitizer/Makefile
+++ b/TPs/TP0/CODE/compute-sanitizer/Makefile
@@ -1,8 +1,8 @@
CC = nvcc -lineinfo
-INCLUDES =
+INCLUDES = -I../common
# the build target executable:
-EXE = memcheck.exe initcheck.exe
+EXE = memcheck.exe initcheck.exe out_of_bound.exe
all: $(EXE)
diff --git a/TPs/TP0/CODE/compute-sanitizer/memcheck.cu b/TPs/TP0/CODE/compute-sanitizer/memcheck.cu
index ad3fc59b16be296946be73084b4f19bb676c52fe..7cccf6ab69581d4cc18f137c90e460279a6be50e 100644
--- a/TPs/TP0/CODE/compute-sanitizer/memcheck.cu
+++ b/TPs/TP0/CODE/compute-sanitizer/memcheck.cu
@@ -4,7 +4,7 @@
__global__ void saxpy(int n, float a, float *x, float *y) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
- y[i] = a * x[i] + y[i];
+ if (i <= n) y[i] = a * x[i] + y[i];
}
int main(void) {
diff --git a/TPs/TP0/CODE/compute-sanitizer/out_of_bound.cu b/TPs/TP0/CODE/compute-sanitizer/out_of_bound.cu
new file mode 100644
index 0000000000000000000000000000000000000000..7ceaad48461a65759886060934cb5dd4fd5a4672
--- /dev/null
+++ b/TPs/TP0/CODE/compute-sanitizer/out_of_bound.cu
@@ -0,0 +1,13 @@
+#include "helper_cuda.h"
+
+__global__ void k1(char *d) { d[41 + 128] = 0; }
+
+int main() {
+ char *d;
+ checkCudaErrors(cudaMalloc(&d, 42));
+
+ k1<<<1, 1>>>(d);
+ checkCudaErrors(cudaGetLastError());
+ checkCudaErrors(cudaDeviceSynchronize());
+
+}
diff --git a/TPs/TP0/CODE/compute-sanitizer/tags b/TPs/TP0/CODE/compute-sanitizer/tags
index f41805ff57b31cfd930077730d066b2145261de1..3ca8122258a249f317cd0d4fe0ea2e13ff376c4e 100644
--- a/TPs/TP0/CODE/compute-sanitizer/tags
+++ b/TPs/TP0/CODE/compute-sanitizer/tags
@@ -4,6 +4,6 @@
!_TAG_PROGRAM_NAME Exuberant Ctags //
!_TAG_PROGRAM_URL http://ctags.sourceforge.net /official site/
!_TAG_PROGRAM_VERSION 5.9~svn20110310 //
-CC Makefile /^CC = nvcc$/;" m
-EXE Makefile /^EXE = memcheck.exe initcheck.exe $/;" m
-INCLUDES Makefile /^INCLUDES = $/;" m
+CC Makefile /^CC = nvcc -lineinfo$/;" m
+EXE Makefile /^EXE = memcheck.exe initcheck.exe out_of_bound.exe $/;" m
+INCLUDES Makefile /^INCLUDES = -I..\/common $/;" m
diff --git a/TPs/TP0/CODE/device_query/Makefile b/TPs/TP0/CODE/device_query/Makefile
index 73b63d518ef203ed6ec1b19cbec0c553bc630a18..32e68d57b8d6fb95157ba9072e94601681308929 100644
--- a/TPs/TP0/CODE/device_query/Makefile
+++ b/TPs/TP0/CODE/device_query/Makefile
@@ -1,5 +1,8 @@
+#***********************************************************
+# Created by Patricio Bulic, Davor Sluga UL FRI on 15/6/2022
+#***********************************************************
CC = nvcc
-INCLUDES = -I../common
+INCLUDES = -I.
# the build target executable:
TARGET = prog
diff --git a/TPs/TP0/CODE/error_checking/exemple1.cu b/TPs/TP0/CODE/error_checking/exemple1.cu
index 2c085a1a20bac6e02fcb0de6368336a9747e6aa9..e896cd4e3cf08c48910498d33a002ea8f40000fe 100755
--- a/TPs/TP0/CODE/error_checking/exemple1.cu
+++ b/TPs/TP0/CODE/error_checking/exemple1.cu
@@ -2,14 +2,19 @@
#include <stdlib.h>
#include "helper_cuda.h"
-#define THREADS 4096
-#define TAB_SIZE 8192
+#define THREADS 2048
+#define TAB_SIZE 100000
__global__ void kernel(int *a, int *b, int *c) {
const int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid < TAB_SIZE) c[tid] = a[tid] + b[tid];
}
+__global__ void init(int *a, int value) {
+ const int tid = blockIdx.x * blockDim.x + threadIdx.x;
+ if (tid < TAB_SIZE) a[tid] = value;
+}
+
int main(int argc, char **argv)
{
int sz_in_bytes = sizeof(int) * TAB_SIZE;
@@ -21,17 +26,23 @@ int main(int argc, char **argv)
// Allocation on host (malloc)
h_c = (int *)malloc(sz_in_bytes);
+ // Kernel configuration
+ dim3 dimBlock(THREADS, 1, 1);
+ dim3 dimGrid(TAB_SIZE / THREADS + 1, 1, 1);
+
// Allocation on device (cudaMalloc)
checkCudaErrors(cudaMalloc((void **)&d_a, sz_in_bytes));
checkCudaErrors(cudaMalloc((void **)&d_b, sz_in_bytes));
checkCudaErrors(cudaMalloc((void **)&d_c, sz_in_bytes));
- checkCudaErrors(cudaMemset(d_a, 1, sz_in_bytes));
- checkCudaErrors(cudaMemset(d_b, 2, sz_in_bytes));
+ init<<<dimGrid, dimBlock>>>(d_a, 1);
+ checkCudaErrors(cudaGetLastError());
+ checkCudaErrors(cudaDeviceSynchronize());
+
+ init<<<dimGrid, dimBlock>>>(d_b, 2);
+ checkCudaErrors(cudaGetLastError());
+ checkCudaErrors(cudaDeviceSynchronize());
- // Kernel configuration
- dim3 dimBlock(THREADS, 1, 1);
- dim3 dimGrid(TAB_SIZE / THREADS + 1, 1, 1);
// Kernel launch
kernel<<<dimGrid, dimBlock>>>(d_a, d_b, d_c);
diff --git a/TPs/TP0/CODE/error_checking/exemple2.cu b/TPs/TP0/CODE/error_checking/exemple2.cu
index bbef844d63fc6c0c0ae42fc93f524c70804a9b83..0a3c2d55e208d75457c8e6b88a3c8d744291ff12 100755
--- a/TPs/TP0/CODE/error_checking/exemple2.cu
+++ b/TPs/TP0/CODE/error_checking/exemple2.cu
@@ -2,59 +2,70 @@
#include <stdlib.h>
#include "helper_cuda.h"
-#define THREADS 256
-#define TAB_SIZE 8192
+#define THREADS 64
+#define TAB_SIZE 1000
__global__ void copy(int *a, int *b) {
- const int tid = blockIdx.x * blockDim.x + threadIdx.x;
- if (tid <= TAB_SIZE) b[tid] = a[tid];
+ const int tid = blockIdx.x * blockDim.x + threadIdx.x;
+ if (tid < TAB_SIZE) b[tid] = a[tid];
}
+__global__ void init(int *a, int value) {
+ const int tid = blockIdx.x * blockDim.x + threadIdx.x;
+ if (tid < TAB_SIZE) a[tid] = value;
+}
+
+
int main(int argc, char **argv)
{
- int sz_in_bytes = sizeof(int) * TAB_SIZE;
+ int sz_in_bytes = sizeof(int) * TAB_SIZE;
+
+ int *h_b;
+ int res = 0;
+ int *d_a, *d_b;
- int *h_b;
- int res = 0;
- int *d_a, *d_b;
+ // Allocation on host (malloc)
+ h_b = (int *)malloc(sz_in_bytes);
- // Allocation on host (malloc)
- h_b = (int *)malloc(sz_in_bytes);
+ // Allocation on device (cudaMalloc)
+ checkCudaErrors(cudaMalloc((void **)&d_a, sz_in_bytes));
+ // suppose we forget this
+ // checkCudaErrors(cudaMalloc((void **)&d_b, sz_in_bytes));
- // Allocation on device (cudaMalloc)
- checkCudaErrors(cudaMalloc((void **)&d_a, sz_in_bytes));
- checkCudaErrors(cudaMalloc((void **)&d_b, sz_in_bytes));
+ // Kernel configuration
+ dim3 dimBlock(THREADS, 1, 1);
+ dim3 dimGrid(TAB_SIZE / THREADS + 1, 1, 1);
- checkCudaErrors(cudaMemset(d_a, 1, sz_in_bytes));
+ init<<<dimGrid, dimBlock>>>(d_a, 1);
- // Kernel configuration
- dim3 dimBlock(THREADS, 1, 1);
- dim3 dimGrid(TAB_SIZE / THREADS + 1, 1, 1);
+ checkCudaErrors(cudaGetLastError());
+ checkCudaErrors(cudaDeviceSynchronize());
- // Kernel launch
- copy<<<dimGrid, dimBlock>>>(d_a, d_b);
- checkCudaErrors(cudaDeviceSynchronize());
+ // Kernel launch
+ copy<<<dimGrid, dimBlock>>>(d_a, d_b);
+ checkCudaErrors(cudaGetLastError());
+ checkCudaErrors(cudaDeviceSynchronize());
- // Retrieving data from device (cudaMemcpy)
- checkCudaErrors(cudaMemcpy(h_b, d_b, sz_in_bytes, cudaMemcpyDeviceToHost));
+ // Retrieving data from device (cudaMemcpy)
+ checkCudaErrors(cudaMemcpy(h_b, d_b, sz_in_bytes, cudaMemcpyDeviceToHost));
- // Freeing on device (cudaFree)
- checkCudaErrors(cudaFree(d_a));
- checkCudaErrors(cudaFree(d_b));
+ // Freeing on device (cudaFree)
+ checkCudaErrors(cudaFree(d_a));
+ checkCudaErrors(cudaFree(d_b));
- // computing sum of tab element
- for (int i = 0; i < TAB_SIZE; i++) res += h_b[i];
+ // computing sum of tab element
+ for (int i = 0; i < TAB_SIZE; i++) res += h_b[i];
- // Verifying if
- if (res == TAB_SIZE) {
- fprintf(stderr, "TEST PASSED !\n");
- }
- else
- {
- fprintf(stderr, "TEST FAILED !\n");
- }
+ // Verifying if
+ if (res == TAB_SIZE) {
+ fprintf(stderr, "TEST PASSED !\n");
+ }
+ else
+ {
+ fprintf(stderr, "TEST FAILED !\n");
+ }
- free(h_b);
+ free(h_b);
- return 0;
+ return 0;
}
diff --git a/TPs/TP0/CODE/error_checking/exemple3.cu b/TPs/TP0/CODE/error_checking/exemple3.cu
index d2a0c2ce1dd4eb0e8f55c29b89402f59146c1ef2..e8cd6a045f54f66412c1d6db9335489da58e666d 100755
--- a/TPs/TP0/CODE/error_checking/exemple3.cu
+++ b/TPs/TP0/CODE/error_checking/exemple3.cu
@@ -2,61 +2,72 @@
#include <stdlib.h>
#include "helper_cuda.h"
-#define THREADS 256
-#define TAB_SIZE 8192
+#define THREADS 64
+#define TAB_SIZE 1000
__global__ void copy(int *a, int *b) {
- const int tid = blockIdx.x * blockDim.x + threadIdx.x;
- if (tid < TAB_SIZE) b[tid] = a[tid];
+ const int tid = blockIdx.x * blockDim.x + threadIdx.x;
+ if (tid < TAB_SIZE) b[tid] = a[tid];
}
+__global__ void init(int *a, int value) {
+ const int tid = blockIdx.x * blockDim.x + threadIdx.x;
+ if (tid < TAB_SIZE) a[tid] = value;
+}
+
+
int main(int argc, char **argv)
{
- int sz_in_bytes = sizeof(int) * TAB_SIZE;
+ int sz_in_bytes = sizeof(int) * TAB_SIZE;
+
+ int *h_b;
+ int res = 0;
+ int *d_a, *d_b;
- int *h_b;
- int res = 0;
- int *d_a, *d_b;
+ // Allocation on host (malloc)
+ h_b = (int *)malloc(sz_in_bytes);
- // Allocation on host (malloc)
- h_b = (int *)malloc(sz_in_bytes);
+ // Too big allocation on device
+ checkCudaErrors(cudaMalloc((void **)&d_a, 100000000000));
- // Too big allocation on device
- checkCudaErrors(cudaMalloc((void **)&d_a, 100000000000));
+ // Allocation on device (cudaMalloc)
+ checkCudaErrors(cudaMalloc((void **)&d_a, sz_in_bytes));
+ checkCudaErrors(cudaMalloc((void **)&d_b, sz_in_bytes));
- // Allocation on device (cudaMalloc)
- checkCudaErrors(cudaMalloc((void **)&d_a, sz_in_bytes));
- checkCudaErrors(cudaMalloc((void **)&d_b, sz_in_bytes));
+ // Kernel configuration
+ dim3 dimBlock(THREADS, 1, 1);
+ dim3 dimGrid(TAB_SIZE / THREADS + 1, 1, 1);
- checkCudaErrors(cudaMemset(d_a, 1, sz_in_bytes));
+ init<<<dimGrid, dimBlock>>>(d_a, 1);
- // Kernel configuration
- dim3 dimBlock(THREADS, 1, 1);
- dim3 dimGrid(TAB_SIZE / THREADS + 1, 1, 1);
+ checkCudaErrors(cudaGetLastError());
+ checkCudaErrors(cudaDeviceSynchronize());
- // Kernel launch
- copy<<<dimGrid, dimBlock>>>(d_a, d_b);
+ // Kernel launch
+ copy<<<dimGrid, dimBlock>>>(d_a, d_b);
+ checkCudaErrors(cudaGetLastError());
+ checkCudaErrors(cudaDeviceSynchronize());
- // Retrieving data from device (cudaMemcpy)
- checkCudaErrors(cudaMemcpy(h_b, d_b, sz_in_bytes, cudaMemcpyDeviceToHost));
+ // Retrieving data from device (cudaMemcpy)
+ checkCudaErrors(cudaMemcpy(h_b, d_b, sz_in_bytes, cudaMemcpyDeviceToHost));
- // Freeing on device (cudaFree)
- checkCudaErrors(cudaFree(d_a));
- checkCudaErrors(cudaFree(d_b));
+ // Freeing on device (cudaFree)
+ checkCudaErrors(cudaFree(d_a));
+ checkCudaErrors(cudaFree(d_b));
- // computing sum of tab element
- for (int i = 0; i < TAB_SIZE; i++) res += h_b[i];
+ // computing sum of tab element
+ for (int i = 0; i < TAB_SIZE; i++) res += h_b[i];
- // Verifying if
- if (res == TAB_SIZE) {
- fprintf(stderr, "TEST PASSED !\n");
- }
- else
- {
- fprintf(stderr, "TEST FAILED !\n");
- }
+ // Verifying if
+ if (res == TAB_SIZE) {
+ fprintf(stderr, "TEST PASSED !\n");
+ }
+ else
+ {
+ fprintf(stderr, "TEST FAILED !\n");
+ }
- free(h_b);
+ free(h_b);
- return 0;
+ return 0;
}
diff --git a/TPs/TP0/CORRECTION/device_query/Makefile b/TPs/TP0/CORRECTION/device_query/Makefile
new file mode 100644
index 0000000000000000000000000000000000000000..5264cf81a352e7036ca6fa7565b7736350881936
--- /dev/null
+++ b/TPs/TP0/CORRECTION/device_query/Makefile
@@ -0,0 +1,17 @@
+CC = nvcc
+INCLUDES = -I.
+
+# the build target executable:
+TARGET = prog
+
+all: $(TARGET)
+
+$(TARGET): $(TARGET).o
+ $(CC) $(TARGET).o -o $@
+
+$(TARGET).o: $(TARGET).cu
+ $(CC) $(INCLUDES) $(TARGET).cu -c -o $@
+
+clean:
+ $(RM) $(TARGET) $(TARGET).o
+
diff --git a/TPs/TP0/CORRECTION/device_query/helper_cuda.h b/TPs/TP0/CORRECTION/device_query/helper_cuda.h
new file mode 100644
index 0000000000000000000000000000000000000000..3dd446dced85de9d10eb9f1597c56099fa358bdb
--- /dev/null
+++ b/TPs/TP0/CORRECTION/device_query/helper_cuda.h
@@ -0,0 +1,183 @@
+/* Copyright (c) 2022, NVIDIA CORPORATION. All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions
+ * are met:
+ * * Redistributions of source code must retain the above copyright
+ * notice, this list of conditions and the following disclaimer.
+ * * Redistributions in binary form must reproduce the above copyright
+ * notice, this list of conditions and the following disclaimer in the
+ * documentation and/or other materials provided with the distribution.
+ * * Neither the name of NVIDIA CORPORATION nor the names of its
+ * contributors may be used to endorse or promote products derived
+ * from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
+ * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
+ * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+
+////////////////////////////////////////////////////////////////////////////////
+// These are CUDA Helper functions for initialization and error checking
+
+#ifndef COMMON_HELPER_CUDA_H_
+#define COMMON_HELPER_CUDA_H_
+
+#pragma once
+
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+
+#ifndef MAX
+#define MAX(a, b) (a > b ? a : b)
+#endif
+
+#ifndef MIN
+#define MIN(a, b) (a < b ? a : b)
+#endif
+
+#define checkCudaErrors(val) check((val), #val, __FILE__, __LINE__)
+#define getLastCudaError(msg) __getLastCudaError(msg, __FILE__, __LINE__)
+
+void check(cudaError_t result, char const *const func, const char *const file,
+ int const line)
+{
+ if (result)
+ {
+ fprintf(stderr, "CUDA error at %s:%d code=%d (%s) \"%s\" \n", file, line, (int)result, cudaGetErrorName(result), func);
+ exit(EXIT_FAILURE);
+ }
+}
+
+inline void __getLastCudaError(const char *errorMessage, const char *file,
+ const int line)
+{
+ cudaError_t err = cudaGetLastError();
+
+ if (cudaSuccess != err)
+ {
+ fprintf(stderr,
+ "%s(%i) : getLastCudaError() CUDA error :"
+ " %s : (%d) %s.\n",
+ file, line, errorMessage, (int)(err),
+ cudaGetErrorString(err));
+ exit(EXIT_FAILURE);
+ }
+}
+
+// Beginning of GPU Architecture definitions
+int _ConvertSMVer2Cores(int major, int minor)
+{
+ // Defines for GPU Architecture types (using the SM version to determine
+ // the # of cores per SM
+ typedef struct
+ {
+ int SM; // 0xMm (hexidecimal notation), M = SM Major version,
+ // and m = SM minor version
+ int Cores;
+ } sSMtoCores;
+
+ sSMtoCores nGpuArchCoresPerSM[] = {
+ {0x30, 192},
+ {0x32, 192},
+ {0x35, 192},
+ {0x37, 192},
+ {0x50, 128},
+ {0x52, 128},
+ {0x53, 128},
+ {0x60, 64},
+ {0x61, 128},
+ {0x62, 128},
+ {0x70, 64},
+ {0x72, 64},
+ {0x75, 64},
+ {0x80, 64},
+ {0x86, 128},
+ {0x87, 128},
+ {-1, -1}};
+
+ int index = 0;
+
+ while (nGpuArchCoresPerSM[index].SM != -1)
+ {
+ if (nGpuArchCoresPerSM[index].SM == ((major << 4) + minor))
+ {
+ return nGpuArchCoresPerSM[index].Cores;
+ }
+
+ index++;
+ }
+
+ // If we don't find the values, we default use the previous one
+ // to run properly
+ printf(
+ "MapSMtoCores for SM %d.%d is undefined."
+ " Default to use %d Cores/SM\n",
+ major, minor, nGpuArchCoresPerSM[index - 1].Cores);
+ return nGpuArchCoresPerSM[index - 1].Cores;
+}
+
+inline const char *_ConvertSMVer2ArchName(int major, int minor)
+{
+ // Defines for GPU Architecture types (using the SM version to determine
+ // the GPU Arch name)
+ typedef struct
+ {
+ int SM; // 0xMm (hexidecimal notation), M = SM Major version,
+ // and m = SM minor version
+ const char *name;
+ } sSMtoArchName;
+
+ sSMtoArchName nGpuArchNameSM[] = {
+ {0x30, "Kepler"},
+ {0x32, "Kepler"},
+ {0x35, "Kepler"},
+ {0x37, "Kepler"},
+ {0x50, "Maxwell"},
+ {0x52, "Maxwell"},
+ {0x53, "Maxwell"},
+ {0x60, "Pascal"},
+ {0x61, "Pascal"},
+ {0x62, "Pascal"},
+ {0x70, "Volta"},
+ {0x72, "Xavier"},
+ {0x75, "Turing"},
+ {0x80, "Ampere"},
+ {0x86, "Ampere"},
+ {-1, "Graphics Device"}};
+
+ int index = 0;
+
+ while (nGpuArchNameSM[index].SM != -1)
+ {
+ if (nGpuArchNameSM[index].SM == ((major << 4) + minor))
+ {
+ return nGpuArchNameSM[index].name;
+ }
+
+ index++;
+ }
+
+ // If we don't find the values, we default use the previous one
+ // to run properly
+ printf(
+ "MapSMtoArchName for SM %d.%d is undefined."
+ " Default to use %s\n",
+ major, minor, nGpuArchNameSM[index - 1].name);
+ return nGpuArchNameSM[index - 1].name;
+}
+// end of GPU Architecture definitions
+
+// end of CUDA Helper Functions
+
+#endif // COMMON_HELPER_CUDA_H_
\ No newline at end of file
diff --git a/TPs/TP0/CORRECTION/device_query/prog.cu b/TPs/TP0/CORRECTION/device_query/prog.cu
new file mode 100644
index 0000000000000000000000000000000000000000..9700ec2d363038a2667245aec5f9706ded410b48
--- /dev/null
+++ b/TPs/TP0/CORRECTION/device_query/prog.cu
@@ -0,0 +1,115 @@
+#include <stdio.h>
+#include <cuda_runtime.h>
+#include <cuda.h>
+
+#include "helper_cuda.h"
+
+int main(int argc, char **argv) {
+
+ // Get number of GPUs
+ int deviceCount = 0;
+ cudaError_t error = cudaGetDeviceCount(&deviceCount);
+
+ if (error != cudaSuccess) {
+ printf("cudaGetDeviceCount error %d\n-> %s\n", error, cudaGetErrorString(error));
+ exit(EXIT_FAILURE);
+ }
+
+ // Get device propreties and print
+ for (int dev = 0; dev < deviceCount; dev++) {
+ struct cudaDeviceProp prop;
+ int value;
+ printf("\n========== cudaDeviceGetProperties ============ \n");
+ cudaGetDeviceProperties(&prop, dev);
+ printf("\nDevice %d: \"%s\"\n", dev, prop.name);
+ printf(" GPU Clock Rate (MHz): %d\n", prop.clockRate/1000);
+ printf(" Memory Clock Rate (MHz): %d\n", prop.memoryClockRate/1000);
+ printf(" Memory Bus Width (bits): %d\n", prop.memoryBusWidth);
+ printf(" Peak Memory Bandwidth (GB/s): %.2f\n",
+ 2.0 * prop.memoryClockRate * (prop.memoryBusWidth / 8) / 1.0e6);
+ printf(" CUDA Cores/MP: %d\n", _ConvertSMVer2Cores(prop.major, prop.minor));
+ printf(" CUDA Cores: %d\n", _ConvertSMVer2Cores(prop.major, prop.minor) *
+ prop.multiProcessorCount);
+ printf(" Peak TFLOPS: %.1f\n",
+ (_ConvertSMVer2Cores(prop.major, prop.minor) *
+ 2.0 * prop.multiProcessorCount * prop.clockRate * 1000.0)/ 1.0e12);
+ printf(" Total amount of global memory: %.0f GB\n", prop.totalGlobalMem / 1073741824.0f);
+ printf(" Total amount of shared memory per block: %zu kB\n",
+ prop.sharedMemPerBlock/1024);
+ printf(" Total number of registers available per block: %d\n",
+ prop.regsPerBlock);
+ printf(" Warp size: %d\n",
+ prop.warpSize);
+ printf(" Maximum number of threads per block: %d\n",
+ prop.maxThreadsPerBlock);
+ printf(" Max dimension size of a thread block (x,y,z): (%d, %d, %d)\n",
+ prop.maxThreadsDim[0], prop.maxThreadsDim[1],
+ prop.maxThreadsDim[2]);
+ printf(" Max dimension size of a grid size (x,y,z): (%d, %d, %d)\n",
+ prop.maxGridSize[0], prop.maxGridSize[1],
+ prop.maxGridSize[2]);
+
+ printf("\n\n========== cudaDeviceGetAttribute ============ \n");
+ printf("\nDevice %d: \"%s\"\n", dev, prop.name);
+ cudaDeviceGetAttribute (&value, cudaDevAttrMaxThreadsPerBlock, dev);
+ printf(" Max number of threads per block: %d\n",
+ value);
+ cudaDeviceGetAttribute (&value, cudaDevAttrMaxBlockDimX, dev);
+ printf(" Max block dimension X: %d\n",
+ value);
+ cudaDeviceGetAttribute (&value, cudaDevAttrMaxBlockDimY, dev);
+ printf(" Max block dimension Y: %d\n",
+ value);
+ cudaDeviceGetAttribute (&value, cudaDevAttrMaxBlockDimZ, dev);
+ printf(" Max block dimension Z: %d\n",
+ value);
+ cudaDeviceGetAttribute (&value, cudaDevAttrMaxGridDimX, dev);
+ printf(" Max grid dimension X: %d\n",
+ value);
+ cudaDeviceGetAttribute (&value, cudaDevAttrMaxGridDimY, dev);
+ printf(" Max grid dimension Y: %d\n",
+ value);
+ cudaDeviceGetAttribute (&value, cudaDevAttrMaxGridDimZ, dev);
+ printf(" Max grid dimension Z: %d\n",
+ value);
+ cudaDeviceGetAttribute (&value, cudaDevAttrMaxSharedMemoryPerBlock, dev);
+ printf(" Max shared memory per block: %d\n",
+ value);
+ cudaDeviceGetAttribute (&value, cudaDevAttrWarpSize, dev);
+ printf(" Warp size: %d\n",
+ value);
+ cudaDeviceGetAttribute (&value, cudaDevAttrClockRate, dev);
+ printf(" Peak clock frequency in kilohertz: %d\n",
+ value);
+ cudaDeviceGetAttribute (&value, cudaDevAttrMemoryClockRate, dev);
+ printf(" Peak memory clock frequency in kilohertz: %d\n",
+ value);
+ cudaDeviceGetAttribute (&value, cudaDevAttrGlobalMemoryBusWidth, dev);
+ printf(" Global memory bus width in bits: %d\n",
+ value);
+ cudaDeviceGetAttribute (&value, cudaDevAttrL2CacheSize, dev);
+ printf(" Size of L2 cache in bytes: %d\n",
+ value);
+ cudaDeviceGetAttribute (&value, cudaDevAttrMaxThreadsPerMultiProcessor, dev);
+ printf(" Maximum resident threads per SM: %d\n",
+ value);
+ cudaDeviceGetAttribute (&value, cudaDevAttrComputeCapabilityMajor, dev);
+ printf(" Major compute capability version number: %d\n",
+ value);
+ cudaDeviceGetAttribute (&value, cudaDevAttrComputeCapabilityMinor, dev);
+ printf(" Minor compute capability version number: %d\n",
+ value);
+ cudaDeviceGetAttribute (&value, cudaDevAttrMaxSharedMemoryPerMultiprocessor, dev);
+ printf(" Max shared memory per SM in bytes: %d\n",
+ value);
+ cudaDeviceGetAttribute (&value, cudaDevAttrMaxRegistersPerMultiprocessor, dev);
+ printf(" Max number of 32-bit registers per SM: %d\n",
+ value);
+ cudaDeviceGetAttribute (&value, cudaDevAttrMaxSharedMemoryPerBlockOptin, dev);
+ printf(" Max per block shared mem size on the device: %d\n",
+ value);
+ cudaDeviceGetAttribute (&value, cudaDevAttrMaxBlocksPerMultiprocessor, dev);
+ printf(" Max thread blocks that can reside on a SM: %d\n",
+ value);
+ }
+}
diff --git a/TPs/TP0/CORRECTION/device_query/tags b/TPs/TP0/CORRECTION/device_query/tags
new file mode 100644
index 0000000000000000000000000000000000000000..74edb7e8f6999f738763c9a1b17486ab326d5c8b
--- /dev/null
+++ b/TPs/TP0/CORRECTION/device_query/tags
@@ -0,0 +1,18 @@
+!_TAG_FILE_FORMAT 2 /extended format; --format=1 will not append ;" to lines/
+!_TAG_FILE_SORTED 1 /0=unsorted, 1=sorted, 2=foldcase/
+!_TAG_PROGRAM_AUTHOR Darren Hiebert /dhiebert@users.sourceforge.net/
+!_TAG_PROGRAM_NAME Exuberant Ctags //
+!_TAG_PROGRAM_URL http://ctags.sourceforge.net /official site/
+!_TAG_PROGRAM_VERSION 5.9~svn20110310 //
+CC Makefile /^CC = nvcc$/;" m
+COMMON_HELPER_CUDA_H_ helper_cuda.h 32;" d
+INCLUDES Makefile /^INCLUDES = -I.$/;" m
+MAX helper_cuda.h 42;" d
+MIN helper_cuda.h 46;" d
+TARGET Makefile /^TARGET = prog$/;" m
+_ConvertSMVer2ArchName helper_cuda.h /^inline const char *_ConvertSMVer2ArchName(int major, int minor)$/;" f
+_ConvertSMVer2Cores helper_cuda.h /^int _ConvertSMVer2Cores(int major, int minor)$/;" f
+__getLastCudaError helper_cuda.h /^inline void __getLastCudaError(const char *errorMessage, const char *file,$/;" f
+check helper_cuda.h /^void check(cudaError_t result, char const *const func, const char *const file,$/;" f
+checkCudaErrors helper_cuda.h 49;" d
+getLastCudaError helper_cuda.h 50;" d
diff --git a/TPs/TP0/CORRECTION/error_checking/exemple2.cu b/TPs/TP0/CORRECTION/error_checking/exemple2.cu
new file mode 100755
index 0000000000000000000000000000000000000000..783478f375efe2ff0eaaf295f1d9fe3bf9d1a595
--- /dev/null
+++ b/TPs/TP0/CORRECTION/error_checking/exemple2.cu
@@ -0,0 +1,72 @@
+#include <stdio.h>
+#include <stdlib.h>
+#include "helper_cuda.h"
+
+#define THREADS 64
+#define TAB_SIZE 1000
+
+__global__ void copy(int *a, int *b) {
+ const int tid = blockIdx.x * blockDim.x + threadIdx.x;
+ if (tid < TAB_SIZE) b[tid] = a[tid];
+}
+
+__global__ void init(int *a, int value) {
+ const int tid = blockIdx.x * blockDim.x + threadIdx.x;
+ if (tid < TAB_SIZE) a[tid] = value;
+}
+
+
+int main(int argc, char **argv)
+{
+ int sz_in_bytes = sizeof(int) * TAB_SIZE;
+
+ int *h_b;
+ int res = 0;
+ int *d_a, *d_b, *d_dummy;
+
+ // Allocation on host (malloc)
+ h_b = (int *)malloc(sz_in_bytes);
+
+ // Allocation on device (cudaMalloc)
+ checkCudaErrors(cudaMalloc((void **)&d_a, sz_in_bytes));
+ // suppose we forget this
+ // checkCudaErrors(cudaMalloc((void **)&d_b, sz_in_bytes));
+
+ // Kernel configuration
+ dim3 dimBlock(THREADS, 1, 1);
+ dim3 dimGrid(TAB_SIZE / THREADS + 1, 1, 1);
+
+ init<<<dimGrid, dimBlock>>>(d_a, 1);
+
+ checkCudaErrors(cudaGetLastError());
+ checkCudaErrors(cudaDeviceSynchronize());
+
+ // Kernel launch
+ copy<<<dimGrid, dimBlock>>>(d_a, d_b);
+ checkCudaErrors(cudaMalloc((void **)&d_dummy, sz_in_bytes));
+ checkCudaErrors(cudaGetLastError());
+ checkCudaErrors(cudaDeviceSynchronize());
+
+ // Retrieving data from device (cudaMemcpy)
+ checkCudaErrors(cudaMemcpy(h_b, d_b, sz_in_bytes, cudaMemcpyDeviceToHost));
+
+ // Freeing on device (cudaFree)
+ checkCudaErrors(cudaFree(d_a));
+ checkCudaErrors(cudaFree(d_b));
+
+ // computing sum of tab element
+ for (int i = 0; i < TAB_SIZE; i++) res += h_b[i];
+
+ // Verifying if
+ if (res == TAB_SIZE) {
+ fprintf(stderr, "TEST PASSED !\n");
+ }
+ else
+ {
+ fprintf(stderr, "TEST FAILED !\n");
+ }
+
+ free(h_b);
+
+ return 0;
+}
diff --git a/TPs/TP0/CORRECTION/tags b/TPs/TP0/CORRECTION/tags
new file mode 100644
index 0000000000000000000000000000000000000000..8de213abb5cdc780fe0f2a444355bf977dbb19d1
--- /dev/null
+++ b/TPs/TP0/CORRECTION/tags
@@ -0,0 +1,18 @@
+!_TAG_FILE_FORMAT 2 /extended format; --format=1 will not append ;" to lines/
+!_TAG_FILE_SORTED 1 /0=unsorted, 1=sorted, 2=foldcase/
+!_TAG_PROGRAM_AUTHOR Darren Hiebert /dhiebert@users.sourceforge.net/
+!_TAG_PROGRAM_NAME Exuberant Ctags //
+!_TAG_PROGRAM_URL http://ctags.sourceforge.net /official site/
+!_TAG_PROGRAM_VERSION 5.9~svn20110310 //
+CC device_query/Makefile /^CC = nvcc$/;" m
+COMMON_HELPER_CUDA_H_ device_query/helper_cuda.h 32;" d
+INCLUDES device_query/Makefile /^INCLUDES = -I.$/;" m
+MAX device_query/helper_cuda.h 42;" d
+MIN device_query/helper_cuda.h 46;" d
+TARGET device_query/Makefile /^TARGET = prog$/;" m
+_ConvertSMVer2ArchName device_query/helper_cuda.h /^inline const char *_ConvertSMVer2ArchName(int major, int minor)$/;" f
+_ConvertSMVer2Cores device_query/helper_cuda.h /^int _ConvertSMVer2Cores(int major, int minor)$/;" f
+__getLastCudaError device_query/helper_cuda.h /^inline void __getLastCudaError(const char *errorMessage, const char *file,$/;" f
+check device_query/helper_cuda.h /^void check(cudaError_t result, char const *const func, const char *const file,$/;" f
+checkCudaErrors device_query/helper_cuda.h 49;" d
+getLastCudaError device_query/helper_cuda.h 50;" d
diff --git a/TPs/TP0/SUJET/tp0.pdf b/TPs/TP0/SUJET/tp0.pdf
index 9061738fe9e2dcb8996cd01a87506887027f81d9..b81d4ccc5c85f52bb976921c8d02fc19c164fc82 100644
Binary files a/TPs/TP0/SUJET/tp0.pdf and b/TPs/TP0/SUJET/tp0.pdf differ