Fixing OpenMP examples

OpenMP/StreamTriad/CMakeLists.txt

@@ -72,6 +72,12 @@ set_target_properties(StreamTriad_par6 PROPERTIES COMPILE_FLAGS "${OpenMPAccel_C
 target_include_directories(StreamTriad_par6 PRIVATE "${CUDA_INCLUDE_DIRS}")
 set_target_properties(StreamTriad_par6 PROPERTIES LINK_FLAGS "${OpenMPAccel_C_FLAGS} ${CUDA_C_FLAGS}")
 
+# Adds build target of stream_triad_par7 with source code files
+add_executable(StreamTriad_par7 StreamTriad_par6.c timer.c timer.h)
+set_target_properties(StreamTriad_par7 PROPERTIES COMPILE_FLAGS "${OpenMPAccel_C_FLAGS}")
+target_include_directories(StreamTriad_par7 PRIVATE "${CUDA_INCLUDE_DIRS}")
+set_target_properties(StreamTriad_par7 PROPERTIES LINK_FLAGS "${OpenMPAccel_C_FLAGS} ${CUDA_C_FLAGS}")
+
 # Cleanup
 add_custom_target(distclean COMMAND rm -rf CMakeCache.txt CMakeFiles
                   Makefile cmake_install.cmake StreamTriad.dSYM ipo_out.optrpt)

OpenMP/StreamTriad/StreamTriad_par5.c

@@ -13,7 +13,7 @@ int main(int argc, char *argv[]){
    cudaMalloc((void *)&b,nsize*sizeof(double));
    cudaMalloc((void *)&c,nsize*sizeof(double));
 
-#pragma target enter data use_device_ptr(a, b, c)
+#pragma omp target enter data use_device_ptr(a, b, c)
 
    struct timespec tstart;
    // initializing data and arrays
@@ -36,7 +36,7 @@ int main(int argc, char *argv[]){
 
    printf("Average runtime for stream triad loop is %lf msecs\n", time_sum/ntimes);
 
-#pragma target exit data use_device_ptr(a, b, c)
+#pragma omp target exit data use_device_ptr(a, b, c)
 
    cudaFree(a);
    cudaFree(b);

OpenMP/StreamTriad/StreamTriad_par6.c

@@ -9,7 +9,7 @@ int main(int argc, char *argv[]){
    int nsize = 20000000, ntimes=16;
 
    double a[nsize], b[nsize], c[nsize];
-#pragma target enter data map(alloc:a[0:nsize], b[0:nsize], c[0:nsize])
+#pragma omp target enter data map(alloc:a[0:nsize], b[0:nsize], c[0:nsize])
 
    struct timespec tstart;
    // initializing data and arrays
@@ -32,7 +32,7 @@ int main(int argc, char *argv[]){
 
    printf("Average runtime for stream triad loop is %lf msecs\n", time_sum/ntimes);
 
-#pragma target exit data map(delete:a[0:nsize], b[0:nsize], c[0:nsize])
+#pragma omp target exit data map(delete:a[0:nsize], b[0:nsize], c[0:nsize])
 
    return(0);
 }

OpenMP/StreamTriad/StreamTriad_par7.c

@@ -0,0 +1,54 @@
+#include <stdio.h>
+#include <stdlib.h>
+#include <time.h>
+#include <cuda_runtime.h>
+#include "timer.h"
+
+int main(int argc, char *argv[]){
+
+   int nsize = 20000000, ntimes=16;
+
+#pragma omp declare target
+   double *a, *b, *c;
+#pragma omp end declare target
+
+#pragma omp target
+   {
+       a = malloc(nsize* sizeof(double);
+       b = malloc(nsize* sizeof(double);
+       c = malloc(nsize* sizeof(double);
+   }
+
+//#pragma omp target enter data map(alloc:a[0:nsize], b[0:nsize], c[0:nsize])
+
+   struct timespec tstart;
+   // initializing data and arrays
+   double scalar = 3.0, time_sum = 0.0;
+#pragma omp target teams distribute parallel for simd is_device_ptr(a, b)
+   for (int i=0; i<nsize; i++) {
+      a[i] = 1.0;
+      b[i] = 2.0;
+   }
+
+   for (int k=0; k<ntimes; k++){
+      cpu_timer_start(&tstart);
+      // stream triad loop 
+#pragma omp target teams distribute parallel for simd is_device_ptr(a, b, c)
+      for (int i=0; i<nsize; i++){
+         c[i] = a[i] + scalar*b[i];
+      }
+      time_sum += cpu_timer_stop(tstart);
+   }
+
+   printf("Average runtime for stream triad loop is %lf msecs\n", time_sum/ntimes);
+
+//#pragma omp target exit data map(delete:a[0:nsize], b[0:nsize], c[0:nsize])
+#pragma omp target
+   {
+       free(a);
+       free(b);
+       free(c);
+   }
+
+   return(0);
+}

OpenMP/StreamTriad/run.sh

@@ -20,3 +20,6 @@ echo "Running StreamTriad_par5 -- only allocating arrays on GPU with cudaMalloc"
 echo ""
 echo "Running StreamTriad_par6 -- map alloc arrays on GPU"
 ./StreamTriad_par6
+echo ""
+echo "Running StreamTriad_par7 -- declare target arrays on GPU"
+./StreamTriad_par7