diff --git a/cuda-omp/cuda/1/classwork_1.cu b/cuda-omp-openacc/cuda/1/classwork_1.cu similarity index 100% rename from cuda-omp/cuda/1/classwork_1.cu rename to cuda-omp-openacc/cuda/1/classwork_1.cu diff --git a/cuda-omp/cuda/1/classwork_2.cu b/cuda-omp-openacc/cuda/1/classwork_2.cu similarity index 100% rename from cuda-omp/cuda/1/classwork_2.cu rename to cuda-omp-openacc/cuda/1/classwork_2.cu diff --git a/cuda-omp/cuda/2/classwork.cu b/cuda-omp-openacc/cuda/2/classwork.cu similarity index 100% rename from cuda-omp/cuda/2/classwork.cu rename to cuda-omp-openacc/cuda/2/classwork.cu diff --git a/cuda-omp/cuda/3/classwork.cu b/cuda-omp-openacc/cuda/3/classwork.cu similarity index 100% rename from cuda-omp/cuda/3/classwork.cu rename to cuda-omp-openacc/cuda/3/classwork.cu diff --git a/cuda-omp/cuda/4/classwork.cu b/cuda-omp-openacc/cuda/4/classwork.cu similarity index 100% rename from cuda-omp/cuda/4/classwork.cu rename to cuda-omp-openacc/cuda/4/classwork.cu diff --git a/cuda-omp/cuda/5/aos_soa.cu b/cuda-omp-openacc/cuda/5/aos_soa.cu similarity index 100% rename from cuda-omp/cuda/5/aos_soa.cu rename to cuda-omp-openacc/cuda/5/aos_soa.cu diff --git a/cuda-omp/cuda/5/linked_list.cu b/cuda-omp-openacc/cuda/5/linked_list.cu similarity index 100% rename from cuda-omp/cuda/5/linked_list.cu rename to cuda-omp-openacc/cuda/5/linked_list.cu diff --git a/cuda-omp/cuda/6/classwork.cu b/cuda-omp-openacc/cuda/6/classwork.cu similarity index 100% rename from cuda-omp/cuda/6/classwork.cu rename to cuda-omp-openacc/cuda/6/classwork.cu diff --git a/cuda-omp/cuda/7/mat_mult.cu b/cuda-omp-openacc/cuda/7/mat_mult.cu similarity index 100% rename from cuda-omp/cuda/7/mat_mult.cu rename to cuda-omp-openacc/cuda/7/mat_mult.cu diff --git a/cuda-omp/cuda/7/mat_mult_block.cu b/cuda-omp-openacc/cuda/7/mat_mult_block.cu similarity index 100% rename from cuda-omp/cuda/7/mat_mult_block.cu rename to cuda-omp-openacc/cuda/7/mat_mult_block.cu diff --git a/cuda-omp/cuda/energy/Makefile b/cuda-omp-openacc/cuda/energy/Makefile similarity index 100% rename from cuda-omp/cuda/energy/Makefile rename to cuda-omp-openacc/cuda/energy/Makefile diff --git a/cuda-omp/cuda/energy/energy/energy_pmt.h b/cuda-omp-openacc/cuda/energy/energy/energy_pmt.h similarity index 100% rename from cuda-omp/cuda/energy/energy/energy_pmt.h rename to cuda-omp-openacc/cuda/energy/energy/energy_pmt.h diff --git a/cuda-omp/cuda/energy/energy/energy_pmt_methods.cpp b/cuda-omp-openacc/cuda/energy/energy/energy_pmt_methods.cpp similarity index 100% rename from cuda-omp/cuda/energy/energy/energy_pmt_methods.cpp rename to cuda-omp-openacc/cuda/energy/energy/energy_pmt_methods.cpp diff --git a/cuda-omp/cuda/energy/energy/energy_pmt_methods.h b/cuda-omp-openacc/cuda/energy/energy/energy_pmt_methods.h similarity index 100% rename from cuda-omp/cuda/energy/energy/energy_pmt_methods.h rename to cuda-omp-openacc/cuda/energy/energy/energy_pmt_methods.h diff --git a/cuda-omp/cuda/energy/mat_mult_block.cu b/cuda-omp-openacc/cuda/energy/mat_mult_block.cu similarity index 100% rename from cuda-omp/cuda/energy/mat_mult_block.cu rename to cuda-omp-openacc/cuda/energy/mat_mult_block.cu diff --git a/cuda-omp/cuda/miscellaneous/coalescing.cu b/cuda-omp-openacc/cuda/miscellaneous/coalescing.cu similarity index 100% rename from cuda-omp/cuda/miscellaneous/coalescing.cu rename to cuda-omp-openacc/cuda/miscellaneous/coalescing.cu diff --git a/cuda-omp/hybrid/hybrid_cublas_omp.c b/cuda-omp-openacc/hybrid/hybrid_cublas_omp.c similarity index 100% rename from cuda-omp/hybrid/hybrid_cublas_omp.c rename to cuda-omp-openacc/hybrid/hybrid_cublas_omp.c diff --git a/cuda-omp/hybrid/hybrid_omp_cublas.c b/cuda-omp-openacc/hybrid/hybrid_omp_cublas.c similarity index 100% rename from cuda-omp/hybrid/hybrid_omp_cublas.c rename to cuda-omp-openacc/hybrid/hybrid_omp_cublas.c diff --git a/cuda-omp/hybrid/hybrid_omp_cuda.cu b/cuda-omp-openacc/hybrid/hybrid_omp_cuda.cu similarity index 100% rename from cuda-omp/hybrid/hybrid_omp_cuda.cu rename to cuda-omp-openacc/hybrid/hybrid_omp_cuda.cu diff --git a/cuda-omp/omp/1/classwork_1.c b/cuda-omp-openacc/omp/1/classwork_1.c similarity index 100% rename from cuda-omp/omp/1/classwork_1.c rename to cuda-omp-openacc/omp/1/classwork_1.c diff --git a/cuda-omp/omp/1/classwork_2.c b/cuda-omp-openacc/omp/1/classwork_2.c similarity index 100% rename from cuda-omp/omp/1/classwork_2.c rename to cuda-omp-openacc/omp/1/classwork_2.c diff --git a/cuda-omp/omp/2/classwork.c b/cuda-omp-openacc/omp/2/classwork.c similarity index 100% rename from cuda-omp/omp/2/classwork.c rename to cuda-omp-openacc/omp/2/classwork.c diff --git a/cuda-omp/omp/3/classwork.c b/cuda-omp-openacc/omp/3/classwork.c similarity index 100% rename from cuda-omp/omp/3/classwork.c rename to cuda-omp-openacc/omp/3/classwork.c diff --git a/cuda-omp/omp/3/classwork_async.c b/cuda-omp-openacc/omp/3/classwork_async.c similarity index 100% rename from cuda-omp/omp/3/classwork_async.c rename to cuda-omp-openacc/omp/3/classwork_async.c diff --git a/cuda-omp/omp/4/classwork.c b/cuda-omp-openacc/omp/4/classwork.c similarity index 100% rename from cuda-omp/omp/4/classwork.c rename to cuda-omp-openacc/omp/4/classwork.c diff --git a/cuda-omp/omp/6/classwork.c b/cuda-omp-openacc/omp/6/classwork.c similarity index 100% rename from cuda-omp/omp/6/classwork.c rename to cuda-omp-openacc/omp/6/classwork.c diff --git a/cuda-omp/omp/7/mat_mult.c b/cuda-omp-openacc/omp/7/mat_mult.c similarity index 100% rename from cuda-omp/omp/7/mat_mult.c rename to cuda-omp-openacc/omp/7/mat_mult.c diff --git a/cuda-omp/omp/7/mat_mult_block.c b/cuda-omp-openacc/omp/7/mat_mult_block.c similarity index 100% rename from cuda-omp/omp/7/mat_mult_block.c rename to cuda-omp-openacc/omp/7/mat_mult_block.c diff --git a/cuda-omp/omp/energy/Makefile b/cuda-omp-openacc/omp/energy/Makefile similarity index 100% rename from cuda-omp/omp/energy/Makefile rename to cuda-omp-openacc/omp/energy/Makefile diff --git a/cuda-omp/omp/energy/energy/energy_pmt.h b/cuda-omp-openacc/omp/energy/energy/energy_pmt.h similarity index 100% rename from cuda-omp/omp/energy/energy/energy_pmt.h rename to cuda-omp-openacc/omp/energy/energy/energy_pmt.h diff --git a/cuda-omp/omp/energy/energy/energy_pmt_methods.cpp b/cuda-omp-openacc/omp/energy/energy/energy_pmt_methods.cpp similarity index 100% rename from cuda-omp/omp/energy/energy/energy_pmt_methods.cpp rename to cuda-omp-openacc/omp/energy/energy/energy_pmt_methods.cpp diff --git a/cuda-omp/omp/energy/energy/energy_pmt_methods.h b/cuda-omp-openacc/omp/energy/energy/energy_pmt_methods.h similarity index 100% rename from cuda-omp/omp/energy/energy/energy_pmt_methods.h rename to cuda-omp-openacc/omp/energy/energy/energy_pmt_methods.h diff --git a/cuda-omp/omp/energy/multiple_devices.c b/cuda-omp-openacc/omp/energy/multiple_devices.c similarity index 100% rename from cuda-omp/omp/energy/multiple_devices.c rename to cuda-omp-openacc/omp/energy/multiple_devices.c diff --git a/cuda-omp/omp/miscellaneous/asynchronous.c b/cuda-omp-openacc/omp/miscellaneous/asynchronous.c similarity index 100% rename from cuda-omp/omp/miscellaneous/asynchronous.c rename to cuda-omp-openacc/omp/miscellaneous/asynchronous.c diff --git a/cuda-omp/omp/miscellaneous/dependencies.c b/cuda-omp-openacc/omp/miscellaneous/dependencies.c similarity index 100% rename from cuda-omp/omp/miscellaneous/dependencies.c rename to cuda-omp-openacc/omp/miscellaneous/dependencies.c diff --git a/cuda-omp/omp/miscellaneous/host_device_address_space.c b/cuda-omp-openacc/omp/miscellaneous/host_device_address_space.c similarity index 100% rename from cuda-omp/omp/miscellaneous/host_device_address_space.c rename to cuda-omp-openacc/omp/miscellaneous/host_device_address_space.c diff --git a/cuda-omp/omp/miscellaneous/map_type_modifier.c b/cuda-omp-openacc/omp/miscellaneous/map_type_modifier.c similarity index 100% rename from cuda-omp/omp/miscellaneous/map_type_modifier.c rename to cuda-omp-openacc/omp/miscellaneous/map_type_modifier.c diff --git a/cuda-omp/omp/miscellaneous/memory_management.c b/cuda-omp-openacc/omp/miscellaneous/memory_management.c similarity index 100% rename from cuda-omp/omp/miscellaneous/memory_management.c rename to cuda-omp-openacc/omp/miscellaneous/memory_management.c diff --git a/cuda-omp/omp/miscellaneous/multiple_devices.c b/cuda-omp-openacc/omp/miscellaneous/multiple_devices.c similarity index 100% rename from cuda-omp/omp/miscellaneous/multiple_devices.c rename to cuda-omp-openacc/omp/miscellaneous/multiple_devices.c diff --git a/cuda-omp/omp/miscellaneous/structure.c b/cuda-omp-openacc/omp/miscellaneous/structure.c similarity index 100% rename from cuda-omp/omp/miscellaneous/structure.c rename to cuda-omp-openacc/omp/miscellaneous/structure.c diff --git a/cuda-omp/omp/miscellaneous/structure_routines.c b/cuda-omp-openacc/omp/miscellaneous/structure_routines.c similarity index 100% rename from cuda-omp/omp/miscellaneous/structure_routines.c rename to cuda-omp-openacc/omp/miscellaneous/structure_routines.c diff --git a/cuda-omp/omp/miscellaneous/task_reductions.c b/cuda-omp-openacc/omp/miscellaneous/task_reductions.c similarity index 100% rename from cuda-omp/omp/miscellaneous/task_reductions.c rename to cuda-omp-openacc/omp/miscellaneous/task_reductions.c diff --git a/cuda-omp/omp/miscellaneous/unified_shared_memory.c b/cuda-omp-openacc/omp/miscellaneous/unified_shared_memory.c similarity index 100% rename from cuda-omp/omp/miscellaneous/unified_shared_memory.c rename to cuda-omp-openacc/omp/miscellaneous/unified_shared_memory.c diff --git a/cuda-omp-openacc/openacc/1/classwork_1.c b/cuda-omp-openacc/openacc/1/classwork_1.c new file mode 100644 index 0000000000000000000000000000000000000000..3780cea57585d7e9aace408f1fb00a6a00653164 --- /dev/null +++ b/cuda-omp-openacc/openacc/1/classwork_1.c @@ -0,0 +1,109 @@ +////////////////////////////////////////////////////////////////////////////////////////////////// +// +// OpenACC GPU Offload is available only on systems with NVIDIA GPUs with compute capability '>= cc70' +// +// Assigment : write an OPENACC-GPU code corresponding to the +// following sequential C code +// +// #include <stdio.h> +// #define N 100 +// int main() +// { +// for (int i=0 ; i<N ; i++) +// printf("%d\n", (i * i)); + +// return 0; +// } +////////////////////////////////////////////////////////////////////////////////////////////////// + +////////////////////////////////////////////////////////////////////////////////////////////////// +// Author: David Goz +// mail : david.goz@inaf.it +// date : 06.07.2024 +// code tested using nvhpc +// +// - Compile the code to run on : +// $ nvc -acc -gpu=ccnative,debug,lineinfo -Minfo=all -v classwork_1.c -o classwork_1_acc +// - Run the code: +// $ ./classwork_1_acc +// - Check the result: +// $ ./classwork_1_acc | tail -n 100 | sort -nk 5 + +////////////////////////////////////////////////////////////////////////////////////////////////// + +#include <stdio.h> +#include <openacc.h> + +#define N 100 +#define NThreads 1024 + +#define NDEBUG + +void GPUkernelSerial(const int size) +{ +#pragma acc parallel seq + { +#if !defined(NDEBUG) + + if (!omp_is_initial_device()) + printf("\n\t GPU is executing GPUkernelSerial\n" ); + else + printf("\n\t CPU is executing GPUkernelSerial\n" ); + +#endif /* NDEBUG */ + + const int whoAmI = omp_get_thread_num(); + + for (int i=0 ; i<size ; i++) + printf("Hello from OPENACC-GPU thread: %d - result %d\n", whoAmI, (i * i)); + } /* omp target region - implicit barrier */ + + return; +} + +void GPUkernelParallel(const int size) +{ +#pragma omp target + { +#if !defined(NDEBUG) + + if (!omp_is_initial_device()) + printf("\n\t GPU is executing GPUkernelSerial\n" ); + else + printf("\n\t CPU is executing GPUkernelSerial\n" ); + +#endif /* NDEBUG */ + + #pragma omp teams distribute parallel for + for (int i=0 ; i<size ; i++) + { + /* get CUDA blockIdx.x */ + const int team = omp_get_team_num(); + + /* get CUDA threadIdx.x */ + const int tid = omp_get_thread_num(); + + /* get CUDA blockDim.x */ + const int nthr = omp_get_num_threads(); + + const int whoAmI = tid + (team * nthr); + + printf("Hello from OMP-GPU thread: %d - result %d\n", whoAmI, (i * i)); + } + } /* omp target region - implicit barrier */ + + return; +} + +int main() +{ + printf("\n\t The host issues the kernel on the GPU in serial\n"); + /* kernel lunch using one GPU thread */ + GPUkernelSerial(N); + + printf("\n\t The host issues the kernel on the GPU in parallel\n"); + /* kernel lunch using N GPU threads */ + GPUkernelParallel(N); + + return 0; +} diff --git a/cuda-omp-openacc/openacc/1/classwork_2.c b/cuda-omp-openacc/openacc/1/classwork_2.c new file mode 100644 index 0000000000000000000000000000000000000000..bc3a7da1c890adc7c6585f620c7d308a7b019258 --- /dev/null +++ b/cuda-omp-openacc/openacc/1/classwork_2.c @@ -0,0 +1,112 @@ +////////////////////////////////////////////////////////////////////////////////////////////////// +// +// OpenMP GPU Offload is available only on systems with NVIDIA GPUs with compute capability '>= cc70' +// +// Assigment : write an OMP-GPU code corresponding to the +// following sequential C code +// +// #include <stdio.h> +// #define N 100 +// int main() +// { +// int A[N]; +// +// for (int i=0 ; i<N ; i++) +// A[i] = (i * i); +// +// return 0; +// } +////////////////////////////////////////////////////////////////////////////////////////////////// + +////////////////////////////////////////////////////////////////////////////////////////////////// +// Author: David Goz +// mail : david.goz@inaf.it +// date : 06.07.2024 +// code tested using nvhpc +// +// - Compile the code: +// $ nvc -mp=gpu -gpu=ccnative,debug,lineinfo -target=gpu -Minfo=all -v classwork_2.c -o classwork_2_omp +// - Run the code: +// $ ./classwork_2_omp + +////////////////////////////////////////////////////////////////////////////////////////////////// + +#include <stdlib.h> +#include <stdio.h> +#include <assert.h> +#include <openacc.h> + +#define N 100 +#define NDEBUG + +void GPUkernel( int *A, + const int size) +{ + /* map A to the address space of the accelerator */ +#pragma omp target data map(from: A[0:size]) + { + /* kernel to be executed on the accelerator */ + #pragma omp target + { + /* create a bunch of teams (CUDA blocks) */ + #pragma omp teams + { + /* distribute the teams over index iterations and */ + /* spawn a bunch of threads within each team */ + #pragma omp distribute parallel for + for (int i=0 ; i<N ; i++) + { +#if !defined(NDEBUG) + + const int team = omp_get_team_num(); + const int nteam = omp_get_num_teams(); + const int tid = omp_get_thread_num(); + const int nthr = omp_get_num_threads(); + + const int whoAmI = tid + (team * nthr); + + if (!omp_is_initial_device()) + { + if (whoAmI == 0) + { + printf("\n\t GPU is executing GPUkernel\n" ); + printf("\n\t team : %d", team); + printf("\n\t nteam: %d", nteam); + printf("\n\t tid : %d", tid); + printf("\n\t nthr : %d\n", nthr); + } + } + else + printf("\n\t CPU is executing GPUkernel\n" ); + +#endif /* NDEBUG */ + + A[i] = (i * i); + } /* omp parallel for */ + } /* omp teams */ + } /* omp target */ + } /* omp target data */ + + return; +} + +int main() +{ + /* host array */ + int *A = (int *)malloc(N * sizeof(*A)); + assert(A != NULL); + + // kernel lunch + GPUkernel(A, N); + + // check the result + printf("\n"); + for (size_t i=0 ; i<N ; i++) + printf("\t A[%d] = %d", i, A[i]); + printf("\n\n"); + + // free host memory + free(A); + + return 0; +}