augustin64/tipe
1
1
Fork 0
mirror of https://github.com/augustin64/projet-tipe synced 2025-01-23 23:26:25 +01:00
Code Issues Packages Projects Releases Wiki Activity

Add reset_3d_array function

Browse Source
This commit is contained in:
augustin64 2023-03-30 18:08:13 +02:00
parent eeff720ae4
commit 2ee1bc4079
3 changed files with 78 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

8
src/include/utils.h
View File

@ -45,4 +45,12 @@ extern "C"
* Copier des valeurs d'un tableau de dimension 3 de mémoire partagée
*/
void copy_3d_array(float*** source, float*** dest, int dimension1, int dimension2, int dimension3);
#ifdef __CUDACC__
extern "C"
#endif
/*
* Remplir un tableau de 0.
*/
void reset_3d_array(float*** source, int dimension1, int dimension2, int dimension3);
#endif

35
src/utils.c
View File

@ -93,3 +93,38 @@ void copy_3d_array(float*** source, float*** dest, int dimension1, int dimension
}
}
#endif
#ifdef __CUDACC__
__global__ void reset_3d_array_kernel(float*** dest, int dimension1, int dimension2, int dimension3) {
int idx = threadIdx.x + blockDim.x*blockIdx.x; // < dimension1
int idy = threadIdx.y + blockDim.y*blockIdx.y; // < dimension2
int idz = threadIdx.z + blockDim.z*blockIdx.z; // < dimension3
if (idx >= dimension1 || idy >= dimension2 || idz >= dimension3) {
return;
}
dest[idx][idy][idz] = 0.;
}
extern "C"
void reset_3d_array(float*** dest, int dimension1, int dimension2, int dimension3) {
dim3 gridSize(i_div_up(dimension1, BLOCKSIZE_x), i_div_up(dimension2, BLOCKSIZE_y), i_div_up(dimension3, BLOCKSIZE_z));
dim3 blockSize(BLOCKSIZE_x, BLOCKSIZE_y, BLOCKSIZE_z);
reset_3d_array_kernel<<<gridSize, blockSize>>>(dest, dimension1, dimension2, dimension3);
gpuErrchk( cudaPeekAtLastError() );
gpuErrchk( cudaDeviceSynchronize() );
}
#else
void reset_3d_array(float*** dest, int dimension1, int dimension2, int dimension3) {
for (int i=0; i < dimension1; i++) {
for (int j=0; j < dimension2; j++) {
for (int k=0; k < dimension3; k++) {
dest[i][j][k] = 0.;
}
}
}
}
#endif

36
src/utils.cu
View File

@ -73,7 +73,6 @@ __global__ void copy_3d_array_kernel(float*** source, float*** dest, int dimensi
dest[idx][idy][idz] = source[idx][idy][idz];
}
extern "C"
void copy_3d_array(float*** source, float*** dest, int dimension1, int dimension2, int dimension3) {
dim3 gridSize(i_div_up(dimension1, BLOCKSIZE_x), i_div_up(dimension2, BLOCKSIZE_y), i_div_up(dimension3, BLOCKSIZE_z));
dim3 blockSize(BLOCKSIZE_x, BLOCKSIZE_y, BLOCKSIZE_z);
@ -94,3 +93,38 @@ void copy_3d_array(float*** source, float*** dest, int dimension1, int dimension
}
}
#endif
#ifdef __CUDACC__
__global__ void reset_3d_array_kernel(float*** dest, int dimension1, int dimension2, int dimension3) {
int idx = threadIdx.x + blockDim.x*blockIdx.x; // < dimension1
int idy = threadIdx.y + blockDim.y*blockIdx.y; // < dimension2
int idz = threadIdx.z + blockDim.z*blockIdx.z; // < dimension3
if (idx >= dimension1 || idy >= dimension2 || idz >= dimension3) {
return;
}
dest[idx][idy][idz] = 0.;
}
extern "C"
void reset_3d_array(float*** dest, int dimension1, int dimension2, int dimension3) {
dim3 gridSize(i_div_up(dimension1, BLOCKSIZE_x), i_div_up(dimension2, BLOCKSIZE_y), i_div_up(dimension3, BLOCKSIZE_z));
dim3 blockSize(BLOCKSIZE_x, BLOCKSIZE_y, BLOCKSIZE_z);
reset_3d_array_kernel<<<gridSize, blockSize>>>(dest, dimension1, dimension2, dimension3);
gpuErrchk( cudaPeekAtLastError() );
gpuErrchk( cudaDeviceSynchronize() );
}
#else
void reset_3d_array(float*** dest, int dimension1, int dimension2, int dimension3) {
for (int i=0; i < dimension1; i++) {
for (int j=0; j < dimension2; j++) {
for (int k=0; k < dimension3; k++) {
dest[i][j][k] = 0.;
}
}
}
}
#endif