Merge pull request #2472 from nawalcopty/development

Update examples in OneAPI Samples Repo to match GPU Opt Guide Repo

Author: jimmytwei

Publications/GPU-Opt-Guide/CMakeLists.txt

@@ -147,6 +147,7 @@ add_subdirectory(multiple-kernel-execution)
 add_subdirectory(work-group-size)
 add_subdirectory(registers)
 add_subdirectory(OpenMP)
+add_subdirectory(optimize-data-transfers)
 add_subdirectory(MPI)
 add_subdirectory(grf-mode-selection)
 add_subdirectory(fp-computations)

Publications/GPU-Opt-Guide/OpenMP/21_omp_target_alloc/test_target_map_f.f90

@@ -55,8 +55,7 @@ program main
 
     nstream_time = omp_get_wtime()
     do iter = 1, iterations
-       !$omp target teams distribute parallel do &
-           map(to: A, B) map(tofrom: C)
+       !$omp target teams distribute parallel do map(to: A, B) map(tofrom: C)
        do i = 1, length
           C(i) = C(i) + A(i) + scalar * B(i)
        end do

Publications/GPU-Opt-Guide/OpenMP/22_mkl_dispatch/dgemm_batch_example_01.cpp

@@ -4,7 +4,6 @@
 // SPDX-License-Identifier: MIT
 // =============================================================
 // clang-format off
-// Snippet begin
 #include <stdio.h>
 #include <stdlib.h>
 #include <math.h>
@@ -95,6 +94,7 @@ int main()
 
     t_start = omp_get_wtime();
 
+// Snippet begin
     // Call cblas_dgemm_batch
     #pragma omp target enter data \
       map(to: A1[0:m*k], B1[0:k*n], C1[0:m*n]) \
@@ -132,6 +132,7 @@ int main()
 
     #pragma omp target exit data \
       map(from: C1[0:m*n], C2[0:m*n])
+// Snippet end
 
     t_end = omp_get_wtime();
 
@@ -216,4 +217,3 @@ int main()
 
     return 0;
 }
-// Snippet end

Publications/GPU-Opt-Guide/OpenMP/22_mkl_dispatch/dgemm_batch_example_02.cpp

@@ -18,7 +18,6 @@
 * License.
 *******************************************************************************/
 
-// Snippet begin
 #include <stdio.h>
 #include <omp.h>
 #include "mkl.h"
@@ -153,6 +152,7 @@ int main() {
     cblas_dgemm_batch(layout, transA, transB, m, n, k, alpha, (const double **) a_array, lda,
                       (const double **) b_array, ldb, beta, c_ref_array, ldc, GROUP_COUNT, group_size);
 
+// Snippet begin
     double *a, *b, *c;
     for (i = 0; i < total_batch_size; i++) {
         a = a_array[i];
@@ -182,6 +182,7 @@ int main() {
         c = c_array[i];
 #pragma omp target exit data map(from:a[0:sizea_array[i]],b[0:sizeb_array[i]],c[0:sizec_array[i]])
     }
+// Snippet end
 
     double computed, reference, diff;
     MKL_INT l;
@@ -254,4 +255,3 @@ int main() {
     mkl_free(alpha); mkl_free(beta);
     return 0;
 }
-// Snippet end

Publications/GPU-Opt-Guide/OpenMP/22_mkl_dispatch/dgemm_dispatch_c.cpp

@@ -4,7 +4,6 @@
 // SPDX-License-Identifier: MIT
 // =============================================================
 // clang-format off
-// Snippet begin
 #include <stdio.h>
 #include <stdlib.h>
 #include <math.h>
@@ -61,12 +60,14 @@ int main()
 
     printf (" Computing matrix product using Intel oneMKL dgemm function via CBLAS interface \n\n");
 
+// Snippet begin
     #pragma omp target data map(to: A[0:m*k], B[0:k*n]) map(tofrom: C[0:m*n])
     {
-       #pragma omp target variant dispatch use_device_ptr(A, B, C)
+       #pragma omp dispatch
        cblas_dgemm(CblasRowMajor, CblasNoTrans, CblasNoTrans,
                    m, n, k, alpha, A, k, B, n, beta, C, n);
     }
+// Snippet end
 
     printf ("\n Top left corner of matrix C: \n");
     for (i=0; i<min(m,6); i++) {
@@ -115,7 +116,7 @@ int main()
     mkl_free(A);
     mkl_free(B);
     mkl_free(C);
+    mkl_free(C_fl);
 
     return fail;
 }
-// Snippet end

Publications/GPU-Opt-Guide/OpenMP/22_mkl_dispatch/dgemm_dispatch_f.f

@@ -3,7 +3,6 @@
         !
         ! SPDX-License-Identifier: MIT
         !=============================================================
-        ! Snippet begin
         include "mkl_omp_offload.f90"
 
         program DGEMM_MAIN
@@ -68,12 +67,14 @@ program DGEMM_MAIN
 
         ! Execute DGEMM on device
 
+! Snippet begin
 !$omp target data map(to: a, b) map(tofrom: c2)
 
 !$omp dispatch
         call DGEMM('N','N',m,n,k,alpha,a,m,b,k,beta,c2,m)
 
 !$omp end target data
+! Snippet end
 
         print *
         print *, 'c2 - After DGEMM device execution'
@@ -86,4 +87,3 @@ program DGEMM_MAIN
  110    format(7x,10(f10.2,2x))
 
         end
-        ! Snippet end

Publications/GPU-Opt-Guide/OpenMP/22_mkl_dispatch/dgemm_example_01.cpp

@@ -4,7 +4,6 @@
 // SPDX-License-Identifier: MIT
 // =============================================================
 // clang-format off
-// Snippet begin
 #include <stdio.h>
 #include <stdlib.h>
 #include <math.h>
@@ -70,6 +69,7 @@ int main()
 
     t_start = omp_get_wtime();
 
+// Snippet begin
     #pragma omp target data				   \
       map(to: A1[0:m*k], B1[0:k*n], A2[0:m*k], B2[0:k*n])  \
       map(tofrom: C1[0:m*n], C2[0:m*n])
@@ -82,6 +82,7 @@ int main()
        cblas_dgemm(CblasRowMajor, CblasNoTrans, CblasNoTrans,
                    m, n, k, alpha, A2, k, B2, n, beta, C2, n);
     }
+// Snippet end
 
     t_end = omp_get_wtime();
 
@@ -166,4 +167,3 @@ int main()
 
     return 0;
 }
-// Snippet end

Publications/GPU-Opt-Guide/OpenMP/22_mkl_dispatch/dgemm_example_02.cpp

@@ -4,7 +4,6 @@
 // SPDX-License-Identifier: MIT
 // =============================================================
 // clang-format off
-// Snippet begin
 #include <stdio.h>
 #include <stdlib.h>
 #include <math.h>
@@ -70,6 +69,7 @@ int main()
 
     t_start = omp_get_wtime();
 
+// Snippet begin
     #pragma omp target data				   \
       map(to: A1[0:m*k], B1[0:k*n], A2[0:m*k], B2[0:k*n])  \
       map(tofrom: C1[0:m*n], C2[0:m*n])
@@ -90,6 +90,7 @@ int main()
           }
        }
     }
+// Snippet end
 
     t_end = omp_get_wtime();
 
@@ -174,4 +175,3 @@ int main()
 
     return 0;
 }
-// Snippet end

Publications/GPU-Opt-Guide/OpenMP/22_mkl_dispatch/dgemm_example_03.cpp

@@ -4,7 +4,6 @@
 // SPDX-License-Identifier: MIT
 // =============================================================
 // clang-format off
-// Snippet begin
 #include <stdio.h>
 #include <stdlib.h>
 #include <math.h>
@@ -70,6 +69,7 @@ int main()
 
     t_start = omp_get_wtime();
 
+// Snippet begin
     #pragma omp target data				   \
       map(to: A1[0:m*k], B1[0:k*n], A2[0:m*k], B2[0:k*n])  \
       map(tofrom: C1[0:m*n], C2[0:m*n])
@@ -84,6 +84,7 @@ int main()
 
        #pragma omp taskwait
     }
+// Snippet end
 
     t_end = omp_get_wtime();
 
@@ -168,4 +169,3 @@ int main()
 
     return 0;
 }
-// Snippet end

Publications/GPU-Opt-Guide/OpenMP/22_mkl_dispatch/dgemm_target_variant_dispatch_c.cpp

@@ -4,7 +4,6 @@
 // SPDX-License-Identifier: MIT
 // =============================================================
 // clang-format off
-// Snippet begin
 #include <stdio.h>
 #include <stdlib.h>
 #include <math.h>
@@ -61,6 +60,7 @@ int main()
 
     printf (" Computing matrix product using Intel oneMKL dgemm function via CBLAS interface \n\n");
 
+// Snippet begin
     #pragma omp target data map(to: A[0:m*k], B[0:k*n]) map(tofrom: C[0:m*n])
     {
        #pragma omp target variant dispatch use_device_ptr(A, B, C)
@@ -69,6 +69,7 @@ int main()
                       m, n, k, alpha, A, k, B, n, beta, C, n);
        }
     }
+// Snippet end
 
     printf ("\n Top left corner of matrix C: \n");
     for (i=0; i<min(m,6); i++) {
@@ -117,7 +118,7 @@ int main()
     mkl_free(A);
     mkl_free(B);
     mkl_free(C);
+    mkl_free(C_fl);
 
     return fail;
 }
-// Snippet end