瀏覽代碼

added license

Mozhgan K. Chimeh 4 年之前
父節點
當前提交
c71fcf096f
共有 42 個文件被更改,包括 758 次插入636 次删除
  1. 5 1
      hpc/nways/Dockerfile
  2. 1 1
      hpc/nways/nways_labs/nways_MD/English/C/jupyter_notebook/cudac/nways_cuda.ipynb
  3. 3 3
      hpc/nways/nways_labs/nways_MD/English/C/jupyter_notebook/openacc/nways_openacc.ipynb
  4. 10 10
      hpc/nways/nways_labs/nways_MD/English/C/jupyter_notebook/openacc/nways_openacc_opt.ipynb
  5. 3 3
      hpc/nways/nways_labs/nways_MD/English/C/jupyter_notebook/openmp/nways_openmp.ipynb
  6. 3 3
      hpc/nways/nways_labs/nways_MD/English/C/jupyter_notebook/serial/rdf_overview.ipynb
  7. 16 0
      hpc/nways/nways_labs/nways_MD/English/C/jupyter_notebook/stdpar/nways_stdpar.ipynb
  8. 40 34
      hpc/nways/nways_labs/nways_MD/English/C/source_code/cudac/SOLUTION/dcdread.h
  9. 1 0
      hpc/nways/nways_labs/nways_MD/English/C/source_code/cudac/SOLUTION/rdf.cu
  10. 40 34
      hpc/nways/nways_labs/nways_MD/English/C/source_code/cudac/dcdread.h
  11. 1 0
      hpc/nways/nways_labs/nways_MD/English/C/source_code/cudac/rdf.cu
  12. 1 0
      hpc/nways/nways_labs/nways_MD/English/C/source_code/cudac/vector_addition_gpu_block_only.cu
  13. 1 0
      hpc/nways/nways_labs/nways_MD/English/C/source_code/cudac/vector_addition_gpu_thread_block.cu
  14. 1 0
      hpc/nways/nways_labs/nways_MD/English/C/source_code/cudac/vector_addition_gpu_thread_only.cu
  15. 129 119
      hpc/nways/nways_labs/nways_MD/English/C/source_code/kokkos/SOLUTION/rdf.cpp
  16. 40 34
      hpc/nways/nways_labs/nways_MD/English/C/source_code/kokkos/dcdread.h
  17. 121 110
      hpc/nways/nways_labs/nways_MD/English/C/source_code/kokkos/rdf.cpp
  18. 40 34
      hpc/nways/nways_labs/nways_MD/English/C/source_code/openacc/SOLUTION/dcdread.h
  19. 1 0
      hpc/nways/nways_labs/nways_MD/English/C/source_code/openacc/SOLUTION/rdf.cpp
  20. 1 0
      hpc/nways/nways_labs/nways_MD/English/C/source_code/openacc/SOLUTION/rdf_collapse.cpp
  21. 2 1
      hpc/nways/nways_labs/nways_MD/English/C/source_code/openacc/SOLUTION/rdf_data_directive.cpp
  22. 1 0
      hpc/nways/nways_labs/nways_MD/English/C/source_code/openacc/SOLUTION/rdf_gang.cpp
  23. 1 0
      hpc/nways/nways_labs/nways_MD/English/C/source_code/openacc/SOLUTION/rdf_gang_vector.cpp
  24. 1 0
      hpc/nways/nways_labs/nways_MD/English/C/source_code/openacc/SOLUTION/rdf_gang_vector_length.cpp
  25. 1 0
      hpc/nways/nways_labs/nways_MD/English/C/source_code/openacc/SOLUTION/rdf_gang_vector_worker.cpp
  26. 1 0
      hpc/nways/nways_labs/nways_MD/English/C/source_code/openacc/SOLUTION/rdf_kernel_directive.cpp
  27. 122 112
      hpc/nways/nways_labs/nways_MD/English/C/source_code/openacc/SOLUTION/rdf_parallel_directive.cpp
  28. 40 34
      hpc/nways/nways_labs/nways_MD/English/C/source_code/openacc/dcdread.h
  29. 1 0
      hpc/nways/nways_labs/nways_MD/English/C/source_code/openacc/rdf.cpp
  30. 40 34
      hpc/nways/nways_labs/nways_MD/English/C/source_code/openmp/SOLUTION/dcdread.h
  31. 1 0
      hpc/nways/nways_labs/nways_MD/English/C/source_code/openmp/SOLUTION/rdf.cpp
  32. 1 0
      hpc/nways/nways_labs/nways_MD/English/C/source_code/openmp/SOLUTION/rdf_offload.cpp
  33. 1 0
      hpc/nways/nways_labs/nways_MD/English/C/source_code/openmp/SOLUTION/rdf_offload_collapse.cpp
  34. 1 0
      hpc/nways/nways_labs/nways_MD/English/C/source_code/openmp/dcdread.h
  35. 1 0
      hpc/nways/nways_labs/nways_MD/English/C/source_code/openmp/rdf.cpp
  36. 40 34
      hpc/nways/nways_labs/nways_MD/English/C/source_code/serial/dcdread.h
  37. 1 0
      hpc/nways/nways_labs/nways_MD/English/C/source_code/serial/rdf.cpp
  38. 40 34
      hpc/nways/nways_labs/nways_MD/English/C/source_code/stdpar/SOLUTION/dcdread.h
  39. 1 0
      hpc/nways/nways_labs/nways_MD/English/C/source_code/stdpar/SOLUTION/rdf.cpp
  40. 1 0
      hpc/nways/nways_labs/nways_MD/English/C/source_code/stdpar/dcdread.h
  41. 1 0
      hpc/nways/nways_labs/nways_MD/English/C/source_code/stdpar/rdf.cpp
  42. 1 1
      hpc/nways/nways_labs/nways_start.ipynb

+ 5 - 1
hpc/nways/Dockerfile

@@ -9,8 +9,9 @@ FROM nvcr.io/nvidia/nvhpc:20.11-devel-cuda_multi-ubuntu20.04
 RUN apt-get -y update && \
 RUN apt-get -y update && \
         DEBIAN_FRONTEND=noninteractive apt-get -yq install --no-install-recommends python3-pip python3-setuptools nginx zip make build-essential libtbb-dev && \
         DEBIAN_FRONTEND=noninteractive apt-get -yq install --no-install-recommends python3-pip python3-setuptools nginx zip make build-essential libtbb-dev && \
         rm -rf /var/lib/apt/lists/* && \
         rm -rf /var/lib/apt/lists/* && \
+        pip3 install --upgrade pip &&\
         pip3 install --no-cache-dir jupyter &&\
         pip3 install --no-cache-dir jupyter &&\
-        mkdir -p /home/openacc/labs
+        pip3 install gdown
 
 
 ############################################
 ############################################
 # NVIDIA nsight-systems-2020.5.1 ,nsight-compute-2
 # NVIDIA nsight-systems-2020.5.1 ,nsight-compute-2
@@ -26,6 +27,9 @@ RUN apt-get update -y && \
 
 
 RUN DEBIAN_FRONTEND=noninteractive apt-get install -y --no-install-recommends nsight-systems-2020.5.1 nsight-compute-2020.2.1 
 RUN DEBIAN_FRONTEND=noninteractive apt-get install -y --no-install-recommends nsight-systems-2020.5.1 nsight-compute-2020.2.1 
 
 
+# TO COPY the data
+COPY nways_labs/ /labs/
+
 RUN python3 /labs/nways_MD/English/C/source_code/dataset.py
 RUN python3 /labs/nways_MD/English/C/source_code/dataset.py
 
 
 #################################################
 #################################################

+ 1 - 1
hpc/nways/nways_labs/nways_MD/English/C/jupyter_notebook/cudac/nways_cuda.ipynb

@@ -308,7 +308,7 @@
     "<img src=\"../images/cuda_profile.png\">\n",
     "<img src=\"../images/cuda_profile.png\">\n",
     "\n",
     "\n",
     "\n",
     "\n",
-    "Feel free to checkout the [solution](../source_code/cudac/SOLUTION/rdf.cu) to help you understand better or compare your implementation with the sample solution.\n",
+    "Feel free to checkout the [solution](../../source_code/cudac/SOLUTION/rdf.cu) to help you understand better or compare your implementation with the sample solution.\n",
     "\n",
     "\n",
     "\n",
     "\n",
     "# CUDA C Analysis\n",
     "# CUDA C Analysis\n",

+ 3 - 3
hpc/nways/nways_labs/nways_MD/English/C/jupyter_notebook/openacc/nways_openacc.ipynb

@@ -286,7 +286,7 @@
     "\n",
     "\n",
     "<img src=\"../images/openacc_construct.png\" width=\"80%\" height=\"80%\">\n",
     "<img src=\"../images/openacc_construct.png\" width=\"80%\" height=\"80%\">\n",
     "\n",
     "\n",
-    "Feel free to checkout the [solution](../source_code/openacc/SOLUTION/rdf_parallel_directive.cpp) to help you understand better."
+    "Feel free to checkout the [solution](../../source_code/openacc/SOLUTION/rdf_parallel_directive.cpp) to help you understand better."
    ]
    ]
   },
   },
   {
   {
@@ -457,7 +457,7 @@
    "source": [
    "source": [
     "Let's checkout the profiler's report. [Download the profiler output](../../source_code/openacc/rdf_kernel.qdrep) and open it via the GUI. Checkout the OpenACC row and hover over OpenACC constructs to see if the detail looks different from when you use parallel directives. Compare the profiler report with the previous section.\n",
     "Let's checkout the profiler's report. [Download the profiler output](../../source_code/openacc/rdf_kernel.qdrep) and open it via the GUI. Checkout the OpenACC row and hover over OpenACC constructs to see if the detail looks different from when you use parallel directives. Compare the profiler report with the previous section.\n",
     "\n",
     "\n",
-    "Feel free to checkout the [solution](../source_code/openacc/SOLUTION/rdf_kernel_directive.cpp) to help you understand better."
+    "Feel free to checkout the [solution](../../source_code/openacc/SOLUTION/rdf_kernel_directive.cpp) to help you understand better."
    ]
    ]
   },
   },
   {
   {
@@ -609,7 +609,7 @@
     "\n",
     "\n",
     "Have a look at the data movements annotated with green color and compare it with the previous versions. We have accelerated the application and reduced the execution time by eliminating the unnecessary data transfers between CPU and GPU.\n",
     "Have a look at the data movements annotated with green color and compare it with the previous versions. We have accelerated the application and reduced the execution time by eliminating the unnecessary data transfers between CPU and GPU.\n",
     "\n",
     "\n",
-    "Feel free to checkout the [solution](../source_code/openacc/SOLUTION/rdf_data_directive.cpp) to help you understand better."
+    "Feel free to checkout the [solution](../../source_code/openacc/SOLUTION/rdf_data_directive.cpp) to help you understand better."
    ]
    ]
   },
   },
   {
   {

+ 10 - 10
hpc/nways/nways_labs/nways_MD/English/C/jupyter_notebook/openacc/nways_openacc_opt.ipynb

@@ -35,9 +35,9 @@
    "source": [
    "source": [
     "In this section, we will optimize the parallel [RDF](../serial/rdf_overview.ipynb) application using OpenACC. Before we begin, feel free to have a look at the parallel version of the code and inspect it once again. \n",
     "In this section, we will optimize the parallel [RDF](../serial/rdf_overview.ipynb) application using OpenACC. Before we begin, feel free to have a look at the parallel version of the code and inspect it once again. \n",
     "\n",
     "\n",
-    "[RDF Parallel Code](../../source_code/solution/rdf_data_directive.cpp)\n",
+    "[RDF Parallel Code](../../source_code/openacc/SOLUTION/rdf_data_directive.cpp)\n",
     "\n",
     "\n",
-    "[File Reader](../../source_code/solution/dcdread.h)\n",
+    "[File Reader](../../source_code/openacc/SOLUTION/dcdread.h)\n",
     "\n",
     "\n",
     "Now, let's compile, and profile it with Nsight Systems first."
     "Now, let's compile, and profile it with Nsight Systems first."
    ]
    ]
@@ -49,7 +49,7 @@
    "outputs": [],
    "outputs": [],
    "source": [
    "source": [
     "#compile the parallel code for Tesla GPU\n",
     "#compile the parallel code for Tesla GPU\n",
-    "!cd ../../source_code/openacc && pgc++ -acc -O3 -w -ldl -ta=tesla:lineinfo  -Minfo=accel -o rdf SOLUTION/rdf_data_directive.cpp -I/opt/nvidia/hpc_sdk/Linux_x86_64/20.9/cuda/11.0/include "
+    "!cd ../../source_code/openacc && nvc++ -acc -ta=tesla:lineinfo  -Minfo=accel -o rdf SOLUTION/rdf_data_directive.cpp "
    ]
    ]
   },
   },
   {
   {
@@ -66,7 +66,7 @@
    "cell_type": "markdown",
    "cell_type": "markdown",
    "metadata": {},
    "metadata": {},
    "source": [
    "source": [
-    "Let's checkout the profiler's report. [Download the profiler output](../../source_code/openacc/SOLUTION/rdf_parallel.qdrep) and open it via the GUI."
+    "Let's checkout the profiler's report. [Download the profiler output](../../source_code/openacc/rdf_parallel.qdrep) and open it via the GUI."
    ]
    ]
   },
   },
   {
   {
@@ -161,7 +161,7 @@
     "\n",
     "\n",
     "How much this optimization will speed-up the code will vary according to the application and the target accelerator, but it is not uncommon to see large speed-ups by using collapse on loop nests. \n",
     "How much this optimization will speed-up the code will vary according to the application and the target accelerator, but it is not uncommon to see large speed-ups by using collapse on loop nests. \n",
     "\n",
     "\n",
-    "Feel free to checkout the [solution](../source_code/openacc/SOLUTION/rdf_gang_vector_length.cpp) to help you understand better."
+    "Feel free to checkout the [solution](../../source_code/openacc/SOLUTION/rdf_gang_vector_length.cpp) to help you understand better."
    ]
    ]
   },
   },
   {
   {
@@ -245,7 +245,7 @@
     "\n",
     "\n",
     "How much this optimization will speed-up the code will vary according to the application and the target accelerator, but it is not uncommon to see large speed-ups by using collapse on loop nests. \n",
     "How much this optimization will speed-up the code will vary according to the application and the target accelerator, but it is not uncommon to see large speed-ups by using collapse on loop nests. \n",
     "\n",
     "\n",
-    "Feel free to checkout the [solution](../source_code/openacc/SOLUTION/rdf_collapse.cpp) to help you understand better.\n",
+    "Feel free to checkout the [solution](../../source_code/openacc/SOLUTION/rdf_collapse.cpp) to help you understand better.\n",
     "\n",
     "\n",
     "\n",
     "\n",
     "There is another clause which may be useful in optimizing loops. *Tile* clause breaks down the next loops into tiles before parallelising and it promotes data locality as the device can then use data from nearby tiles. \n",
     "There is another clause which may be useful in optimizing loops. *Tile* clause breaks down the next loops into tiles before parallelising and it promotes data locality as the device can then use data from nearby tiles. \n",
@@ -263,7 +263,7 @@
     "\n",
     "\n",
     "We did not cover this in the labs but this is something you can explore and compare with the explained methods of loop optimization.\n",
     "We did not cover this in the labs but this is something you can explore and compare with the explained methods of loop optimization.\n",
     "\n",
     "\n",
-    "Now, let's dig deeper and profile the kernel with the Nsight Compute profiler. To do this, let's use the [solution](../source_code/openacc/SOLUTION/rdf_collapse.cpp) as a reference to get a similar report from Nsight Compute. Run the application, and profile it with the Nsight Systems first."
+    "Now, let's dig deeper and profile the kernel with the Nsight Compute profiler. To do this, let's use the [solution](../../source_code/openacc/SOLUTION/rdf_collapse.cpp) as a reference to get a similar report from Nsight Compute. Run the application, and profile it with the Nsight Systems first."
    ]
    ]
   },
   },
   {
   {
@@ -273,7 +273,7 @@
    "outputs": [],
    "outputs": [],
    "source": [
    "source": [
     "#compile the solution for Tesla GPU\n",
     "#compile the solution for Tesla GPU\n",
-    "!cd ../../source_code/openacc && pgc++ -acc -O3 -w -ldl -ta=tesla:lineinfo  -Minfo=accel -o rdf SOLUTION/rdf_collapse.cpp -I/opt/nvidia/hpc_sdk/Linux_x86_64/20.9/cuda/11.0/include "
+    "!cd ../../source_code/openacc && nvc++ -acc -ta=tesla:lineinfo  -Minfo=accel -o rdf SOLUTION/rdf_collapse.cpp "
    ]
    ]
   },
   },
   {
   {
@@ -290,7 +290,7 @@
    "cell_type": "markdown",
    "cell_type": "markdown",
    "metadata": {},
    "metadata": {},
    "source": [
    "source": [
-    "Let's checkout the profiler's report. [Download the profiler output](../../source_code/openacc/SOLUTION/rdf_collapse_solution.qdrep) and open it via the GUI. Now, right click on the kernel `pair_gpu_182_gpu` and click on \"Analyze the Selected Kernel with NVIDIA Nsight Compute\" (see below screenshot). \n",
+    "Let's checkout the profiler's report. [Download the profiler output](../../source_code/openacc/rdf_collapse_solution.qdrep) and open it via the GUI. Now, right click on the kernel `pair_gpu_182_gpu` and click on \"Analyze the Selected Kernel with NVIDIA Nsight Compute\" (see below screenshot). \n",
     "\n",
     "\n",
     "<img src=\"../images/compute_analyz.png\">\n",
     "<img src=\"../images/compute_analyz.png\">\n",
     "\n",
     "\n",
@@ -319,7 +319,7 @@
    "cell_type": "markdown",
    "cell_type": "markdown",
    "metadata": {},
    "metadata": {},
    "source": [
    "source": [
-    "Let's checkout the Nsight Compute profiler's report together. [Download the profiler output](../../source_code/openacc/rdf_collapse.ncu-rep) and open it via the GUI. Let's checkout the first section called \"GPU Speed Of Light\". This section gives an overview of the utilization for compute and memory resources on the GPU. As you can see from the below screenshot, the Speed of Light (SOL) reports the achieved percentage of utilization of 30.58% for SM and 19.47% for memory. \n",
+    "Let's checkout the Nsight Compute profiler's report together. [Download the profiler output](../../source_code/openacc/rdf_collapse_solution.ncu-rep) and open it via the GUI. Let's checkout the first section called \"GPU Speed Of Light\". This section gives an overview of the utilization for compute and memory resources on the GPU. As you can see from the below screenshot, the Speed of Light (SOL) reports the achieved percentage of utilization of 30.58% for SM and 19.47% for memory. \n",
     "\n",
     "\n",
     "\n",
     "\n",
     "<img src=\"../images/sol.png\">\n",
     "<img src=\"../images/sol.png\">\n",

+ 3 - 3
hpc/nways/nways_labs/nways_MD/English/C/jupyter_notebook/openmp/nways_openmp.ipynb

@@ -323,11 +323,11 @@
    "cell_type": "markdown",
    "cell_type": "markdown",
    "metadata": {},
    "metadata": {},
    "source": [
    "source": [
-    "Let's checkout the profiler's report. [Download the profiler output](../source_code/openmp/rdf_multicore.qdrep) and open it via the GUI. Have a look at the example expected profiler report below:\n",
+    "Let's checkout the profiler's report. [Download the profiler output](../../source_code/openmp/rdf_multicore.qdrep) and open it via the GUI. Have a look at the example expected profiler report below:\n",
     "\n",
     "\n",
     "<img src=\"../images/openmp_multicore.png\">\n",
     "<img src=\"../images/openmp_multicore.png\">\n",
     "\n",
     "\n",
-    "Feel free to checkout the [solution](../source_code/openmp/SOLUTION/rdf_offload.cpp) to help you understand better."
+    "Feel free to checkout the [solution](../../source_code/openmp/SOLUTION/rdf_offload.cpp) to help you understand better."
    ]
    ]
   },
   },
   {
   {
@@ -527,7 +527,7 @@
     "\n",
     "\n",
     "Compare the execution time for the `Pair_Calculation` from the NVTX row under CUDA row (annotated in Red rectangle in the example screenshot) with the previous section. It is clear the using collapse clause improved the performance by extracting more parallelisim.\n",
     "Compare the execution time for the `Pair_Calculation` from the NVTX row under CUDA row (annotated in Red rectangle in the example screenshot) with the previous section. It is clear the using collapse clause improved the performance by extracting more parallelisim.\n",
     "\n",
     "\n",
-    "Feel free to checkout the [solution](../source_code/openmp/SOLUTION/rdf_offload_collapse.cpp) to help you understand better."
+    "Feel free to checkout the [solution](../../source_code/openmp/SOLUTION/rdf_offload_collapse.cpp) to help you understand better."
    ]
    ]
   },
   },
   {
   {

+ 3 - 3
hpc/nways/nways_labs/nways_MD/English/C/jupyter_notebook/serial/rdf_overview.ipynb

@@ -83,13 +83,13 @@
     "\n",
     "\n",
     "\n",
     "\n",
     "# Links and Resources\n",
     "# Links and Resources\n",
-    "[OpenACC API guide](https://www.openacc.org/sites/default/files/inline-files/OpenACC%20API%202.6%20Reference%20Guide.pdf)\n",
+    "<!--[OpenACC API guide](https://www.openacc.org/sites/default/files/inline-files/OpenACC%20API%202.6%20Reference%20Guide.pdf)-->\n",
     "\n",
     "\n",
     "[NVIDIA Nsight System](https://docs.nvidia.com/nsight-systems/)\n",
     "[NVIDIA Nsight System](https://docs.nvidia.com/nsight-systems/)\n",
     "\n",
     "\n",
-    "[NVIDIA Nsight Compute](https://developer.nvidia.com/nsight-compute)\n",
+    "<!--[NVIDIA Nsight Compute](https://developer.nvidia.com/nsight-compute)-->\n",
     "\n",
     "\n",
-    "[CUDA Toolkit Download](https://developer.nvidia.com/cuda-downloads)\n",
+    "<!--[CUDA Toolkit Download](https://developer.nvidia.com/cuda-downloads)-->\n",
     "\n",
     "\n",
     "[Profiling timelines with NVTX](https://devblogs.nvidia.com/cuda-pro-tip-generate-custom-application-profile-timelines-nvtx/)\n",
     "[Profiling timelines with NVTX](https://devblogs.nvidia.com/cuda-pro-tip-generate-custom-application-profile-timelines-nvtx/)\n",
     "\n",
     "\n",

+ 16 - 0
hpc/nways/nways_labs/nways_MD/English/C/jupyter_notebook/stdpar/nways_stdpar.ipynb

@@ -4,6 +4,22 @@
    "cell_type": "markdown",
    "cell_type": "markdown",
    "metadata": {},
    "metadata": {},
    "source": [
    "source": [
+    "Before we begin, let's execute the cell below to display information about the CUDA driver and GPUs running on the server by running the `nvidia-smi` command. To do this, execute the cell block below by giving it focus (clicking on it with your mouse), and hitting Ctrl-Enter, or pressing the play button in the toolbar above. If all goes well, you should see some output returned below the grey cell."
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "!nvidia-smi"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
     "## Learning objectives\n",
     "## Learning objectives\n",
     "The **goal** of this lab is to:\n",
     "The **goal** of this lab is to:\n",
     "\n",
     "\n",

+ 40 - 34
hpc/nways/nways_labs/nways_MD/English/C/source_code/cudac/SOLUTION/dcdread.h

@@ -1,43 +1,49 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved.
 using namespace std;
 using namespace std;
 
 
-void dcdreadhead(int *natom, int *nframes, std::istream &infile) {
+void dcdreadhead(int *natom, int *nframes, std::istream &infile)
+{
 
 
- infile.seekg(8,ios::beg);
- infile.read((char*)nframes, sizeof(int));
- infile.seekg(64*4,ios::cur);
- infile.read((char*)natom, sizeof(int));
- infile.seekg(1*8,ios::cur);
- return;
+    infile.seekg(8, ios::beg);
+    infile.read((char *)nframes, sizeof(int));
+    infile.seekg(64 * 4, ios::cur);
+    infile.read((char *)natom, sizeof(int));
+    infile.seekg(1 * 8, ios::cur);
+    return;
 }
 }
 
 
 void dcdreadframe(double *x, double *y, double *z, std::istream &infile,
 void dcdreadframe(double *x, double *y, double *z, std::istream &infile,
-    int natom, double& xbox,double& ybox,double& zbox){
+                  int natom, double &xbox, double &ybox, double &zbox)
+{
 
 
-  double d[6];
-  for (int i=0;i<6;i++) {
-      infile.read((char*)&d[i], sizeof(double));
-  }
-  xbox=d[0];
-  ybox=d[2];
-  zbox=d[5];
-  float a,b,c;
-  infile.seekg(1*8,ios::cur);
-  for (int i=0;i<natom;i++){
-      infile.read((char*)&a, sizeof(float));
-      x[i]=a;
-  }
-  infile.seekg(1*8,ios::cur);
-  for (int i=0;i<natom;i++) {
-      infile.read((char*)&b, sizeof(float));
-      y[i]=b;
-  }
-  infile.seekg(1*8,ios::cur);
-  for (int i=0;i<natom;i++) {
-      infile.read((char*)&c, sizeof(float));
-      z[i]=c;
-  }
-  infile.seekg(1*8,ios::cur);
+    double d[6];
+    for (int i = 0; i < 6; i++)
+    {
+        infile.read((char *)&d[i], sizeof(double));
+    }
+    xbox = d[0];
+    ybox = d[2];
+    zbox = d[5];
+    float a, b, c;
+    infile.seekg(1 * 8, ios::cur);
+    for (int i = 0; i < natom; i++)
+    {
+        infile.read((char *)&a, sizeof(float));
+        x[i] = a;
+    }
+    infile.seekg(1 * 8, ios::cur);
+    for (int i = 0; i < natom; i++)
+    {
+        infile.read((char *)&b, sizeof(float));
+        y[i] = b;
+    }
+    infile.seekg(1 * 8, ios::cur);
+    for (int i = 0; i < natom; i++)
+    {
+        infile.read((char *)&c, sizeof(float));
+        z[i] = c;
+    }
+    infile.seekg(1 * 8, ios::cur);
 
 
-  return;
+    return;
 }
 }
-           

+ 1 - 0
hpc/nways/nways_labs/nways_MD/English/C/source_code/cudac/SOLUTION/rdf.cu

@@ -1,3 +1,4 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved. 
 #include <stdio.h>
 #include <stdio.h>
 #include <iostream>
 #include <iostream>
 #include <fstream>
 #include <fstream>

+ 40 - 34
hpc/nways/nways_labs/nways_MD/English/C/source_code/cudac/dcdread.h

@@ -1,43 +1,49 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved.
 using namespace std;
 using namespace std;
 
 
-void dcdreadhead(int *natom, int *nframes, std::istream &infile) {
+void dcdreadhead(int *natom, int *nframes, std::istream &infile)
+{
 
 
- infile.seekg(8,ios::beg);
- infile.read((char*)nframes, sizeof(int));
- infile.seekg(64*4,ios::cur);
- infile.read((char*)natom, sizeof(int));
- infile.seekg(1*8,ios::cur);
- return;
+    infile.seekg(8, ios::beg);
+    infile.read((char *)nframes, sizeof(int));
+    infile.seekg(64 * 4, ios::cur);
+    infile.read((char *)natom, sizeof(int));
+    infile.seekg(1 * 8, ios::cur);
+    return;
 }
 }
 
 
 void dcdreadframe(double *x, double *y, double *z, std::istream &infile,
 void dcdreadframe(double *x, double *y, double *z, std::istream &infile,
-    int natom, double& xbox,double& ybox,double& zbox){
+                  int natom, double &xbox, double &ybox, double &zbox)
+{
 
 
-  double d[6];
-  for (int i=0;i<6;i++) {
-      infile.read((char*)&d[i], sizeof(double));
-  }
-  xbox=d[0];
-  ybox=d[2];
-  zbox=d[5];
-  float a,b,c;
-  infile.seekg(1*8,ios::cur);
-  for (int i=0;i<natom;i++){
-      infile.read((char*)&a, sizeof(float));
-      x[i]=a;
-  }
-  infile.seekg(1*8,ios::cur);
-  for (int i=0;i<natom;i++) {
-      infile.read((char*)&b, sizeof(float));
-      y[i]=b;
-  }
-  infile.seekg(1*8,ios::cur);
-  for (int i=0;i<natom;i++) {
-      infile.read((char*)&c, sizeof(float));
-      z[i]=c;
-  }
-  infile.seekg(1*8,ios::cur);
+    double d[6];
+    for (int i = 0; i < 6; i++)
+    {
+        infile.read((char *)&d[i], sizeof(double));
+    }
+    xbox = d[0];
+    ybox = d[2];
+    zbox = d[5];
+    float a, b, c;
+    infile.seekg(1 * 8, ios::cur);
+    for (int i = 0; i < natom; i++)
+    {
+        infile.read((char *)&a, sizeof(float));
+        x[i] = a;
+    }
+    infile.seekg(1 * 8, ios::cur);
+    for (int i = 0; i < natom; i++)
+    {
+        infile.read((char *)&b, sizeof(float));
+        y[i] = b;
+    }
+    infile.seekg(1 * 8, ios::cur);
+    for (int i = 0; i < natom; i++)
+    {
+        infile.read((char *)&c, sizeof(float));
+        z[i] = c;
+    }
+    infile.seekg(1 * 8, ios::cur);
 
 
-  return;
+    return;
 }
 }
-           

+ 1 - 0
hpc/nways/nways_labs/nways_MD/English/C/source_code/cudac/rdf.cu

@@ -1,3 +1,4 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved. 
 #include <stdio.h>
 #include <stdio.h>
 #include <iostream>
 #include <iostream>
 #include <fstream>
 #include <fstream>

+ 1 - 0
hpc/nways/nways_labs/nways_MD/English/C/source_code/cudac/vector_addition_gpu_block_only.cu

@@ -1,3 +1,4 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved. 
 #include<stdio.h>
 #include<stdio.h>
 #include<stdlib.h>
 #include<stdlib.h>
 
 

+ 1 - 0
hpc/nways/nways_labs/nways_MD/English/C/source_code/cudac/vector_addition_gpu_thread_block.cu

@@ -1,3 +1,4 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved. 
 #include<stdio.h>
 #include<stdio.h>
 #include<stdlib.h>
 #include<stdlib.h>
 
 

+ 1 - 0
hpc/nways/nways_labs/nways_MD/English/C/source_code/cudac/vector_addition_gpu_thread_only.cu

@@ -1,3 +1,4 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved. 
 #include<stdio.h>
 #include<stdio.h>
 #include<stdlib.h>
 #include<stdlib.h>
 
 

+ 129 - 119
hpc/nways/nways_labs/nways_MD/English/C/source_code/kokkos/SOLUTION/rdf.cpp

@@ -1,3 +1,4 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved.
 #include <stdio.h>
 #include <stdio.h>
 #include <iostream>
 #include <iostream>
 #include <fstream>
 #include <fstream>
@@ -10,85 +11,90 @@
 #include <Kokkos_Core.hpp>
 #include <Kokkos_Core.hpp>
 #include <nvToolsExt.h>
 #include <nvToolsExt.h>
 
 
-int l_round(float num); 
+int l_round(float num);
 
 
 typedef Kokkos::View<double *> view_type_double;
 typedef Kokkos::View<double *> view_type_double;
 typedef Kokkos::View<unsigned long long int *> view_type_long;
 typedef Kokkos::View<unsigned long long int *> view_type_long;
 typedef view_type_double::HostMirror host_view_type_double;
 typedef view_type_double::HostMirror host_view_type_double;
 typedef view_type_long::HostMirror host_view_type_long;
 typedef view_type_long::HostMirror host_view_type_long;
 
 
-void pair_gpu(view_type_double d_x, view_type_double d_y, view_type_double d_z, 
-		view_type_long d_g2, int numatm, int nconf, 
-		const double xbox, const double ybox, const double zbox,
-		int d_bin);
+void pair_gpu(view_type_double d_x, view_type_double d_y, view_type_double d_z,
+			  view_type_long d_g2, int numatm, int nconf,
+			  const double xbox, const double ybox, const double zbox,
+			  int d_bin);
 
 
-int main(int argc , char* argv[] )
+int main(int argc, char *argv[])
 {
 {
-	Kokkos::initialize( argc, argv ); {
+	Kokkos::initialize(argc, argv);
+	{
 
 
 		printf("Default  Kokkos execution space %s\n",
 		printf("Default  Kokkos execution space %s\n",
-				typeid(Kokkos::DefaultExecutionSpace).name());
+			   typeid(Kokkos::DefaultExecutionSpace).name());
 
 
-		double xbox,ybox,zbox;
+		double xbox, ybox, zbox;
 		int nbin;
 		int nbin;
-		int numatm,nconf,inconf;
+		int numatm, nconf, inconf;
 		string file;
 		string file;
 
 
 		///////////////////////////////////////////////////////////////
 		///////////////////////////////////////////////////////////////
 
 
 		inconf = 10;
 		inconf = 10;
-		nbin=2000;
+		nbin = 2000;
 		file = "../input/alk.traj.dcd";
 		file = "../input/alk.traj.dcd";
 		///////////////////////////////////////
 		///////////////////////////////////////
 		std::ifstream infile;
 		std::ifstream infile;
 		infile.open(file.c_str());
 		infile.open(file.c_str());
-		if(!infile){
-			cout<<"file "<<file.c_str()<<" not found\n";
+		if (!infile)
+		{
+			cout << "file " << file.c_str() << " not found\n";
 			return 1;
 			return 1;
 		}
 		}
 		assert(infile);
 		assert(infile);
 
 
-
-		ofstream pairfile,stwo;
+		ofstream pairfile, stwo;
 		pairfile.open("RDF.dat");
 		pairfile.open("RDF.dat");
 		stwo.open("Pair_entropy.dat");
 		stwo.open("Pair_entropy.dat");
 
 
 		/////////////////////////////////////////////////////////
 		/////////////////////////////////////////////////////////
-		dcdreadhead(&numatm,&nconf,infile);
-		cout<<"Dcd file has "<< numatm << " atoms and " << nconf << " frames"<<endl;
-		if (inconf>nconf) cout << "nconf is reset to "<< nconf <<endl;
+		dcdreadhead(&numatm, &nconf, infile);
+		cout << "Dcd file has " << numatm << " atoms and " << nconf << " frames" << endl;
+		if (inconf > nconf)
+			cout << "nconf is reset to " << nconf << endl;
 		else
 		else
-		{nconf=inconf;}
-		cout<<"Calculating RDF for " << nconf << " frames"<<endl;
+		{
+			nconf = inconf;
+		}
+		cout << "Calculating RDF for " << nconf << " frames" << endl;
 		////////////////////////////////////////////////////////
 		////////////////////////////////////////////////////////
 
 
-		view_type_double x("x",  nconf*numatm);
-		view_type_double y("y", nconf*numatm);
-		view_type_double z("z",  nconf*numatm);
-		view_type_long g2("g2",  nbin);
+		view_type_double x("x", nconf * numatm);
+		view_type_double y("y", nconf * numatm);
+		view_type_double z("z", nconf * numatm);
+		view_type_long g2("g2", nbin);
 
 
 		host_view_type_double h_x = Kokkos::create_mirror_view(x);
 		host_view_type_double h_x = Kokkos::create_mirror_view(x);
 		host_view_type_double h_y = Kokkos::create_mirror_view(y);
 		host_view_type_double h_y = Kokkos::create_mirror_view(y);
 		host_view_type_double h_z = Kokkos::create_mirror_view(z);
 		host_view_type_double h_z = Kokkos::create_mirror_view(z);
 		host_view_type_long h_g2 = Kokkos::create_mirror_view(g2);
 		host_view_type_long h_g2 = Kokkos::create_mirror_view(g2);
 
 
-
 		/////////reading cordinates//////////////////////////////////////////////
 		/////////reading cordinates//////////////////////////////////////////////
 		nvtxRangePush("Read_File");
 		nvtxRangePush("Read_File");
-		double ax[numatm],ay[numatm],az[numatm];
-		for (int i=0;i<nconf;i++) {
-			dcdreadframe(ax,ay,az,infile,numatm,xbox,ybox,zbox);
-			for (int j=0;j<numatm;j++){
-				h_x(i*numatm+j)=ax[j];
-				h_y(i*numatm+j)=ay[j];
-				h_z(i*numatm+j)=az[j];
+		double ax[numatm], ay[numatm], az[numatm];
+		for (int i = 0; i < nconf; i++)
+		{
+			dcdreadframe(ax, ay, az, infile, numatm, xbox, ybox, zbox);
+			for (int j = 0; j < numatm; j++)
+			{
+				h_x(i * numatm + j) = ax[j];
+				h_y(i * numatm + j) = ay[j];
+				h_z(i * numatm + j) = az[j];
 			}
 			}
 		}
 		}
-		for(int i=0;i<nbin;i++)
+		for (int i = 0; i < nbin; i++)
 			h_g2(0) = 0;
 			h_g2(0) = 0;
 
 
 		nvtxRangePop(); //pop for Reading file
 		nvtxRangePop(); //pop for Reading file
-		cout<<"Reading of input file is completed"<<endl;
+		cout << "Reading of input file is completed" << endl;
 
 
 		nvtxRangePush("Pair_Calculation");
 		nvtxRangePush("Pair_Calculation");
 		Kokkos::deep_copy(x, h_x);
 		Kokkos::deep_copy(x, h_x);
@@ -99,112 +105,116 @@ int main(int argc , char* argv[] )
 		pair_gpu(x, y, z, g2, numatm, nconf, xbox, ybox, zbox, nbin);
 		pair_gpu(x, y, z, g2, numatm, nconf, xbox, ybox, zbox, nbin);
 		Kokkos::deep_copy(h_g2, g2);
 		Kokkos::deep_copy(h_g2, g2);
 		nvtxRangePop(); //Pop for Pair Calculation
 		nvtxRangePop(); //Pop for Pair Calculation
-		double pi=acos(-1.0l);
-		double rho=(numatm)/(xbox*ybox*zbox);
-		double norm=(4.0l*pi*rho)/3.0l;
-		double rl,ru,nideal;
+		double pi = acos(-1.0l);
+		double rho = (numatm) / (xbox * ybox * zbox);
+		double norm = (4.0l * pi * rho) / 3.0l;
+		double rl, ru, nideal;
 		double t_g2[nbin];
 		double t_g2[nbin];
-		double r,gr,lngr,lngrbond,s2=0.0l,s2bond=0.0l;
-		double box=min(xbox,ybox);
-		box=min(box,zbox);
-		double del=box/(2.0l*nbin);
+		double r, gr, lngr, lngrbond, s2 = 0.0l, s2bond = 0.0l;
+		double box = min(xbox, ybox);
+		box = min(box, zbox);
+		double del = box / (2.0l * nbin);
 		nvtxRangePush("Entropy_Calculation");
 		nvtxRangePush("Entropy_Calculation");
-		for (int i=0;i<nbin;i++) {
-			rl=(i)*del;
-			ru=rl+del;
-			nideal=norm*(ru*ru*ru-rl*rl*rl);
-			t_g2[i]=(double)h_g2(i)/((double)nconf*(double)numatm*nideal);
-			r=(i)*del;
-			pairfile<<(i+0.5l)*del<<" "<<t_g2[i]<<endl;
-			if (r<2.0l) {
-				gr=0.0l;
+		for (int i = 0; i < nbin; i++)
+		{
+			rl = (i)*del;
+			ru = rl + del;
+			nideal = norm * (ru * ru * ru - rl * rl * rl);
+			t_g2[i] = (double)h_g2(i) / ((double)nconf * (double)numatm * nideal);
+			r = (i)*del;
+			pairfile << (i + 0.5l) * del << " " << t_g2[i] << endl;
+			if (r < 2.0l)
+			{
+				gr = 0.0l;
 			}
 			}
-			else {
-				gr=t_g2[i];
+			else
+			{
+				gr = t_g2[i];
 			}
 			}
-			if (gr<1e-5) {
-				lngr=0.0l;
+			if (gr < 1e-5)
+			{
+				lngr = 0.0l;
 			}
 			}
-			else {
-				lngr=log(gr);
+			else
+			{
+				lngr = log(gr);
 			}
 			}
 
 
-			if (t_g2[i]<1e-6) {
-				lngrbond=0.0l;
+			if (t_g2[i] < 1e-6)
+			{
+				lngrbond = 0.0l;
 			}
 			}
-			else {
-				lngrbond=log(t_g2[i]);
+			else
+			{
+				lngrbond = log(t_g2[i]);
 			}
 			}
-			s2=s2-2.0l*pi*rho*((gr*lngr)-gr+1.0l)*del*r*r;
-			s2bond=s2bond-2.0l*pi*rho*((t_g2[i]*lngrbond)-t_g2[i]+1.0l)*del*r*r;
-
+			s2 = s2 - 2.0l * pi * rho * ((gr * lngr) - gr + 1.0l) * del * r * r;
+			s2bond = s2bond - 2.0l * pi * rho * ((t_g2[i] * lngrbond) - t_g2[i] + 1.0l) * del * r * r;
 		}
 		}
 		nvtxRangePop(); //Pop for Entropy Calculation
 		nvtxRangePop(); //Pop for Entropy Calculation
-		stwo<<"s2 value is "<<s2<<endl;
-		stwo<<"s2bond value is "<<s2bond<<endl;
+		stwo << "s2 value is " << s2 << endl;
+		stwo << "s2bond value is " << s2bond << endl;
 
 
+		cout << "#Freeing Host memory" << endl;
 
 
-
-		cout<<"#Freeing Host memory"<<endl;
-
-		cout<<"#Number of atoms processed: "<<numatm<<endl<<endl;
-		cout<<"#Number of confs processed: "<<nconf<<endl<<endl;
+		cout << "#Number of atoms processed: " << numatm << endl
+			 << endl;
+		cout << "#Number of confs processed: " << nconf << endl
+			 << endl;
 
 
 	} // Kokkos Initialize ends here
 	} // Kokkos Initialize ends here
 	Kokkos::finalize();
 	Kokkos::finalize();
 	return 0;
 	return 0;
 }
 }
-int l_round(float num) 
-{ 
-	return num < 0 ? num - 0.5 : num + 0.5; 
-} 
-
-void pair_gpu(view_type_double d_x, view_type_double d_y, view_type_double d_z, 
-		view_type_long d_g2, int numatm, int nconf, 
-		const double xbox, const double ybox, const double zbox,
-		int d_bin)
+int l_round(float num)
+{
+	return num < 0 ? num - 0.5 : num + 0.5;
+}
+
+void pair_gpu(view_type_double d_x, view_type_double d_y, view_type_double d_z,
+			  view_type_long d_g2, int numatm, int nconf,
+			  const double xbox, const double ybox, const double zbox,
+			  int d_bin)
 {
 {
 
 
-	printf("\n %d %d ", nconf,numatm);
-	for (int frame=0;frame<nconf;frame++){
+	printf("\n %d %d ", nconf, numatm);
+	for (int frame = 0; frame < nconf; frame++)
+	{
 		printf("\n %d  ", frame);
 		printf("\n %d  ", frame);
-		Kokkos::parallel_for( (numatm*numatm), KOKKOS_LAMBDA (const int index) {
-				int id1 = index / numatm;
-				int id2 = index % numatm;
-				double r,cut,dx,dy,dz;
-				int ig2;
-				double box;
-				int myround;
-				float num;
-				box=min(xbox,ybox);
-				box=min(box,zbox);
-				double del=box/(2.0*d_bin);
-				cut=box*0.5;
-
-				dx=d_x(frame*numatm+id1)-d_x(frame*numatm+id2);
-				dy=d_y(frame*numatm+id1)-d_y(frame*numatm+id2);
-				dz=d_z(frame*numatm+id1)-d_z(frame*numatm+id2);
-
-				num = dx/xbox;
-				myround = num < 0 ? num - 0.5 : num + 0.5;
-				dx=dx-xbox*myround;
-
-				num = dy/ybox;
-				myround = num < 0 ? num - 0.5 : num + 0.5;
-				dy=dy-ybox*myround;
-
-				num = dz/zbox;
-				myround = num < 0 ? num - 0.5 : num + 0.5;
-				dz=dz-zbox*myround;
-				r=sqrtf(dx*dx+dy*dy+dz*dz);
-				if (r<cut) {
-					ig2=(int)(r/del);
-					Kokkos::atomic_increment(&d_g2(ig2));
-				}
+		Kokkos::parallel_for((numatm * numatm), KOKKOS_LAMBDA(const int index) {
+			int id1 = index / numatm;
+			int id2 = index % numatm;
+			double r, cut, dx, dy, dz;
+			int ig2;
+			double box;
+			int myround;
+			float num;
+			box = min(xbox, ybox);
+			box = min(box, zbox);
+			double del = box / (2.0 * d_bin);
+			cut = box * 0.5;
+
+			dx = d_x(frame * numatm + id1) - d_x(frame * numatm + id2);
+			dy = d_y(frame * numatm + id1) - d_y(frame * numatm + id2);
+			dz = d_z(frame * numatm + id1) - d_z(frame * numatm + id2);
+
+			num = dx / xbox;
+			myround = num < 0 ? num - 0.5 : num + 0.5;
+			dx = dx - xbox * myround;
+
+			num = dy / ybox;
+			myround = num < 0 ? num - 0.5 : num + 0.5;
+			dy = dy - ybox * myround;
+
+			num = dz / zbox;
+			myround = num < 0 ? num - 0.5 : num + 0.5;
+			dz = dz - zbox * myround;
+			r = sqrtf(dx * dx + dy * dy + dz * dz);
+			if (r < cut)
+			{
+				ig2 = (int)(r / del);
+				Kokkos::atomic_increment(&d_g2(ig2));
+			}
 		});
 		});
 	}
 	}
 }
 }
-
-
-
-

+ 40 - 34
hpc/nways/nways_labs/nways_MD/English/C/source_code/kokkos/dcdread.h

@@ -1,43 +1,49 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved.
 using namespace std;
 using namespace std;
 
 
-void dcdreadhead(int *natom, int *nframes, std::istream &infile) {
+void dcdreadhead(int *natom, int *nframes, std::istream &infile)
+{
 
 
- infile.seekg(8,ios::beg);
- infile.read((char*)nframes, sizeof(int));
- infile.seekg(64*4,ios::cur);
- infile.read((char*)natom, sizeof(int));
- infile.seekg(1*8,ios::cur);
- return;
+    infile.seekg(8, ios::beg);
+    infile.read((char *)nframes, sizeof(int));
+    infile.seekg(64 * 4, ios::cur);
+    infile.read((char *)natom, sizeof(int));
+    infile.seekg(1 * 8, ios::cur);
+    return;
 }
 }
 
 
 void dcdreadframe(double *x, double *y, double *z, std::istream &infile,
 void dcdreadframe(double *x, double *y, double *z, std::istream &infile,
-    int natom, double& xbox,double& ybox,double& zbox){
+                  int natom, double &xbox, double &ybox, double &zbox)
+{
 
 
-  double d[6];
-  for (int i=0;i<6;i++) {
-      infile.read((char*)&d[i], sizeof(double));
-  }
-  xbox=d[0];
-  ybox=d[2];
-  zbox=d[5];
-  float a,b,c;
-  infile.seekg(1*8,ios::cur);
-  for (int i=0;i<natom;i++){
-      infile.read((char*)&a, sizeof(float));
-      x[i]=a;
-  }
-  infile.seekg(1*8,ios::cur);
-  for (int i=0;i<natom;i++) {
-      infile.read((char*)&b, sizeof(float));
-      y[i]=b;
-  }
-  infile.seekg(1*8,ios::cur);
-  for (int i=0;i<natom;i++) {
-      infile.read((char*)&c, sizeof(float));
-      z[i]=c;
-  }
-  infile.seekg(1*8,ios::cur);
+    double d[6];
+    for (int i = 0; i < 6; i++)
+    {
+        infile.read((char *)&d[i], sizeof(double));
+    }
+    xbox = d[0];
+    ybox = d[2];
+    zbox = d[5];
+    float a, b, c;
+    infile.seekg(1 * 8, ios::cur);
+    for (int i = 0; i < natom; i++)
+    {
+        infile.read((char *)&a, sizeof(float));
+        x[i] = a;
+    }
+    infile.seekg(1 * 8, ios::cur);
+    for (int i = 0; i < natom; i++)
+    {
+        infile.read((char *)&b, sizeof(float));
+        y[i] = b;
+    }
+    infile.seekg(1 * 8, ios::cur);
+    for (int i = 0; i < natom; i++)
+    {
+        infile.read((char *)&c, sizeof(float));
+        z[i] = c;
+    }
+    infile.seekg(1 * 8, ios::cur);
 
 
-  return;
+    return;
 }
 }
-           

+ 121 - 110
hpc/nways/nways_labs/nways_MD/English/C/source_code/kokkos/rdf.cpp

@@ -1,3 +1,4 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved.
 #include <stdio.h>
 #include <stdio.h>
 #include <iostream>
 #include <iostream>
 #include <fstream>
 #include <fstream>
@@ -10,67 +11,71 @@
 #include <Kokkos_Core.hpp> // Note:: Included the Kokkos core library
 #include <Kokkos_Core.hpp> // Note:: Included the Kokkos core library
 #include <nvToolsExt.h>
 #include <nvToolsExt.h>
 
 
-int l_round(float num); 
+int l_round(float num);
 
 
 //Todo: Fill the correct data type and dimensions in the code
 //Todo: Fill the correct data type and dimensions in the code
-typedef Kokkos::View<Fill here > view_type_double;
+typedef Kokkos::View<Fill here> view_type_double;
 typedef Kokkos::View<Fill here> view_type_long;
 typedef Kokkos::View<Fill here> view_type_long;
 
 
 typedef view_type_double::HostMirror host_view_type_double;
 typedef view_type_double::HostMirror host_view_type_double;
 typedef view_type_long::HostMirror host_view_type_long;
 typedef view_type_long::HostMirror host_view_type_long;
 
 
-void pair_gpu(view_type_double d_x, view_type_double d_y, view_type_double d_z, 
-		view_type_long d_g2, int numatm, int nconf, 
-		const double xbox, const double ybox, const double zbox,
-		int d_bin);
+void pair_gpu(view_type_double d_x, view_type_double d_y, view_type_double d_z,
+			  view_type_long d_g2, int numatm, int nconf,
+			  const double xbox, const double ybox, const double zbox,
+			  int d_bin);
 
 
-int main(int argc , char* argv[] )
+int main(int argc, char *argv[])
 {
 {
 	//Note:: We are initailizing the Kokkos library before calling any Kokkos API
 	//Note:: We are initailizing the Kokkos library before calling any Kokkos API
-	Kokkos::initialize( argc, argv ); {
+	Kokkos::initialize(argc, argv);
+	{
 
 
 		//Note: This will print the default execution space with which Kokkos library was built
 		//Note: This will print the default execution space with which Kokkos library was built
 		printf("Default  Kokkos execution space %s\n",
 		printf("Default  Kokkos execution space %s\n",
-				typeid(Kokkos::DefaultExecutionSpace).name());
+			   typeid(Kokkos::DefaultExecutionSpace).name());
 
 
-		double xbox,ybox,zbox;
+		double xbox, ybox, zbox;
 		int nbin;
 		int nbin;
-		int numatm,nconf,inconf;
+		int numatm, nconf, inconf;
 		string file;
 		string file;
 
 
 		///////////////////////////////////////////////////////////////
 		///////////////////////////////////////////////////////////////
 
 
 		inconf = 10;
 		inconf = 10;
-		nbin=2000;
+		nbin = 2000;
 		file = "../input/alk.traj.dcd";
 		file = "../input/alk.traj.dcd";
 		///////////////////////////////////////
 		///////////////////////////////////////
 		std::ifstream infile;
 		std::ifstream infile;
 		infile.open(file.c_str());
 		infile.open(file.c_str());
-		if(!infile){
-			cout<<"file "<<file.c_str()<<" not found\n";
+		if (!infile)
+		{
+			cout << "file " << file.c_str() << " not found\n";
 			return 1;
 			return 1;
 		}
 		}
 		assert(infile);
 		assert(infile);
 
 
-
-		ofstream pairfile,stwo;
+		ofstream pairfile, stwo;
 		pairfile.open("RDF.dat");
 		pairfile.open("RDF.dat");
 		stwo.open("Pair_entropy.dat");
 		stwo.open("Pair_entropy.dat");
 
 
 		/////////////////////////////////////////////////////////
 		/////////////////////////////////////////////////////////
-		dcdreadhead(&numatm,&nconf,infile);
-		cout<<"Dcd file has "<< numatm << " atoms and " << nconf << " frames"<<endl;
-		if (inconf>nconf) cout << "nconf is reset to "<< nconf <<endl;
+		dcdreadhead(&numatm, &nconf, infile);
+		cout << "Dcd file has " << numatm << " atoms and " << nconf << " frames" << endl;
+		if (inconf > nconf)
+			cout << "nconf is reset to " << nconf << endl;
 		else
 		else
-		{nconf=inconf;}
-		cout<<"Calculating RDF for " << nconf << " frames"<<endl;
+		{
+			nconf = inconf;
+		}
+		cout << "Calculating RDF for " << nconf << " frames" << endl;
 		////////////////////////////////////////////////////////
 		////////////////////////////////////////////////////////
 
 
 		//Todo: Fill the correct dimension is view type. This is where the allocation on default Memory space will occur
 		//Todo: Fill the correct dimension is view type. This is where the allocation on default Memory space will occur
-		view_type_double x("x",  Fill here);
+		view_type_double x("x", Fill here);
 		view_type_double y("y", Fill here);
 		view_type_double y("y", Fill here);
-		view_type_double z("z",  Fill here);
-		view_type_long g2("g2",  Fill here);
+		view_type_double z("z", Fill here);
+		view_type_long g2("g2", Fill here);
 
 
 		//Todo : Fill the right mirror image variabe here
 		//Todo : Fill the right mirror image variabe here
 		host_view_type_double h_x = Kokkos::create_mirror_view(x);
 		host_view_type_double h_x = Kokkos::create_mirror_view(x);
@@ -78,144 +83,150 @@ int main(int argc , char* argv[] )
 		host_view_type_double h_z = Kokkos::create_mirror_view(Fill here);
 		host_view_type_double h_z = Kokkos::create_mirror_view(Fill here);
 		host_view_type_long h_g2 = Kokkos::create_mirror_view(Fill here);
 		host_view_type_long h_g2 = Kokkos::create_mirror_view(Fill here);
 
 
-
 		/////////reading cordinates//////////////////////////////////////////////
 		/////////reading cordinates//////////////////////////////////////////////
 		nvtxRangePush("Read_File");
 		nvtxRangePush("Read_File");
-		double ax[numatm],ay[numatm],az[numatm];
-		for (int i=0;i<nconf;i++) {
-			dcdreadframe(ax,ay,az,infile,numatm,xbox,ybox,zbox);
-			for (int j=0;j<numatm;j++){
-				h_x(i*numatm+j)=ax[j];
-				h_y(i*numatm+j)=ay[j];
-				h_z(i*numatm+j)=az[j];
+		double ax[numatm], ay[numatm], az[numatm];
+		for (int i = 0; i < nconf; i++)
+		{
+			dcdreadframe(ax, ay, az, infile, numatm, xbox, ybox, zbox);
+			for (int j = 0; j < numatm; j++)
+			{
+				h_x(i * numatm + j) = ax[j];
+				h_y(i * numatm + j) = ay[j];
+				h_z(i * numatm + j) = az[j];
 			}
 			}
 		}
 		}
-		for(int i=0;i<nbin;i++)
+		for (int i = 0; i < nbin; i++)
 			h_g2(0) = 0;
 			h_g2(0) = 0;
 
 
 		nvtxRangePop(); //pop for Reading file
 		nvtxRangePop(); //pop for Reading file
-		cout<<"Reading of input file is completed"<<endl;
+		cout << "Reading of input file is completed" << endl;
 
 
 		nvtxRangePush("Pair_Calculation");
 		nvtxRangePush("Pair_Calculation");
 		//Todo: Copy from Host to device h_x->x,h_y->y, h_z-> z and h_g2->g2
 		//Todo: Copy from Host to device h_x->x,h_y->y, h_z-> z and h_g2->g2
-		Kokkos::deep_copy( Fill Destination View , Fill Source View );
-		Kokkos::deep_copy(Fill Destination View , Fill Source View );
-		Kokkos::deep_copy(Fill Destination View , Fill Source View );
-		Kokkos::deep_copy(Fill Destination View , Fill Source View );
+		Kokkos::deep_copy(Fill Destination View, Fill Source View);
+		Kokkos::deep_copy(Fill Destination View, Fill Source View);
+		Kokkos::deep_copy(Fill Destination View, Fill Source View);
+		Kokkos::deep_copy(Fill Destination View, Fill Source View);
 		//////////////////////////////////////////////////////////////////////////
 		//////////////////////////////////////////////////////////////////////////
 		pair_gpu(x, y, z, g2, numatm, nconf, xbox, ybox, zbox, nbin);
 		pair_gpu(x, y, z, g2, numatm, nconf, xbox, ybox, zbox, nbin);
 		//Todo: Copy from Device to host g2 -> h_g2 before being used on host
 		//Todo: Copy from Device to host g2 -> h_g2 before being used on host
-		Kokkos::deep_copy(Fill Destination View , Fill Source View );
+		Kokkos::deep_copy(Fill Destination View, Fill Source View);
 		nvtxRangePop(); //Pop for Pair Calculation
 		nvtxRangePop(); //Pop for Pair Calculation
-		double pi=acos(-1.0l);
-		double rho=(numatm)/(xbox*ybox*zbox);
-		double norm=(4.0l*pi*rho)/3.0l;
-		double rl,ru,nideal;
+		double pi = acos(-1.0l);
+		double rho = (numatm) / (xbox * ybox * zbox);
+		double norm = (4.0l * pi * rho) / 3.0l;
+		double rl, ru, nideal;
 		double t_g2[nbin];
 		double t_g2[nbin];
-		double r,gr,lngr,lngrbond,s2=0.0l,s2bond=0.0l;
-		double box=min(xbox,ybox);
-		box=min(box,zbox);
-		double del=box/(2.0l*nbin);
+		double r, gr, lngr, lngrbond, s2 = 0.0l, s2bond = 0.0l;
+		double box = min(xbox, ybox);
+		box = min(box, zbox);
+		double del = box / (2.0l * nbin);
 		nvtxRangePush("Entropy_Calculation");
 		nvtxRangePush("Entropy_Calculation");
-		for (int i=0;i<nbin;i++) {
-			rl=(i)*del;
-			ru=rl+del;
-			nideal=norm*(ru*ru*ru-rl*rl*rl);
-			t_g2[i]=(double)h_g2(i)/((double)nconf*(double)numatm*nideal);
-			r=(i)*del;
-			pairfile<<(i+0.5l)*del<<" "<<t_g2[i]<<endl;
-			if (r<2.0l) {
-				gr=0.0l;
+		for (int i = 0; i < nbin; i++)
+		{
+			rl = (i)*del;
+			ru = rl + del;
+			nideal = norm * (ru * ru * ru - rl * rl * rl);
+			t_g2[i] = (double)h_g2(i) / ((double)nconf * (double)numatm * nideal);
+			r = (i)*del;
+			pairfile << (i + 0.5l) * del << " " << t_g2[i] << endl;
+			if (r < 2.0l)
+			{
+				gr = 0.0l;
 			}
 			}
-			else {
-				gr=t_g2[i];
+			else
+			{
+				gr = t_g2[i];
 			}
 			}
-			if (gr<1e-5) {
-				lngr=0.0l;
+			if (gr < 1e-5)
+			{
+				lngr = 0.0l;
 			}
 			}
-			else {
-				lngr=log(gr);
+			else
+			{
+				lngr = log(gr);
 			}
 			}
 
 
-			if (t_g2[i]<1e-6) {
-				lngrbond=0.0l;
+			if (t_g2[i] < 1e-6)
+			{
+				lngrbond = 0.0l;
 			}
 			}
-			else {
-				lngrbond=log(t_g2[i]);
+			else
+			{
+				lngrbond = log(t_g2[i]);
 			}
 			}
-			s2=s2-2.0l*pi*rho*((gr*lngr)-gr+1.0l)*del*r*r;
-			s2bond=s2bond-2.0l*pi*rho*((t_g2[i]*lngrbond)-t_g2[i]+1.0l)*del*r*r;
-
+			s2 = s2 - 2.0l * pi * rho * ((gr * lngr) - gr + 1.0l) * del * r * r;
+			s2bond = s2bond - 2.0l * pi * rho * ((t_g2[i] * lngrbond) - t_g2[i] + 1.0l) * del * r * r;
 		}
 		}
 		nvtxRangePop(); //Pop for Entropy Calculation
 		nvtxRangePop(); //Pop for Entropy Calculation
-		stwo<<"s2 value is "<<s2<<endl;
-		stwo<<"s2bond value is "<<s2bond<<endl;
+		stwo << "s2 value is " << s2 << endl;
+		stwo << "s2bond value is " << s2bond << endl;
 
 
+		cout << "#Freeing Host memory" << endl;
 
 
-
-		cout<<"#Freeing Host memory"<<endl;
-
-		cout<<"#Number of atoms processed: "<<numatm<<endl<<endl;
-		cout<<"#Number of confs processed: "<<nconf<<endl<<endl;
+		cout << "#Number of atoms processed: " << numatm << endl
+			 << endl;
+		cout << "#Number of confs processed: " << nconf << endl
+			 << endl;
 
 
 	} // Kokkos Initialize ends here
 	} // Kokkos Initialize ends here
 	//Note:: Free up the memory
 	//Note:: Free up the memory
 	Kokkos::finalize();
 	Kokkos::finalize();
 	return 0;
 	return 0;
 }
 }
-int l_round(float num) 
-{ 
-	return num < 0 ? num - 0.5 : num + 0.5; 
-} 
-
-void pair_gpu(view_type_double d_x, view_type_double d_y, view_type_double d_z, 
-		view_type_long d_g2, int numatm, int nconf, 
-		const double xbox, const double ybox, const double zbox,
-		int d_bin)
+int l_round(float num)
 {
 {
+	return num < 0 ? num - 0.5 : num + 0.5;
+}
 
 
-	printf("\n %d %d ", nconf,numatm);
-	for (int frame=0;frame<nconf;frame++){
+void pair_gpu(view_type_double d_x, view_type_double d_y, view_type_double d_z,
+			  view_type_long d_g2, int numatm, int nconf,
+			  const double xbox, const double ybox, const double zbox,
+			  int d_bin)
+{
+
+	printf("\n %d %d ", nconf, numatm);
+	for (int frame = 0; frame < nconf; frame++)
+	{
 		printf("\n %d  ", frame);
 		printf("\n %d  ", frame);
 		//Fill here the pattern we intend to use along with loop size
 		//Fill here the pattern we intend to use along with loop size
-		Kokkos::Fill_Here( Fill the loop size here, KOKKOS_LAMBDA (const int index) {
+		Kokkos::Fill_Here(
+			Fill the loop size here, KOKKOS_LAMBDA(const int index) {
 				int id1 = index / numatm;
 				int id1 = index / numatm;
 				int id2 = index % numatm;
 				int id2 = index % numatm;
-				double r,cut,dx,dy,dz;
+				double r, cut, dx, dy, dz;
 				int ig2;
 				int ig2;
 				double box;
 				double box;
 				int myround;
 				int myround;
 				float num;
 				float num;
-				box=min(xbox,ybox);
-				box=min(box,zbox);
-				double del=box/(2.0*d_bin);
-				cut=box*0.5;
+				box = min(xbox, ybox);
+				box = min(box, zbox);
+				double del = box / (2.0 * d_bin);
+				cut = box * 0.5;
 
 
-				dx=d_x(frame*numatm+id1)-d_x(frame*numatm+id2);
-				dy=d_y(frame*numatm+id1)-d_y(frame*numatm+id2);
-				dz=d_z(frame*numatm+id1)-d_z(frame*numatm+id2);
+				dx = d_x(frame * numatm + id1) - d_x(frame * numatm + id2);
+				dy = d_y(frame * numatm + id1) - d_y(frame * numatm + id2);
+				dz = d_z(frame * numatm + id1) - d_z(frame * numatm + id2);
 
 
-				num = dx/xbox;
+				num = dx / xbox;
 				myround = num < 0 ? num - 0.5 : num + 0.5;
 				myround = num < 0 ? num - 0.5 : num + 0.5;
-				dx=dx-xbox*myround;
+				dx = dx - xbox * myround;
 
 
-				num = dy/ybox;
+				num = dy / ybox;
 				myround = num < 0 ? num - 0.5 : num + 0.5;
 				myround = num < 0 ? num - 0.5 : num + 0.5;
-				dy=dy-ybox*myround;
+				dy = dy - ybox * myround;
 
 
-				num = dz/zbox;
+				num = dz / zbox;
 				myround = num < 0 ? num - 0.5 : num + 0.5;
 				myround = num < 0 ? num - 0.5 : num + 0.5;
-				dz=dz-zbox*myround;
-				r=sqrtf(dx*dx+dy*dy+dz*dz);
-				if (r<cut) {
-					ig2=(int)(r/del);
+				dz = dz - zbox * myround;
+				r = sqrtf(dx * dx + dy * dy + dz * dz);
+				if (r < cut)
+				{
+					ig2 = (int)(r / del);
 					//Note:  We are using a atomic increment here
 					//Note:  We are using a atomic increment here
 					Kokkos::atomic_increment(&d_g2(ig2));
 					Kokkos::atomic_increment(&d_g2(ig2));
 				}
 				}
-		});
+			});
 	}
 	}
 }
 }
-
-
-
-

+ 40 - 34
hpc/nways/nways_labs/nways_MD/English/C/source_code/openacc/SOLUTION/dcdread.h

@@ -1,43 +1,49 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved.
 using namespace std;
 using namespace std;
 
 
-void dcdreadhead(int *natom, int *nframes, std::istream &infile) {
+void dcdreadhead(int *natom, int *nframes, std::istream &infile)
+{
 
 
- infile.seekg(8,ios::beg);
- infile.read((char*)nframes, sizeof(int));
- infile.seekg(64*4,ios::cur);
- infile.read((char*)natom, sizeof(int));
- infile.seekg(1*8,ios::cur);
- return;
+    infile.seekg(8, ios::beg);
+    infile.read((char *)nframes, sizeof(int));
+    infile.seekg(64 * 4, ios::cur);
+    infile.read((char *)natom, sizeof(int));
+    infile.seekg(1 * 8, ios::cur);
+    return;
 }
 }
 
 
 void dcdreadframe(double *x, double *y, double *z, std::istream &infile,
 void dcdreadframe(double *x, double *y, double *z, std::istream &infile,
-    int natom, double& xbox,double& ybox,double& zbox){
+                  int natom, double &xbox, double &ybox, double &zbox)
+{
 
 
-  double d[6];
-  for (int i=0;i<6;i++) {
-      infile.read((char*)&d[i], sizeof(double));
-  }
-  xbox=d[0];
-  ybox=d[2];
-  zbox=d[5];
-  float a,b,c;
-  infile.seekg(1*8,ios::cur);
-  for (int i=0;i<natom;i++){
-      infile.read((char*)&a, sizeof(float));
-      x[i]=a;
-  }
-  infile.seekg(1*8,ios::cur);
-  for (int i=0;i<natom;i++) {
-      infile.read((char*)&b, sizeof(float));
-      y[i]=b;
-  }
-  infile.seekg(1*8,ios::cur);
-  for (int i=0;i<natom;i++) {
-      infile.read((char*)&c, sizeof(float));
-      z[i]=c;
-  }
-  infile.seekg(1*8,ios::cur);
+    double d[6];
+    for (int i = 0; i < 6; i++)
+    {
+        infile.read((char *)&d[i], sizeof(double));
+    }
+    xbox = d[0];
+    ybox = d[2];
+    zbox = d[5];
+    float a, b, c;
+    infile.seekg(1 * 8, ios::cur);
+    for (int i = 0; i < natom; i++)
+    {
+        infile.read((char *)&a, sizeof(float));
+        x[i] = a;
+    }
+    infile.seekg(1 * 8, ios::cur);
+    for (int i = 0; i < natom; i++)
+    {
+        infile.read((char *)&b, sizeof(float));
+        y[i] = b;
+    }
+    infile.seekg(1 * 8, ios::cur);
+    for (int i = 0; i < natom; i++)
+    {
+        infile.read((char *)&c, sizeof(float));
+        z[i] = c;
+    }
+    infile.seekg(1 * 8, ios::cur);
 
 
-  return;
+    return;
 }
 }
-           

+ 1 - 0
hpc/nways/nways_labs/nways_MD/English/C/source_code/openacc/SOLUTION/rdf.cpp

@@ -1,3 +1,4 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved.
 #include <stdio.h>
 #include <stdio.h>
 #include <iostream>
 #include <iostream>
 #include <fstream>
 #include <fstream>

+ 1 - 0
hpc/nways/nways_labs/nways_MD/English/C/source_code/openacc/SOLUTION/rdf_collapse.cpp

@@ -1,3 +1,4 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved.
 #include <stdio.h>
 #include <stdio.h>
 #include <iostream>
 #include <iostream>
 #include <fstream>
 #include <fstream>

+ 2 - 1
hpc/nways/nways_labs/nways_MD/English/C/source_code/openacc/SOLUTION/rdf_data_directive.cpp

@@ -1,3 +1,4 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved.
 #include <stdio.h>
 #include <stdio.h>
 #include <iostream>
 #include <iostream>
 #include <fstream>
 #include <fstream>
@@ -91,7 +92,7 @@ int main(int argc, char *argv[])
 		nvtxRangePush("Pair_Calculation");
 		nvtxRangePush("Pair_Calculation");
 		pair_gpu(h_x, h_y, h_z, h_g2, numatm, nconf, xbox, ybox, zbox, nbin);
 		pair_gpu(h_x, h_y, h_z, h_g2, numatm, nconf, xbox, ybox, zbox, nbin);
 		nvtxRangePop(); //Pop for Pair Calculation
 		nvtxRangePop(); //Pop for Pair Calculation
-		//////////////////////////////////////////////////////////////////////////
+						//////////////////////////////////////////////////////////////////////////
 	}
 	}
 
 
 	double pi = acos(-1.0l);
 	double pi = acos(-1.0l);

+ 1 - 0
hpc/nways/nways_labs/nways_MD/English/C/source_code/openacc/SOLUTION/rdf_gang.cpp

@@ -1,3 +1,4 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved. 
 #include <stdio.h>
 #include <stdio.h>
 #include <iostream>
 #include <iostream>
 #include <fstream>
 #include <fstream>

+ 1 - 0
hpc/nways/nways_labs/nways_MD/English/C/source_code/openacc/SOLUTION/rdf_gang_vector.cpp

@@ -1,3 +1,4 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved.
 #include <stdio.h>
 #include <stdio.h>
 #include <iostream>
 #include <iostream>
 #include <fstream>
 #include <fstream>

+ 1 - 0
hpc/nways/nways_labs/nways_MD/English/C/source_code/openacc/SOLUTION/rdf_gang_vector_length.cpp

@@ -1,3 +1,4 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved.
 #include <stdio.h>
 #include <stdio.h>
 #include <iostream>
 #include <iostream>
 #include <fstream>
 #include <fstream>

+ 1 - 0
hpc/nways/nways_labs/nways_MD/English/C/source_code/openacc/SOLUTION/rdf_gang_vector_worker.cpp

@@ -1,3 +1,4 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved.
 #include <stdio.h>
 #include <stdio.h>
 #include <iostream>
 #include <iostream>
 #include <fstream>
 #include <fstream>

+ 1 - 0
hpc/nways/nways_labs/nways_MD/English/C/source_code/openacc/SOLUTION/rdf_kernel_directive.cpp

@@ -1,3 +1,4 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved.
 #include <stdio.h>
 #include <stdio.h>
 #include <iostream>
 #include <iostream>
 #include <fstream>
 #include <fstream>

+ 122 - 112
hpc/nways/nways_labs/nways_MD/English/C/source_code/openacc/SOLUTION/rdf_parallel_directive.cpp

@@ -1,8 +1,9 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved.
 #include <stdio.h>
 #include <stdio.h>
 #include <iostream>
 #include <iostream>
 #include <fstream>
 #include <fstream>
-#include<stdlib.h>
-#include<accelmath.h>
+#include <stdlib.h>
+#include <accelmath.h>
 #include <cstring>
 #include <cstring>
 #include <cstdio>
 #include <cstdio>
 #include <iomanip>
 #include <iomanip>
@@ -11,182 +12,191 @@
 #include <omp.h>
 #include <omp.h>
 #include <nvToolsExt.h>
 #include <nvToolsExt.h>
 
 
-void pair_gpu(const double* d_x, const double* d_y, const double* d_z, 
-		unsigned long long int *d_g2, int numatm, int nconf, 
-		const double xbox, const double ybox, const double zbox,
-		int d_bin);
+void pair_gpu(const double *d_x, const double *d_y, const double *d_z,
+			  unsigned long long int *d_g2, int numatm, int nconf,
+			  const double xbox, const double ybox, const double zbox,
+			  int d_bin);
 
 
-int main(int argc , char* argv[] )
+int main(int argc, char *argv[])
 {
 {
-	double xbox,ybox,zbox;
-	double* h_x,*h_y,*h_z;
+	double xbox, ybox, zbox;
+	double *h_x, *h_y, *h_z;
 	unsigned long long int *h_g2;
 	unsigned long long int *h_g2;
 	int nbin;
 	int nbin;
-	int numatm,nconf,inconf;
+	int numatm, nconf, inconf;
 	string file;
 	string file;
-	// double start; 
-	//double end; 
-
+	// double start;
+	//double end;
 
 
 	///////////////////////////////////////////////////////////////
 	///////////////////////////////////////////////////////////////
 
 
 	inconf = 10;
 	inconf = 10;
-	nbin=2000;
+	nbin = 2000;
 	file = "../input/alk.traj.dcd";
 	file = "../input/alk.traj.dcd";
 	///////////////////////////////////////
 	///////////////////////////////////////
 	std::ifstream infile;
 	std::ifstream infile;
 	infile.open(file.c_str());
 	infile.open(file.c_str());
-	if(!infile){
-		cout<<"file "<<file.c_str()<<" not found\n";
+	if (!infile)
+	{
+		cout << "file " << file.c_str() << " not found\n";
 		return 1;
 		return 1;
 	}
 	}
 	assert(infile);
 	assert(infile);
 
 
-
-	ofstream pairfile,stwo;
+	ofstream pairfile, stwo;
 	pairfile.open("RDF.dat");
 	pairfile.open("RDF.dat");
 	stwo.open("Pair_entropy.dat");
 	stwo.open("Pair_entropy.dat");
 
 
 	/////////////////////////////////////////////////////////
 	/////////////////////////////////////////////////////////
-	dcdreadhead(&numatm,&nconf,infile);
-	cout<<"Dcd file has "<< numatm << " atoms and " << nconf << " frames"<<endl;
-	if (inconf>nconf) cout << "nconf is reset to "<< nconf <<endl;
+	dcdreadhead(&numatm, &nconf, infile);
+	cout << "Dcd file has " << numatm << " atoms and " << nconf << " frames" << endl;
+	if (inconf > nconf)
+		cout << "nconf is reset to " << nconf << endl;
 	else
 	else
-	{nconf=inconf;}
-	cout<<"Calculating RDF for " << nconf << " frames"<<endl;
+	{
+		nconf = inconf;
+	}
+	cout << "Calculating RDF for " << nconf << " frames" << endl;
 	////////////////////////////////////////////////////////
 	////////////////////////////////////////////////////////
 
 
-	unsigned long long int sizef= nconf*numatm*sizeof(double);
-	unsigned long long int sizebin= nbin*sizeof(unsigned long long int);
+	unsigned long long int sizef = nconf * numatm * sizeof(double);
+	unsigned long long int sizebin = nbin * sizeof(unsigned long long int);
 
 
-	h_x = (double*) malloc (sizef);
-	h_y = (double*) malloc (sizef); 
-	h_z = (double*) malloc (sizef);
-	h_g2 = (unsigned long long int*) malloc (sizebin);
+	h_x = (double *)malloc(sizef);
+	h_y = (double *)malloc(sizef);
+	h_z = (double *)malloc(sizef);
+	h_g2 = (unsigned long long int *)malloc(sizebin);
 
 
-	memset(h_g2,0,sizebin);
+	memset(h_g2, 0, sizebin);
 
 
 	/////////reading cordinates//////////////////////////////////////////////
 	/////////reading cordinates//////////////////////////////////////////////
 	nvtxRangePush("Read_File");
 	nvtxRangePush("Read_File");
-	double ax[numatm],ay[numatm],az[numatm];
-	for (int i=0;i<nconf;i++) {
-		dcdreadframe(ax,ay,az,infile,numatm,xbox,ybox,zbox);
-		for (int j=0;j<numatm;j++){
-			h_x[i*numatm+j]=ax[j];
-			h_y[i*numatm+j]=ay[j];
-			h_z[i*numatm+j]=az[j];
+	double ax[numatm], ay[numatm], az[numatm];
+	for (int i = 0; i < nconf; i++)
+	{
+		dcdreadframe(ax, ay, az, infile, numatm, xbox, ybox, zbox);
+		for (int j = 0; j < numatm; j++)
+		{
+			h_x[i * numatm + j] = ax[j];
+			h_y[i * numatm + j] = ay[j];
+			h_z[i * numatm + j] = az[j];
 		}
 		}
 	}
 	}
 	nvtxRangePop(); //pop for REading file
 	nvtxRangePop(); //pop for REading file
-	cout<<"Reading of input file is completed"<<endl;
+	cout << "Reading of input file is completed" << endl;
 	//////////////////////////////////////////////////////////////////////////
 	//////////////////////////////////////////////////////////////////////////
 	nvtxRangePush("Pair_Calculation");
 	nvtxRangePush("Pair_Calculation");
 	pair_gpu(h_x, h_y, h_z, h_g2, numatm, nconf, xbox, ybox, zbox, nbin);
 	pair_gpu(h_x, h_y, h_z, h_g2, numatm, nconf, xbox, ybox, zbox, nbin);
 	nvtxRangePop(); //Pop for Pair Calculation
 	nvtxRangePop(); //Pop for Pair Calculation
 	//////////////////////////////////////////////////////////////////////////
 	//////////////////////////////////////////////////////////////////////////
 
 
-	double pi=acos(-1.0l);
-	double rho=(numatm)/(xbox*ybox*zbox);
-	double norm=(4.0l*pi*rho)/3.0l;
-	double rl,ru,nideal;
+	double pi = acos(-1.0l);
+	double rho = (numatm) / (xbox * ybox * zbox);
+	double norm = (4.0l * pi * rho) / 3.0l;
+	double rl, ru, nideal;
 	double g2[nbin];
 	double g2[nbin];
-	double r,gr,lngr,lngrbond,s2=0.0l,s2bond=0.0l;
-	double box=min(xbox,ybox);
-	box=min(box,zbox);
-	double del=box/(2.0l*nbin);
+	double r, gr, lngr, lngrbond, s2 = 0.0l, s2bond = 0.0l;
+	double box = min(xbox, ybox);
+	box = min(box, zbox);
+	double del = box / (2.0l * nbin);
 	nvtxRangePush("Entropy_Calculation");
 	nvtxRangePush("Entropy_Calculation");
-	for (int i=0;i<nbin;i++) {
+	for (int i = 0; i < nbin; i++)
+	{
 		//      cout<<i+1<<" "<<h_g2[i]<<endl;
 		//      cout<<i+1<<" "<<h_g2[i]<<endl;
-		rl=(i)*del;
-		ru=rl+del;
-		nideal=norm*(ru*ru*ru-rl*rl*rl);
-		g2[i]=(double)h_g2[i]/((double)nconf*(double)numatm*nideal);
-		r=(i)*del;
-		pairfile<<(i+0.5l)*del<<" "<<g2[i]<<endl;
-		if (r<2.0l) {
-			gr=0.0l;
+		rl = (i)*del;
+		ru = rl + del;
+		nideal = norm * (ru * ru * ru - rl * rl * rl);
+		g2[i] = (double)h_g2[i] / ((double)nconf * (double)numatm * nideal);
+		r = (i)*del;
+		pairfile << (i + 0.5l) * del << " " << g2[i] << endl;
+		if (r < 2.0l)
+		{
+			gr = 0.0l;
 		}
 		}
-		else {
-			gr=g2[i];
+		else
+		{
+			gr = g2[i];
 		}
 		}
-		if (gr<1e-5) {
-			lngr=0.0l;
+		if (gr < 1e-5)
+		{
+			lngr = 0.0l;
 		}
 		}
-		else {
-			lngr=log(gr);
+		else
+		{
+			lngr = log(gr);
 		}
 		}
 
 
-		if (g2[i]<1e-6) {
-			lngrbond=0.0l;
+		if (g2[i] < 1e-6)
+		{
+			lngrbond = 0.0l;
 		}
 		}
-		else {
-			lngrbond=log(g2[i]);
+		else
+		{
+			lngrbond = log(g2[i]);
 		}
 		}
-		s2=s2-2.0l*pi*rho*((gr*lngr)-gr+1.0l)*del*r*r;
-		s2bond=s2bond-2.0l*pi*rho*((g2[i]*lngrbond)-g2[i]+1.0l)*del*r*r;
-
+		s2 = s2 - 2.0l * pi * rho * ((gr * lngr) - gr + 1.0l) * del * r * r;
+		s2bond = s2bond - 2.0l * pi * rho * ((g2[i] * lngrbond) - g2[i] + 1.0l) * del * r * r;
 	}
 	}
 	nvtxRangePop(); //Pop for Entropy Calculation
 	nvtxRangePop(); //Pop for Entropy Calculation
-	stwo<<"s2 value is "<<s2<<endl;
-	stwo<<"s2bond value is "<<s2bond<<endl;
-
-
-
-	cout<<"#Freeing Host memory"<<endl;
-	free ( h_x ) ;
-	free ( h_y ) ;
-	free ( h_z ) ;
-	free ( h_g2 ) ;
-
-	cout<<"#Number of atoms processed: "<<numatm<<endl<<endl;
-	cout<<"#Number of confs processed: "<<nconf<<endl<<endl;
+	stwo << "s2 value is " << s2 << endl;
+	stwo << "s2bond value is " << s2bond << endl;
+
+	cout << "#Freeing Host memory" << endl;
+	free(h_x);
+	free(h_y);
+	free(h_z);
+	free(h_g2);
+
+	cout << "#Number of atoms processed: " << numatm << endl
+		 << endl;
+	cout << "#Number of confs processed: " << nconf << endl
+		 << endl;
 	return 0;
 	return 0;
 }
 }
 #pragma acc routine seq
 #pragma acc routine seq
-int round(float num) 
-{ 
-	return num < 0 ? num - 0.5 : num + 0.5; 
-} 
-void pair_gpu( const double* d_x, const double* d_y, const double* d_z, 
-		unsigned long long int *d_g2, int numatm, int nconf, 
-		const double xbox,const double ybox,const double zbox,int d_bin) 
+int round(float num)
+{
+	return num < 0 ? num - 0.5 : num + 0.5;
+}
+void pair_gpu(const double *d_x, const double *d_y, const double *d_z,
+			  unsigned long long int *d_g2, int numatm, int nconf,
+			  const double xbox, const double ybox, const double zbox, int d_bin)
 {
 {
-	double r,cut,dx,dy,dz;
+	double r, cut, dx, dy, dz;
 	int ig2;
 	int ig2;
 	double box;
 	double box;
 	//int myround;
 	//int myround;
-	box=min(xbox,ybox);
-	box=min(box,zbox);
+	box = min(xbox, ybox);
+	box = min(box, zbox);
 
 
-	double del=box/(2.0*d_bin);
-	cut=box*0.5;
+	double del = box / (2.0 * d_bin);
+	cut = box * 0.5;
 
 
-	printf("\n %d %d ", nconf,numatm);
-	for (int frame=0;frame<nconf;frame++){
+	printf("\n %d %d ", nconf, numatm);
+	for (int frame = 0; frame < nconf; frame++)
+	{
 		printf("\n %d  ", frame);
 		printf("\n %d  ", frame);
 #pragma acc parallel loop
 #pragma acc parallel loop
-		for(int id1=0;id1<numatm;id1++)
+		for (int id1 = 0; id1 < numatm; id1++)
 		{
 		{
-			for(int id2=0;id2<numatm;id2++)
+			for (int id2 = 0; id2 < numatm; id2++)
 			{
 			{
-				dx=d_x[frame*numatm+id1]-d_x[frame*numatm+id2];
-				dy=d_y[frame*numatm+id1]-d_y[frame*numatm+id2];
-				dz=d_z[frame*numatm+id1]-d_z[frame*numatm+id2];
-
-				dx=dx-xbox*(round(dx/xbox));
-				dy=dy-ybox*(round(dy/ybox));
-				dz=dz-zbox*(round(dz/zbox));
-
-				r=sqrtf(dx*dx+dy*dy+dz*dz);
-				if (r<cut) {
-					ig2=(int)(r/del);
+				dx = d_x[frame * numatm + id1] - d_x[frame * numatm + id2];
+				dy = d_y[frame * numatm + id1] - d_y[frame * numatm + id2];
+				dz = d_z[frame * numatm + id1] - d_z[frame * numatm + id2];
+
+				dx = dx - xbox * (round(dx / xbox));
+				dy = dy - ybox * (round(dy / ybox));
+				dz = dz - zbox * (round(dz / zbox));
+
+				r = sqrtf(dx * dx + dy * dy + dz * dz);
+				if (r < cut)
+				{
+					ig2 = (int)(r / del);
 #pragma acc atomic
 #pragma acc atomic
-					d_g2[ig2] = d_g2[ig2] +1 ;
+					d_g2[ig2] = d_g2[ig2] + 1;
 				}
 				}
 			}
 			}
 		}
 		}
 	}
 	}
 }
 }
-
-
-

+ 40 - 34
hpc/nways/nways_labs/nways_MD/English/C/source_code/openacc/dcdread.h

@@ -1,43 +1,49 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved.
 using namespace std;
 using namespace std;
 
 
-void dcdreadhead(int *natom, int *nframes, std::istream &infile) {
+void dcdreadhead(int *natom, int *nframes, std::istream &infile)
+{
 
 
- infile.seekg(8,ios::beg);
- infile.read((char*)nframes, sizeof(int));
- infile.seekg(64*4,ios::cur);
- infile.read((char*)natom, sizeof(int));
- infile.seekg(1*8,ios::cur);
- return;
+    infile.seekg(8, ios::beg);
+    infile.read((char *)nframes, sizeof(int));
+    infile.seekg(64 * 4, ios::cur);
+    infile.read((char *)natom, sizeof(int));
+    infile.seekg(1 * 8, ios::cur);
+    return;
 }
 }
 
 
 void dcdreadframe(double *x, double *y, double *z, std::istream &infile,
 void dcdreadframe(double *x, double *y, double *z, std::istream &infile,
-    int natom, double& xbox,double& ybox,double& zbox){
+                  int natom, double &xbox, double &ybox, double &zbox)
+{
 
 
-  double d[6];
-  for (int i=0;i<6;i++) {
-      infile.read((char*)&d[i], sizeof(double));
-  }
-  xbox=d[0];
-  ybox=d[2];
-  zbox=d[5];
-  float a,b,c;
-  infile.seekg(1*8,ios::cur);
-  for (int i=0;i<natom;i++){
-      infile.read((char*)&a, sizeof(float));
-      x[i]=a;
-  }
-  infile.seekg(1*8,ios::cur);
-  for (int i=0;i<natom;i++) {
-      infile.read((char*)&b, sizeof(float));
-      y[i]=b;
-  }
-  infile.seekg(1*8,ios::cur);
-  for (int i=0;i<natom;i++) {
-      infile.read((char*)&c, sizeof(float));
-      z[i]=c;
-  }
-  infile.seekg(1*8,ios::cur);
+    double d[6];
+    for (int i = 0; i < 6; i++)
+    {
+        infile.read((char *)&d[i], sizeof(double));
+    }
+    xbox = d[0];
+    ybox = d[2];
+    zbox = d[5];
+    float a, b, c;
+    infile.seekg(1 * 8, ios::cur);
+    for (int i = 0; i < natom; i++)
+    {
+        infile.read((char *)&a, sizeof(float));
+        x[i] = a;
+    }
+    infile.seekg(1 * 8, ios::cur);
+    for (int i = 0; i < natom; i++)
+    {
+        infile.read((char *)&b, sizeof(float));
+        y[i] = b;
+    }
+    infile.seekg(1 * 8, ios::cur);
+    for (int i = 0; i < natom; i++)
+    {
+        infile.read((char *)&c, sizeof(float));
+        z[i] = c;
+    }
+    infile.seekg(1 * 8, ios::cur);
 
 
-  return;
+    return;
 }
 }
-           

+ 1 - 0
hpc/nways/nways_labs/nways_MD/English/C/source_code/openacc/rdf.cpp

@@ -1,3 +1,4 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved.
 #include <stdio.h>
 #include <stdio.h>
 #include <iostream>
 #include <iostream>
 #include <fstream>
 #include <fstream>

+ 40 - 34
hpc/nways/nways_labs/nways_MD/English/C/source_code/openmp/SOLUTION/dcdread.h

@@ -1,43 +1,49 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved.
 using namespace std;
 using namespace std;
 
 
-void dcdreadhead(int *natom, int *nframes, std::istream &infile) {
+void dcdreadhead(int *natom, int *nframes, std::istream &infile)
+{
 
 
- infile.seekg(8,ios::beg);
- infile.read((char*)nframes, sizeof(int));
- infile.seekg(64*4,ios::cur);
- infile.read((char*)natom, sizeof(int));
- infile.seekg(1*8,ios::cur);
- return;
+    infile.seekg(8, ios::beg);
+    infile.read((char *)nframes, sizeof(int));
+    infile.seekg(64 * 4, ios::cur);
+    infile.read((char *)natom, sizeof(int));
+    infile.seekg(1 * 8, ios::cur);
+    return;
 }
 }
 
 
 void dcdreadframe(double *x, double *y, double *z, std::istream &infile,
 void dcdreadframe(double *x, double *y, double *z, std::istream &infile,
-    int natom, double& xbox,double& ybox,double& zbox){
+                  int natom, double &xbox, double &ybox, double &zbox)
+{
 
 
-  double d[6];
-  for (int i=0;i<6;i++) {
-      infile.read((char*)&d[i], sizeof(double));
-  }
-  xbox=d[0];
-  ybox=d[2];
-  zbox=d[5];
-  float a,b,c;
-  infile.seekg(1*8,ios::cur);
-  for (int i=0;i<natom;i++){
-      infile.read((char*)&a, sizeof(float));
-      x[i]=a;
-  }
-  infile.seekg(1*8,ios::cur);
-  for (int i=0;i<natom;i++) {
-      infile.read((char*)&b, sizeof(float));
-      y[i]=b;
-  }
-  infile.seekg(1*8,ios::cur);
-  for (int i=0;i<natom;i++) {
-      infile.read((char*)&c, sizeof(float));
-      z[i]=c;
-  }
-  infile.seekg(1*8,ios::cur);
+    double d[6];
+    for (int i = 0; i < 6; i++)
+    {
+        infile.read((char *)&d[i], sizeof(double));
+    }
+    xbox = d[0];
+    ybox = d[2];
+    zbox = d[5];
+    float a, b, c;
+    infile.seekg(1 * 8, ios::cur);
+    for (int i = 0; i < natom; i++)
+    {
+        infile.read((char *)&a, sizeof(float));
+        x[i] = a;
+    }
+    infile.seekg(1 * 8, ios::cur);
+    for (int i = 0; i < natom; i++)
+    {
+        infile.read((char *)&b, sizeof(float));
+        y[i] = b;
+    }
+    infile.seekg(1 * 8, ios::cur);
+    for (int i = 0; i < natom; i++)
+    {
+        infile.read((char *)&c, sizeof(float));
+        z[i] = c;
+    }
+    infile.seekg(1 * 8, ios::cur);
 
 
-  return;
+    return;
 }
 }
-           

+ 1 - 0
hpc/nways/nways_labs/nways_MD/English/C/source_code/openmp/SOLUTION/rdf.cpp

@@ -1,3 +1,4 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved.
 #include <stdio.h>
 #include <stdio.h>
 #include <iostream>
 #include <iostream>
 #include <fstream>
 #include <fstream>

+ 1 - 0
hpc/nways/nways_labs/nways_MD/English/C/source_code/openmp/SOLUTION/rdf_offload.cpp

@@ -1,3 +1,4 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved.
 #include <stdio.h>
 #include <stdio.h>
 #include <iostream>
 #include <iostream>
 #include <fstream>
 #include <fstream>

+ 1 - 0
hpc/nways/nways_labs/nways_MD/English/C/source_code/openmp/SOLUTION/rdf_offload_collapse.cpp

@@ -1,3 +1,4 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved.
 #include <stdio.h>
 #include <stdio.h>
 #include <iostream>
 #include <iostream>
 #include <fstream>
 #include <fstream>

+ 1 - 0
hpc/nways/nways_labs/nways_MD/English/C/source_code/openmp/dcdread.h

@@ -1,3 +1,4 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved. 
 using namespace std;
 using namespace std;
 
 
 void dcdreadhead(int *natom, int *nframes, std::istream &infile) {
 void dcdreadhead(int *natom, int *nframes, std::istream &infile) {

+ 1 - 0
hpc/nways/nways_labs/nways_MD/English/C/source_code/openmp/rdf.cpp

@@ -1,3 +1,4 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved.
 #include <stdio.h>
 #include <stdio.h>
 #include <iostream>
 #include <iostream>
 #include <fstream>
 #include <fstream>

+ 40 - 34
hpc/nways/nways_labs/nways_MD/English/C/source_code/serial/dcdread.h

@@ -1,43 +1,49 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved.
 using namespace std;
 using namespace std;
 
 
-void dcdreadhead(int *natom, int *nframes, std::istream &infile) {
+void dcdreadhead(int *natom, int *nframes, std::istream &infile)
+{
 
 
- infile.seekg(8,ios::beg);
- infile.read((char*)nframes, sizeof(int));
- infile.seekg(64*4,ios::cur);
- infile.read((char*)natom, sizeof(int));
- infile.seekg(1*8,ios::cur);
- return;
+    infile.seekg(8, ios::beg);
+    infile.read((char *)nframes, sizeof(int));
+    infile.seekg(64 * 4, ios::cur);
+    infile.read((char *)natom, sizeof(int));
+    infile.seekg(1 * 8, ios::cur);
+    return;
 }
 }
 
 
 void dcdreadframe(double *x, double *y, double *z, std::istream &infile,
 void dcdreadframe(double *x, double *y, double *z, std::istream &infile,
-    int natom, double& xbox,double& ybox,double& zbox){
+                  int natom, double &xbox, double &ybox, double &zbox)
+{
 
 
-  double d[6];
-  for (int i=0;i<6;i++) {
-      infile.read((char*)&d[i], sizeof(double));
-  }
-  xbox=d[0];
-  ybox=d[2];
-  zbox=d[5];
-  float a,b,c;
-  infile.seekg(1*8,ios::cur);
-  for (int i=0;i<natom;i++){
-      infile.read((char*)&a, sizeof(float));
-      x[i]=a;
-  }
-  infile.seekg(1*8,ios::cur);
-  for (int i=0;i<natom;i++) {
-      infile.read((char*)&b, sizeof(float));
-      y[i]=b;
-  }
-  infile.seekg(1*8,ios::cur);
-  for (int i=0;i<natom;i++) {
-      infile.read((char*)&c, sizeof(float));
-      z[i]=c;
-  }
-  infile.seekg(1*8,ios::cur);
+    double d[6];
+    for (int i = 0; i < 6; i++)
+    {
+        infile.read((char *)&d[i], sizeof(double));
+    }
+    xbox = d[0];
+    ybox = d[2];
+    zbox = d[5];
+    float a, b, c;
+    infile.seekg(1 * 8, ios::cur);
+    for (int i = 0; i < natom; i++)
+    {
+        infile.read((char *)&a, sizeof(float));
+        x[i] = a;
+    }
+    infile.seekg(1 * 8, ios::cur);
+    for (int i = 0; i < natom; i++)
+    {
+        infile.read((char *)&b, sizeof(float));
+        y[i] = b;
+    }
+    infile.seekg(1 * 8, ios::cur);
+    for (int i = 0; i < natom; i++)
+    {
+        infile.read((char *)&c, sizeof(float));
+        z[i] = c;
+    }
+    infile.seekg(1 * 8, ios::cur);
 
 
-  return;
+    return;
 }
 }
-           

+ 1 - 0
hpc/nways/nways_labs/nways_MD/English/C/source_code/serial/rdf.cpp

@@ -1,3 +1,4 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved.
 #include <stdio.h>
 #include <stdio.h>
 #include <stdlib.h>
 #include <stdlib.h>
 #include <iostream>
 #include <iostream>

+ 40 - 34
hpc/nways/nways_labs/nways_MD/English/C/source_code/stdpar/SOLUTION/dcdread.h

@@ -1,43 +1,49 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved.
 using namespace std;
 using namespace std;
 
 
-void dcdreadhead(int *natom, int *nframes, std::istream &infile) {
+void dcdreadhead(int *natom, int *nframes, std::istream &infile)
+{
 
 
- infile.seekg(8,ios::beg);
- infile.read((char*)nframes, sizeof(int));
- infile.seekg(64*4,ios::cur);
- infile.read((char*)natom, sizeof(int));
- infile.seekg(1*8,ios::cur);
- return;
+    infile.seekg(8, ios::beg);
+    infile.read((char *)nframes, sizeof(int));
+    infile.seekg(64 * 4, ios::cur);
+    infile.read((char *)natom, sizeof(int));
+    infile.seekg(1 * 8, ios::cur);
+    return;
 }
 }
 
 
 void dcdreadframe(double *x, double *y, double *z, std::istream &infile,
 void dcdreadframe(double *x, double *y, double *z, std::istream &infile,
-    int natom, double& xbox,double& ybox,double& zbox){
+                  int natom, double &xbox, double &ybox, double &zbox)
+{
 
 
-  double d[6];
-  for (int i=0;i<6;i++) {
-      infile.read((char*)&d[i], sizeof(double));
-  }
-  xbox=d[0];
-  ybox=d[2];
-  zbox=d[5];
-  float a,b,c;
-  infile.seekg(1*8,ios::cur);
-  for (int i=0;i<natom;i++){
-      infile.read((char*)&a, sizeof(float));
-      x[i]=a;
-  }
-  infile.seekg(1*8,ios::cur);
-  for (int i=0;i<natom;i++) {
-      infile.read((char*)&b, sizeof(float));
-      y[i]=b;
-  }
-  infile.seekg(1*8,ios::cur);
-  for (int i=0;i<natom;i++) {
-      infile.read((char*)&c, sizeof(float));
-      z[i]=c;
-  }
-  infile.seekg(1*8,ios::cur);
+    double d[6];
+    for (int i = 0; i < 6; i++)
+    {
+        infile.read((char *)&d[i], sizeof(double));
+    }
+    xbox = d[0];
+    ybox = d[2];
+    zbox = d[5];
+    float a, b, c;
+    infile.seekg(1 * 8, ios::cur);
+    for (int i = 0; i < natom; i++)
+    {
+        infile.read((char *)&a, sizeof(float));
+        x[i] = a;
+    }
+    infile.seekg(1 * 8, ios::cur);
+    for (int i = 0; i < natom; i++)
+    {
+        infile.read((char *)&b, sizeof(float));
+        y[i] = b;
+    }
+    infile.seekg(1 * 8, ios::cur);
+    for (int i = 0; i < natom; i++)
+    {
+        infile.read((char *)&c, sizeof(float));
+        z[i] = c;
+    }
+    infile.seekg(1 * 8, ios::cur);
 
 
-  return;
+    return;
 }
 }
-           

+ 1 - 0
hpc/nways/nways_labs/nways_MD/English/C/source_code/stdpar/SOLUTION/rdf.cpp

@@ -1,3 +1,4 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved.
 #include <stdio.h>
 #include <stdio.h>
 #include <iostream>
 #include <iostream>
 #include <fstream>
 #include <fstream>

+ 1 - 0
hpc/nways/nways_labs/nways_MD/English/C/source_code/stdpar/dcdread.h

@@ -1,3 +1,4 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved. 
 using namespace std;
 using namespace std;
 
 
 void dcdreadhead(int *natom, int *nframes, std::istream &infile) {
 void dcdreadhead(int *natom, int *nframes, std::istream &infile) {

+ 1 - 0
hpc/nways/nways_labs/nways_MD/English/C/source_code/stdpar/rdf.cpp

@@ -1,3 +1,4 @@
+// Copyright (c) 2021 NVIDIA Corporation.  All rights reserved.
 #include <stdio.h>
 #include <stdio.h>
 #include <iostream>
 #include <iostream>
 #include <fstream>
 #include <fstream>

+ 1 - 1
hpc/nways/nways_labs/nways_start.ipynb

@@ -31,7 +31,7 @@
    "source": [
    "source": [
     "### Tutorial Outline\n",
     "### Tutorial Outline\n",
     "\n",
     "\n",
-    "During this lab, we will be working on porting mini applications in MD and/or CFD domains to GPUs. You can choose to work with either of this application. Please click on one of the below links to start:\n",
+    "During this lab, we will be working on porting mini applications in MD domain to GPUs. You can choose to work with either of this application. Please click on one of the below links to start:\n",
     "\n",
     "\n",
     "- N Ways to GPU Programming in [MD](nways_MD/English/nways_MD_start.ipynb) domain\n"
     "- N Ways to GPU Programming in [MD](nways_MD/English/nways_MD_start.ipynb) domain\n"
    ]
    ]