From 88557f41fc4eabdcb622fdb5530e94c68025478e Mon Sep 17 00:00:00 2001 From: avillarr Date: Mon, 13 Jul 2020 19:43:54 -0700 Subject: [PATCH 01/17] Initial Commit to Open Source Repo Signed-off-by: avillarr --- .../iso2dfd_dpcpp/CMakeLists.txt | 27 ++ .../StructuredGrids/iso2dfd_dpcpp/License.txt | 7 + .../StructuredGrids/iso2dfd_dpcpp/README.md | 140 +++++++ .../StructuredGrids/iso2dfd_dpcpp/iso2dfd.sln | 25 ++ .../iso2dfd_dpcpp/iso2dfd.vcxproj | 147 +++++++ .../iso2dfd_dpcpp/iso2dfd.vcxproj.filters | 22 + .../iso2dfd_dpcpp/iso2dfd.vcxproj.user | 11 + .../StructuredGrids/iso2dfd_dpcpp/sample.json | 30 ++ .../iso2dfd_dpcpp/src/iso2dfd.cpp | 380 ++++++++++++++++++ 9 files changed, 789 insertions(+) create mode 100644 DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/CMakeLists.txt create mode 100644 DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/License.txt create mode 100644 DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md create mode 100644 DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/iso2dfd.sln create mode 100644 DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/iso2dfd.vcxproj create mode 100644 DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/iso2dfd.vcxproj.filters create mode 100644 DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/iso2dfd.vcxproj.user create mode 100644 DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/sample.json create mode 100644 DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/src/iso2dfd.cpp diff --git a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/CMakeLists.txt b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/CMakeLists.txt new file mode 100644 index 0000000000..9dd05e922e --- /dev/null +++ b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/CMakeLists.txt @@ -0,0 +1,27 @@ +# required cmake version +cmake_minimum_required(VERSION 3.5) + +# CMakeLists.txt for ISO2DFD_DPCPP project +project (iso2dfd_dpcpp) + +set(CMAKE_CXX_COMPILER "dpcpp") + +# Set default build type to RelWithDebInfo if not specified +if (NOT CMAKE_BUILD_TYPE) + message (STATUS "Default CMAKE_BUILD_TYPE not set using Release with Debug Info") + set (CMAKE_BUILD_TYPE "RelWithDebInfo" CACHE + STRING "Choose the type of build, options are: None Debug Release RelWithDebInfo MinSizeRel" + FORCE) +endif() + +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -O3 -fsycl -std=c++17") + +set(CMAKE_EXE_LINKER_FLAGS "${CMAKE_EXE_LINKER_FLAGS} -lOpenCL -lsycl") + +add_executable (iso2dfd src/iso2dfd.cpp) + +add_custom_target (run + COMMAND iso2dfd 1000 1000 2000 + WORKING_DIRECTORY ${CMAKE_PROJECT_DIR} +) + diff --git a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/License.txt b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/License.txt new file mode 100644 index 0000000000..0578223382 --- /dev/null +++ b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/License.txt @@ -0,0 +1,7 @@ +Copyright 2019 Intel Corporation + +Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights to use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies of the Software, and to permit persons to whom the Software is furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. diff --git a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md new file mode 100644 index 0000000000..427847d11a --- /dev/null +++ b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md @@ -0,0 +1,140 @@ +# ISO2DFD sample + +ISO2DFD: Intel® oneAPI DPC++ Language Basics Using +2D-Finite-Difference-Wave Propagation + +The ISO2DFD sample refers to Two-Dimensional Finite-Difference Wave Propagation in Isotropic Media. It is a two-dimensional stencil to simulate a wave propagating in a 2D isotropic medium and illustrates the basics of the DPC++ programming language using direct programming. + +A complete code walk-through for this sample can be found at: +https://software.intel.com/en-us/articles/code-sample-two-dimensional-finite-difference-wave-propagation-in-isotropic-media-iso2dfd + +For comprehensive instructions regarding DPC++ Programming, go to +https://software.intel.com/en-us/oneapi-programming-guide +and search based on relevant terms noted in the comments. + + +| Optimized for | Description +|:--- |:--- +| OS | Linux Ubuntu 18.04 +| Hardware | Skylake with GEN9 or newer +| Software | Intel® oneAPI DPC++ Compiler (beta); Intel C++ Compiler (beta) +| What you will learn | How to offload the computation to GPU using Intel DPC++ compiler +| Time to complete | 10 minutes + + +## Purpose + +ISO2DFD is a finite difference stencil kernel for solving the 2D acoustic isotropic wave equation. In +this sample, we chose the problem of solving a Partial Differential Equation (PDE), using a +finite-difference method, to illustrate the essential elements of the DPC++ programming language: +queues, buffers/accessors, and kernels. Use it as an entry point to start programming in DPC++ or as a +proxy to develop or better understand complicated code for similar problems. + +Using Data Parallel C++, the sample will explicitly run on the GPU as well as CPU to calculate a +result. The output will include GPU device name. The results from the two devices are compared and, if +the sample ran correctly, report a success message. The output of the wavefield can be plotted using +the SU Seismic processing library, which has utilities to display seismic wavefields and can be +downloaded from John Stockwell’s SeisUnix GitHub* (https://github.com/JohnWStockwellJr/SeisUnix/wiki/ +Seismic-Unix-install-on-Ubuntu) + + +## Key implementation details + +SYCL implementation explained. + +* DPC++ queues (including device selectors and exception handlers). +* DPC++ buffers and accessors. +* The ability to call a function inside a kernel definition and pass accessor arguments as pointers. A +function called inside the kernel performs a computation (it updates a grid point specified by the +global ID variable) for a single time step. + + +## License + +This code sample is licensed under MIT license. + + +## Building the `iso2dfd` Program for CPU and GPU + +### Running Samples In DevCloud + +If running a sample in the Intel DevCloud, remember that you must specify the compute node (CPU, GPU, +FPGA) as well whether to run in batch or interactive mode. For more information see the Intel® oneAPI +Base Toolkit Get Started Guide (https://devcloud.intel.com/oneapi/get-started/base-toolkit/) + +### On a Linux* System +Perform the following steps: +1. Build the program using the following `cmake` commands. + + ``` + cd iso2dfd_dpcpp && + mkdir build && + cd build && + cmake .. && + make -j + ``` + +2. Run the program on Gen9 + + ``` + make run + ``` + +3. Clean the program + + ``` + make clean + ``` + + +### On a Windows* System Using Visual Studio* Version 2017 or Newer +* Build the program using VS2017 or VS2019 + Right click on the solution file and open using either VS2017 or VS2019 IDE. + Right click on the project in Solution explorer and select Rebuild. + From top menu select Debug -> Start without Debugging. + +* Build the program using MSBuild + Open "x64 Native Tools Command Prompt for VS2017" or "x64 Native Tools Command Prompt for VS2019" + Run - MSBuild iso2dfd.sln /t:Rebuild /p:Configuration="Release" + + +## Running the Sample + +### Application Parameters + +You can execute the code with different parameters. For example the following command will run the iso2dfd executable using a 1000x1000 grid size and it will iterate over 2000 time steps. + + ``` + ./iso2dfd 1000 1000 2000 + ``` + + Usage: ./iso2dfd n1 n2 Iterations + + n1 n2 : Grid sizes for the stencil + Iterations : Number of timesteps. + + * Find graphical output for sample execution in the online tutorial at: + https://software.intel.com/en-us/articles/code-sample-two-dimensional-finite-difference-wave-propagation-in-isotropic-media-iso2dfd + +### Example of Output + + ``` + Initializing ... + Grid Sizes: 1000 1000 + Iterations: 2000 + + Computing wavefield in device .. + Running on Intel(R) Gen9 HD Graphics NEO + The Device Max Work Group Size is : 256 + The Device Max EUCount is : 24 + SYCL time: 3282 ms + + Computing wavefield in CPU .. + Initializing ... + CPU time: 8846 ms + + Final wavefields from device and CPU are equivalent: Success + Final wavefields (from device and CPU) written to disk + Finished. + [100%] Built target run + ``` diff --git a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/iso2dfd.sln b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/iso2dfd.sln new file mode 100644 index 0000000000..174faa6896 --- /dev/null +++ b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/iso2dfd.sln @@ -0,0 +1,25 @@ + +Microsoft Visual Studio Solution File, Format Version 12.00 +# Visual Studio 15 +VisualStudioVersion = 15.0.28307.960 +MinimumVisualStudioVersion = 10.0.40219.1 +Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "iso2dfd", "iso2dfd.vcxproj", "{1AE3DD06-C3F0-4746-B126-EEB6A94CF35C}" +EndProject +Global + GlobalSection(SolutionConfigurationPlatforms) = preSolution + Debug|x64 = Debug|x64 + Release|x64 = Release|x64 + EndGlobalSection + GlobalSection(ProjectConfigurationPlatforms) = postSolution + {1AE3DD06-C3F0-4746-B126-EEB6A94CF35C}.Debug|x64.ActiveCfg = Debug|x64 + {1AE3DD06-C3F0-4746-B126-EEB6A94CF35C}.Debug|x64.Build.0 = Debug|x64 + {1AE3DD06-C3F0-4746-B126-EEB6A94CF35C}.Release|x64.ActiveCfg = Release|x64 + {1AE3DD06-C3F0-4746-B126-EEB6A94CF35C}.Release|x64.Build.0 = Release|x64 + EndGlobalSection + GlobalSection(SolutionProperties) = preSolution + HideSolutionNode = FALSE + EndGlobalSection + GlobalSection(ExtensibilityGlobals) = postSolution + SolutionGuid = {AC4985B6-FFDE-4420-B533-7D4318863288} + EndGlobalSection +EndGlobal diff --git a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/iso2dfd.vcxproj b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/iso2dfd.vcxproj new file mode 100644 index 0000000000..c7e2a1d29f --- /dev/null +++ b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/iso2dfd.vcxproj @@ -0,0 +1,147 @@ + + + + + Debug + x64 + + + Release + x64 + + + + + + + 15.0 + {1ae3dd06-c3f0-4746-b126-eeb6a94cf35c} + Win32Proj + iso2dfd + $(WindowsSDKVersion.Replace("\","")) + + + + Application + true + Intel(R) oneAPI DPC++ Compiler + Unicode + + + Application + false + Intel(R) oneAPI DPC++ Compiler + true + Unicode + + + Application + true + Intel(R) oneAPI DPC++ Compiler + Unicode + + + Application + false + Intel(R) oneAPI DPC++ Compiler + true + Unicode + + + + + + + + + + + + + + + + + + + + + true + + + true + + + false + + + false + + + + Use + Level3 + Disabled + true + true + pch.h + + + Console + true + + + + + Use + Level3 + Disabled + true + true + pch.h + + + Console + true + + + + + Use + Level3 + MaxSpeed + true + true + true + true + pch.h + + + Console + true + true + true + + + + + Use + Level3 + MaxSpeed + true + true + true + true + pch.h + + + Console + true + true + true + + + + + + \ No newline at end of file diff --git a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/iso2dfd.vcxproj.filters b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/iso2dfd.vcxproj.filters new file mode 100644 index 0000000000..1b7c40576f --- /dev/null +++ b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/iso2dfd.vcxproj.filters @@ -0,0 +1,22 @@ + + + + + {4FC737F1-C7A5-4376-A066-2A32D752A2FF} + cpp;c;cc;cxx;def;odl;idl;hpj;bat;asm;asmx + + + {93995380-89BD-4b04-88EB-625FBE52EBFB} + h;hh;hpp;hxx;hm;inl;inc;ipp;xsd + + + {67DA6AB6-F800-4c08-8B7A-83BB121AAD01} + rc;ico;cur;bmp;dlg;rc2;rct;bin;rgs;gif;jpg;jpeg;jpe;resx;tiff;tif;png;wav;mfcribbon-ms + + + + + Source Files + + + \ No newline at end of file diff --git a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/iso2dfd.vcxproj.user b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/iso2dfd.vcxproj.user new file mode 100644 index 0000000000..f5d6e260eb --- /dev/null +++ b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/iso2dfd.vcxproj.user @@ -0,0 +1,11 @@ + + + + 1000 1000 2000 + WindowsLocalDebugger + + + 1000 1000 2000 + WindowsLocalDebugger + + \ No newline at end of file diff --git a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/sample.json b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/sample.json new file mode 100644 index 0000000000..f97a3bd596 --- /dev/null +++ b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/sample.json @@ -0,0 +1,30 @@ +{ + "guid": "9483C0F0-7D63-4E99-86C5-37C40F77B2AE" , + "name": "iso2dfd_dpcpp", + "categories": [ "Toolkit/Intel® oneAPI HPC Toolkit" ], + "description": "ISO2DFD: Intel® oneAPI DPC++ Language Basics Using 2D Finite-Difference-Wave Propagation", + "toolchain": [ "dpcpp" ], + "targetDevice": [ "CPU", "GPU" ], + "languages": [ { "cpp": {} } ], + "os": [ "linux", "windows"], + "builder": [ "ide", "cmake" ], + "ciTests": { + "linux": [{ + "steps": [ + "mkdir build", + "cd build", + "cmake ..", + "make", + "make run" + ] + }], + "windows": [{ + "steps": [ + "MSBuild iso2dfd.sln /t:Rebuild /p:Configuration=\"Release\"", + "cd x64/Release", + "iso2dfd.exe 1000 1000 2000" + ] + }] + + } +} diff --git a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/src/iso2dfd.cpp b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/src/iso2dfd.cpp new file mode 100644 index 0000000000..e4638ba703 --- /dev/null +++ b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/src/iso2dfd.cpp @@ -0,0 +1,380 @@ +//============================================================== +// Copyright © 2019 Intel Corporation +// +// SPDX-License-Identifier: MIT +// ============================================================= + +// ISO2DFD: Intel® oneAPI DPC++ Language Basics Using 2D-Finite-Difference-Wave +// Propagation +// +// ISO2DFD is a finite difference stencil kernel for solving the 2D acoustic +// isotropic wave equation. Kernels in this sample are implemented as 2nd order +// in space, 2nd order in time scheme without boundary conditions. Using Data +// Parallel C++, the sample will explicitly run on the GPU as well as CPU to +// calculate a result. If successful, the output will include GPU device name. +// +// A complete online tutorial for this code sample can be found at : +// https://software.intel.com/en-us/articles/code-sample-two-dimensional-finite-difference-wave-propagation-in-isotropic-media-iso2dfd +// +// For comprehensive instructions regarding DPC++ Programming, go to +// https://software.intel.com/en-us/oneapi-programming-guide +// and search based on relevant terms noted in the comments. +// +// DPC++ material used in this code sample: +// +// Basic structures of DPC++: +// DPC++ Queues (including device selectors and exception handlers) +// DPC++ Buffers and accessors (communicate data between the host and the device) +// DPC++ Kernels (including parallel_for function and range<2> objects) +// + +#include +#include +#include +#include +#include +#include +#include + +#include "dpc_common.hpp" + +using namespace cl::sycl; +using namespace std; + +/* + * Parameters to define coefficients + * half_length: Radius of the stencil + * Sample source code is tested for half_length=1 resulting in + * 2nd order Stencil finite difference kernel + */ + +constexpr float DT = 0.002f; +constexpr float DXY = 20.0f; +constexpr unsigned int half_length = 1; + +/* + * Host-Code + * Utility function to display input arguments + */ +void Usage(const string &program_name) { + cout << " Incorrect parameters\n"; + cout << " Usage: "; + cout << program_name << " n1 n2 Iterations\n\n"; + cout << " n1 n2 : Grid sizes for the stencil\n"; + cout << " Iterations : No. of timesteps.\n"; +} + +/* + * Host-Code + * Function used for initialization + */ +void Initialize(float* ptr_prev, float* ptr_next, float* ptr_vel, size_t n_rows, + size_t n_cols) { + cout << "Initializing ...\n"; + + // Define source wavelet + float wavelet[12] = {0.016387336, -0.041464937, -0.067372555, 0.386110067, + 0.812723635, 0.416998396, 0.076488599, -0.059434419, + 0.023680172, 0.005611435, 0.001823209, -0.000720549}; + + // Initialize arrays + for (size_t i = 0; i < n_rows; i++) { + size_t offset = i * n_cols; + + for (int k = 0; k < n_cols; k++) { + ptr_prev[offset + k] = 0.0f; + ptr_next[offset + k] = 0.0f; + // pre-compute squared value of sample wave velocity v*v (v = 1500 m/s) + ptr_vel[offset + k] = (1500.0f * 1500.0f); + } + } + // Add a source to initial wavefield as an initial condition + for (int s = 11; s >= 0; s--) { + for (int i = n_rows / 2 - s; i < n_rows / 2 + s; i++) { + size_t offset = i * n_cols; + for (int k = n_cols / 2 - s; k < n_cols / 2 + s; k++) { + ptr_prev[offset + k] = wavelet[s]; + } + } + } +} + +/* + * Host-Code + * Utility function to print device info + */ +void PrintTargetInfo(queue& q) { + auto device = q.get_device(); + auto max_block_size = + device.get_info(); + + auto max_EU_count = + device.get_info(); + + cout<< " Running on " << device.get_info()<<"\n"; + cout<< " The Device Max Work Group Size is : "<< max_block_size<<"\n"; + cout<< " The Device Max EUCount is : " << max_EU_count<<"\n"; +} + +/* + * Host-Code + * Utility function to calculate L2-norm between resulting buffer and reference + * buffer + */ +bool WithinEpsilon(float* output, float* reference, const size_t dim_x, + const size_t dim_y, const unsigned int radius, + const float delta = 0.01f) { + ofstream err_file; + err_file.open("error_diff.txt"); + + bool error = false; + double norm2 = 0; + + for (size_t iy = 0; iy < dim_y; iy++) { + for (size_t ix = 0; ix < dim_x; ix++) { + if (ix >= radius && ix < (dim_x - radius) && iy >= radius && + iy < (dim_y - radius)) { + float difference = fabsf(*reference - *output); + norm2 += difference * difference; + if (difference > delta) { + error = true; + err_file<<" ERROR: "< it, float* next, float* prev, + float* vel, const float dtDIVdxy, int n_rows, + int n_cols) { + float value = 0.0; + + // Compute global id + // We can use the get.global.id() function of the item variable + // to compute global id. The 2D array is laid out in memory in row major + // order. + size_t gid_row = it.get(0); + size_t gid_col = it.get(1); + size_t gid = (gid_row)*n_cols + gid_col; + + // Computation to solve wave equation in 2D + // First check if gid is inside the effective grid (not in halo) + if ((gid_col >= half_length && gid_col < n_cols - half_length) && + (gid_row >= half_length && gid_row < n_rows - half_length)) { + // Stencil code to update grid point at position given by global id (gid) + // New time step for grid point is computed based on the values of the + // the immediate neighbors in both the horizontal and vertical + // directions, as well as the value of grid point at a previous time step + value = 0.0; + value += prev[gid + 1] - 2.0 * prev[gid] + prev[gid - 1]; + value += prev[gid + n_cols] - 2.0 * prev[gid] + prev[gid - n_cols]; + value *= dtDIVdxy * vel[gid]; + next[gid] = 2.0f * prev[gid] - next[gid] + value; + } +} + +int main(int argc, char* argv[]) { + // Arrays used to update the wavefield + float* prev_base; + float* next_base; + float* next_cpu; + // Array to store wave velocity + float* vel_base; + + bool error = false; + + size_t n_rows, n_cols; + unsigned int n_iterations; + + // Read parameters + try { + n_rows = stoi(argv[1]); + n_cols = stoi(argv[2]); + n_iterations = stoi(argv[3]); + } + + catch (...) { + Usage(argv[0]); + return 1; + } + + // Compute the total size of grid + size_t n_size = n_rows * n_cols; + + // Allocate arrays to hold wavefield and velocity + prev_base = new float[n_size]; + next_base = new float[n_size]; + next_cpu = new float[n_size]; + vel_base = new float[n_size]; + + // Compute constant value (delta t)^2 (delta x)^2. To be used in wavefield + // update + float dtDIVdxy = (DT * DT) / (DXY * DXY); + + // Initialize arrays and introduce initial conditions (source) + Initialize(prev_base, next_base, vel_base, n_rows, n_cols); + + cout << "Grid Sizes: " << n_rows << " " << n_cols << "\n"; + cout << "Iterations: " << n_iterations << "\n\n"; + + // Define device selector as 'default' + default_selector device_selector; + + // Create a device queue using DPC++ class queue + queue q(device_selector, dpc_common::exception_handler); + + cout << "Computing wavefield in device ..\n"; + // Display info about device + PrintTargetInfo(q); + + // Start timer + dpc_common::TimeInterval t_offload; + + { // Begin buffer scope + // Create buffers using DPC++ class buffer + buffer b_next(next_base, range(n_size)); + buffer b_prev(prev_base, range(n_size)); + buffer b_vel(vel_base, range(n_size)); + + // Iterate over time steps + for (unsigned int k = 0; k < n_iterations; k += 1) { + // Submit command group for execution + q.submit([&](auto &h) { + // Create accessors + auto next = b_next.get_access(h); + auto prev = b_prev.get_access(h); + auto vel = b_vel.get_access(h); + + // Define local and global range + auto global_range = range<2>(n_rows, n_cols); + + // Send a DPC++ kernel (lambda) for parallel execution + // The function that executes a single iteration is called + // "iso_2dfd_iteration_global" + // alternating the 'next' and 'prev' parameters which effectively + // swaps their content at every iteration. + if (k % 2 == 0) + h.parallel_for(global_range, [=](id<2> it) { + Iso2dfdIterationGlobal(it, next.get_pointer(), + prev.get_pointer(), vel.get_pointer(), + dtDIVdxy, n_rows, n_cols); + }); + else + h.parallel_for(global_range, [=](id<2> it) { + Iso2dfdIterationGlobal(it, prev.get_pointer(), + next.get_pointer(), vel.get_pointer(), + dtDIVdxy, n_rows, n_cols); + }); + }); + + } // end for + + } // buffer scope + + // Wait for commands to complete. Enforce synchronization on the command queue + q.wait_and_throw(); + + // Compute and display time used by device + auto time = t_offload.Elapsed(); + + cout << "Offload time: " << time << " ms\n\n"; + + // Output final wavefield (computed by device) to binary file + ofstream out_file; + out_file.open("wavefield_snapshot.bin", ios::out | ios::binary); + out_file.write(reinterpret_cast(next_base), n_size * sizeof(float)); + out_file.close(); + + // Compute wavefield on CPU (for validation) + + cout << "Computing wavefield in CPU ..\n"; + // Re-initialize arrays + Initialize(prev_base, next_cpu, vel_base, n_rows, n_cols); + + // Compute wavefield on CPU + // Start timer for CPU + dpc_common::TimeInterval t_cpu; + + Iso2dfdIterationCpu(next_cpu, prev_base, vel_base, dtDIVdxy, n_rows, n_cols, + n_iterations); + + // Compute and display time used by CPU + time = t_cpu.Elapsed(); + + cout << "CPU time: " << time << " ms\n\n"; + + // Compute error (difference between final wavefields computed in device and + // CPU) + error = WithinEpsilon(next_base, next_cpu, n_rows, n_cols, half_length, 0.1f); + + // If error greater than threshold (last parameter in error function), report + if (error) + cout << "Final wavefields from device and CPU are different: Error\n"; + else + cout << "Final wavefields from device and CPU are equivalent: Success\n"; + + // Output final wavefield (computed by CPU) to binary file + out_file.open("wavefield_snapshot_cpu.bin", ios::out | ios::binary); + out_file.write(reinterpret_cast(next_cpu), n_size * sizeof(float)); + out_file.close(); + + cout << "Final wavefields (from device and CPU) written to disk\n"; + cout << "Finished.\n"; + + // Cleanup + delete[] prev_base; + delete[] next_base; + delete[] vel_base; + + return error ? 1 : 0; +} From 5ec7f2bbaabb5ca42a95212612e3c382013e6066 Mon Sep 17 00:00:00 2001 From: avillarr Date: Tue, 14 Jul 2020 16:31:39 -0700 Subject: [PATCH 02/17] Adding additional inlclude directories to VS project file Signed-off-by: avillarr --- .../DPC++/StructuredGrids/iso2dfd_dpcpp/iso2dfd.vcxproj | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/iso2dfd.vcxproj b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/iso2dfd.vcxproj index c7e2a1d29f..7d258692a9 100644 --- a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/iso2dfd.vcxproj +++ b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/iso2dfd.vcxproj @@ -85,6 +85,7 @@ true true pch.h + $(ONEAPI_ROOT)dev-utilities\latest\include Console @@ -99,6 +100,7 @@ true true pch.h + $(ONEAPI_ROOT)dev-utilities\latest\include Console @@ -115,6 +117,7 @@ true true pch.h + $(ONEAPI_ROOT)dev-utilities\latest\include Console @@ -133,6 +136,7 @@ true true pch.h + $(ONEAPI_ROOT)dev-utilities\latest\include Console @@ -144,4 +148,4 @@ - \ No newline at end of file + From 03421339922742d1d31ea2347285e754b11cd198 Mon Sep 17 00:00:00 2001 From: avillarr Date: Wed, 15 Jul 2020 19:48:45 -0700 Subject: [PATCH 03/17] Adding statement about dpc_common file and where to find them in README file. Change copyright to 2020 Signed-off-by: avillarr --- .../DPC++/StructuredGrids/iso2dfd_dpcpp/License.txt | 2 +- .../DPC++/StructuredGrids/iso2dfd_dpcpp/README.md | 7 +++++++ 2 files changed, 8 insertions(+), 1 deletion(-) diff --git a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/License.txt b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/License.txt index 0578223382..6e9524bd74 100644 --- a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/License.txt +++ b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/License.txt @@ -1,4 +1,4 @@ -Copyright 2019 Intel Corporation +Copyright 2020 Intel Corporation Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights to use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies of the Software, and to permit persons to whom the Software is furnished to do so, subject to the following conditions: diff --git a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md index 427847d11a..1a5217baa2 100644 --- a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md +++ b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md @@ -93,6 +93,13 @@ Perform the following steps: Right click on the project in Solution explorer and select Rebuild. From top menu select Debug -> Start without Debugging. +>If you see the following error message when compiling this sample: +> +``` +Error 'dpc_common.hpp' file not found +``` +>You need to add the following directory to the list of include folders, that are required by your project, in your project's Visual Studio project property panel. The missing include folder is located at `%ONEAPI_ROOT%\dev-utilities\latest\include` on your development system. + * Build the program using MSBuild Open "x64 Native Tools Command Prompt for VS2017" or "x64 Native Tools Command Prompt for VS2019" Run - MSBuild iso2dfd.sln /t:Rebuild /p:Configuration="Release" From 88b994df604999bed8888fc51596ed3f0ae31601 Mon Sep 17 00:00:00 2001 From: avillarr Date: Wed, 15 Jul 2020 20:04:23 -0700 Subject: [PATCH 04/17] Adding statement about dpc_common file and where to find them in README file. Change copyright to 2020 Signed-off-by: avillarr --- .../DPC++/StructuredGrids/iso2dfd_dpcpp/README.md | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md index 1a5217baa2..c5abea2028 100644 --- a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md +++ b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md @@ -93,12 +93,12 @@ Perform the following steps: Right click on the project in Solution explorer and select Rebuild. From top menu select Debug -> Start without Debugging. ->If you see the following error message when compiling this sample: -> +If you see the following error message when compiling this sample: + ``` Error 'dpc_common.hpp' file not found ``` ->You need to add the following directory to the list of include folders, that are required by your project, in your project's Visual Studio project property panel. The missing include folder is located at `%ONEAPI_ROOT%\dev-utilities\latest\include` on your development system. +You need to add the following directory to the list of include folders, that are required by your project, in your project's Visual Studio project property panel. The missing include folder is located at `%ONEAPI_ROOT%\dev-utilities\latest\include` on your development system. * Build the program using MSBuild Open "x64 Native Tools Command Prompt for VS2017" or "x64 Native Tools Command Prompt for VS2019" From 1bc0350ec470f7c5311763129a096b7a7f30ddde Mon Sep 17 00:00:00 2001 From: avillarr Date: Wed, 15 Jul 2020 21:40:20 -0700 Subject: [PATCH 05/17] Adding statement about dpc_common file and where to find them in README file. Change copyright to 2020 Signed-off-by: avillarr --- .../DPC++/StructuredGrids/iso2dfd_dpcpp/README.md | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md index c5abea2028..861dbc697e 100644 --- a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md +++ b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md @@ -93,12 +93,12 @@ Perform the following steps: Right click on the project in Solution explorer and select Rebuild. From top menu select Debug -> Start without Debugging. -If you see the following error message when compiling this sample: - +>If you see the following error message when compiling this sample: +> ``` Error 'dpc_common.hpp' file not found ``` -You need to add the following directory to the list of include folders, that are required by your project, in your project's Visual Studio project property panel. The missing include folder is located at `%ONEAPI_ROOT%\dev-utilities\latest\include` on your development system. +>You need to add the following directory to the list of include folders, that are required by your project, in your project's Visual Studio project property panel. The missing include folder is located at `%ONEAPI_ROOT%\dev-utilities\latest\include` on your development system. * Build the program using MSBuild Open "x64 Native Tools Command Prompt for VS2017" or "x64 Native Tools Command Prompt for VS2019" @@ -106,7 +106,6 @@ You need to add the following directory to the list of include folders, that are ## Running the Sample - ### Application Parameters You can execute the code with different parameters. For example the following command will run the iso2dfd executable using a 1000x1000 grid size and it will iterate over 2000 time steps. From 7e37fcf4a62967001ea10d607e5686bc5c93248f Mon Sep 17 00:00:00 2001 From: avillarr Date: Wed, 15 Jul 2020 21:43:37 -0700 Subject: [PATCH 06/17] Adding statement about dpc_common file and where to find them in README file. Change copyright to 2020 Signed-off-by: avillarr --- DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md | 1 - 1 file changed, 1 deletion(-) diff --git a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md index 861dbc697e..1dcc4b57d0 100644 --- a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md +++ b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md @@ -86,7 +86,6 @@ Perform the following steps: make clean ``` - ### On a Windows* System Using Visual Studio* Version 2017 or Newer * Build the program using VS2017 or VS2019 Right click on the solution file and open using either VS2017 or VS2019 IDE. From 0e08ef743e853d5a25dfcfaed2b330d0bd751e6c Mon Sep 17 00:00:00 2001 From: avillarr Date: Wed, 15 Jul 2020 21:45:48 -0700 Subject: [PATCH 07/17] Adding statement about dpc_common file and where to find them in README file. Change copyright to 2020 Signed-off-by: avillarr --- DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md index 1dcc4b57d0..604dd14b56 100644 --- a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md +++ b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md @@ -66,7 +66,7 @@ Base Toolkit Get Started Guide (https://devcloud.intel.com/oneapi/get-started/ba Perform the following steps: 1. Build the program using the following `cmake` commands. - ``` + ``` cd iso2dfd_dpcpp && mkdir build && cd build && From 24033eeead6fac23ec15a95a5d6ed887d7c33e71 Mon Sep 17 00:00:00 2001 From: avillarr Date: Mon, 20 Jul 2020 15:04:29 -0700 Subject: [PATCH 08/17] Adding particle-diffusion directory Signed-off-by: avillarr --- .../particle-diffusion/CMakeLists.txt | 19 ++ .../particle-diffusion/License.txt | 7 + .../particle-diffusion/Particle_Diffusion.sln | 25 ++ .../Particle_Diffusion.vcxproj | 165 +++++++++++ .../Particle_Diffusion.vcxproj.filters | 28 ++ .../Particle_Diffusion.vcxproj.user | 15 + .../particle-diffusion/README.md | 172 ++++++++++++ .../StructuredGrids/particle-diffusion/draft | 73 +++++ .../particle-diffusion/sample.json | 30 ++ .../particle-diffusion/src/CMakeLists.txt | 11 + .../particle-diffusion/src/motionsim.cpp | 264 ++++++++++++++++++ 11 files changed, 809 insertions(+) create mode 100644 DirectProgramming/DPC++/StructuredGrids/particle-diffusion/CMakeLists.txt create mode 100644 DirectProgramming/DPC++/StructuredGrids/particle-diffusion/License.txt create mode 100644 DirectProgramming/DPC++/StructuredGrids/particle-diffusion/Particle_Diffusion.sln create mode 100644 DirectProgramming/DPC++/StructuredGrids/particle-diffusion/Particle_Diffusion.vcxproj create mode 100644 DirectProgramming/DPC++/StructuredGrids/particle-diffusion/Particle_Diffusion.vcxproj.filters create mode 100644 DirectProgramming/DPC++/StructuredGrids/particle-diffusion/Particle_Diffusion.vcxproj.user create mode 100644 DirectProgramming/DPC++/StructuredGrids/particle-diffusion/README.md create mode 100644 DirectProgramming/DPC++/StructuredGrids/particle-diffusion/draft create mode 100644 DirectProgramming/DPC++/StructuredGrids/particle-diffusion/sample.json create mode 100644 DirectProgramming/DPC++/StructuredGrids/particle-diffusion/src/CMakeLists.txt create mode 100644 DirectProgramming/DPC++/StructuredGrids/particle-diffusion/src/motionsim.cpp diff --git a/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/CMakeLists.txt b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/CMakeLists.txt new file mode 100644 index 0000000000..2c12a7cbeb --- /dev/null +++ b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/CMakeLists.txt @@ -0,0 +1,19 @@ +cmake_minimum_required (VERSION 3.0) + +if(WIN32) + set(CMAKE_CXX_COMPILER "dpcpp-cl") +else() + set(CMAKE_CXX_COMPILER "dpcpp") + +endif() + +# Set default build type to RelWithDebInfo if not specified +if (NOT CMAKE_BUILD_TYPE) + message (STATUS "Default CMAKE_BUILD_TYPE not set using Release with Debug Info") + set (CMAKE_BUILD_TYPE "RelWithDebInfo" CACHE + STRING "Choose the type of build, options are: None Debug Release RelWithDebInfo MinSizeRel" + FORCE) +endif() + +project (Particle_Diffusion) +add_subdirectory (src) diff --git a/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/License.txt b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/License.txt new file mode 100644 index 0000000000..da5f7c1888 --- /dev/null +++ b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/License.txt @@ -0,0 +1,7 @@ +Copyright 2019 Intel Corporation + +Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights to use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies of the Software, and to permit persons to whom the Software is furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. diff --git a/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/Particle_Diffusion.sln b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/Particle_Diffusion.sln new file mode 100644 index 0000000000..1814266556 --- /dev/null +++ b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/Particle_Diffusion.sln @@ -0,0 +1,25 @@ + +Microsoft Visual Studio Solution File, Format Version 12.00 +# Visual Studio 15 +VisualStudioVersion = 15.0.28307.136 +MinimumVisualStudioVersion = 10.0.40219.1 +Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "Particle_Diffusion", "Particle_Diffusion.vcxproj", "{9A9AAD41-D608-401F-9339-FECF2C3CD8E4}" +EndProject +Global + GlobalSection(SolutionConfigurationPlatforms) = preSolution + Debug|x64 = Debug|x64 + Release|x64 = Release|x64 + EndGlobalSection + GlobalSection(ProjectConfigurationPlatforms) = postSolution + {9A9AAD41-D608-401F-9339-FECF2C3CD8E4}.Debug|x64.ActiveCfg = Debug|x64 + {9A9AAD41-D608-401F-9339-FECF2C3CD8E4}.Debug|x64.Build.0 = Debug|x64 + {9A9AAD41-D608-401F-9339-FECF2C3CD8E4}.Release|x64.ActiveCfg = Release|x64 + {9A9AAD41-D608-401F-9339-FECF2C3CD8E4}.Release|x64.Build.0 = Release|x64 + EndGlobalSection + GlobalSection(SolutionProperties) = preSolution + HideSolutionNode = FALSE + EndGlobalSection + GlobalSection(ExtensibilityGlobals) = postSolution + SolutionGuid = {8E9143FB-0D63-4518-8E00-05D9E8F5CACB} + EndGlobalSection +EndGlobal diff --git a/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/Particle_Diffusion.vcxproj b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/Particle_Diffusion.vcxproj new file mode 100644 index 0000000000..bd8c54ed7f --- /dev/null +++ b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/Particle_Diffusion.vcxproj @@ -0,0 +1,165 @@ + + + + + Debug + x64 + + + Release + x64 + + + + 15.0 + {9a9aad41-d608-401f-9339-fecf2c3cd8e4} + Win32Proj + Particle_Diffusion + $(WindowsSDKVersion.Replace("\","")) + + + + Application + true + oneAPI Data Parallel C++ Compiler + Unicode + + + Application + false + oneAPI Data Parallel C++ Compiler + true + Unicode + + + Application + true + Intel(R) oneAPI DPC++ Compiler + Unicode + No + + + Application + false + Intel(R) oneAPI DPC++ Compiler + true + Unicode + No + + + + + + + + + + + + + + + + + + + + + true + + + true + + + false + + + false + + + + Use + Level3 + Disabled + true + true + pch.h + + + Console + true + + + + + Use + Level3 + Disabled + true + true + pch.h + DisableAllWarnings + $(ONEAPI_ROOT)\mkl\latest\include;$(ONEAPI_ROOT)\mkl\latest\include\rng; %(AdditionalIncludeDirectories) + + + Console + true + /link /libpath:"$(ONEAPI_ROOT)\mkl\latest\lib\intel64" mkl_intel_ilp64.lib mkl_sequential.lib mkl_core.lib mkl_sycl.lib opencl.lib %(AdditionalOptions) + -fsycl-device-code-split;%(SpecifyDevCmplAdditionalOptions) + $(ONEAPI_ROOT)/mkl/latest/lib/intel64/mkl_sycl.lib + + + + + Use + Level3 + MaxSpeed + true + true + true + true + pch.h + + + Console + true + true + true + + + + + Use + Level3 + MaxSpeed + true + true + true + true + pch.h + DisableAllWarnings + $(ONEAPI_ROOT)\mkl\latest\include;$(ONEAPI_ROOT)\mkl\latest\include\rng; %(AdditionalIncludeDirectories) + + + Console + true + true + true + /link /libpath:"$(ONEAPI_ROOT)\mkl\latest\lib\intel64" mkl_intel_ilp64.lib mkl_sequential.lib mkl_core.lib mkl_sycl.lib opencl.lib %(AdditionalOptions) + -fsycl-device-code-split;%(SpecifyDevCmplAdditionalOptions) + $(ONEAPI_ROOT)/mkl/latest/lib/intel64/mkl_sycl.lib + + + + + + + + + + + + + + + \ No newline at end of file diff --git a/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/Particle_Diffusion.vcxproj.filters b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/Particle_Diffusion.vcxproj.filters new file mode 100644 index 0000000000..79fc277815 --- /dev/null +++ b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/Particle_Diffusion.vcxproj.filters @@ -0,0 +1,28 @@ + + + + + {4FC737F1-C7A5-4376-A066-2A32D752A2FF} + cpp;c;cc;cxx;def;odl;idl;hpj;bat;asm;asmx + + + {93995380-89BD-4b04-88EB-625FBE52EBFB} + h;hh;hpp;hxx;hm;inl;inc;ipp;xsd + + + {67DA6AB6-F800-4c08-8B7A-83BB121AAD01} + rc;ico;cur;bmp;dlg;rc2;rct;bin;rgs;gif;jpg;jpeg;jpe;resx;tiff;tif;png;wav;mfcribbon-ms + + + + + Source Files + + + + + + + + + \ No newline at end of file diff --git a/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/Particle_Diffusion.vcxproj.user b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/Particle_Diffusion.vcxproj.user new file mode 100644 index 0000000000..be132d4a55 --- /dev/null +++ b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/Particle_Diffusion.vcxproj.user @@ -0,0 +1,15 @@ + + + + 10000 777 + WindowsLocalDebugger + + + + + 10000 777 + WindowsLocalDebugger + + + + \ No newline at end of file diff --git a/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/README.md b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/README.md new file mode 100644 index 0000000000..3d138c967e --- /dev/null +++ b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/README.md @@ -0,0 +1,172 @@ +# water molecule diffusion sample + +motionsim: Intel® oneAPI DPC++ Language Basics Using a Monte Carlo Simulation + +This code sample implements a simple example of a Monte Carlo simulation of the diffusion of water molecules in tissue. IT reflects basid DPC++ programming as well as some techniques for optimization (API-based programming and Atomic Functions). + +For comprehensive instructions regarding DPC++ Programming, go to +https://software.intel.com/en-us/oneapi-programming-guide +and search based on relevant terms noted in the comments. + + For more information and details: https://software.intel.com/en-us/articles/vectorization-of-monte-carlo-simulation-for-diffusion-weighted-imaging-on-intel-xeon + +| Optimized for | Description +|:--- |:--- +| OS | Linux Ubuntu 18.04; Windows 10 or Windows Server 2017 +| Hardware | Kaby Lake with GEN9 or newer +| Software | Intel Data Parallel C++ Compiler (beta) +| What you will learn | How to offload the computation to GPU using Intel DPC++ compiler +| Time to complete | 15 minutes + +Performance number tabulation [if applicable] + +| motionsim sample | Performance data +|:--- |:--- +| Scalar baseline -O2 | 1.0 +| SYCL | +| OpenMP offload | + +## Purpose + +Using the Monte Carlo simulation, the Particle Diffusion sample provides simulation of the +diffusion of water molecules in tissue This kind of computational experiment can be used to +simulate acquisition of a diffusion signal for dMRI. + +The model for the simulation consists of water molecules moving through a 2D array of cells in a +tissue sample (water molecule diffusion). In this code sample, we use a uniform rectilinear 2D +array of digital cells, where cells are spaced regularly along each direction and are represented +by circles. + +Water molecule diffusion is simulated by defining a number of particles P (simulated water +molecules) at random positions in the grid, followed by random walks of these particles in the +ensemble of cells in the grid. During the random walks, particles can move randomly inside or +outside simulated cells. The positions of these particles at every time step in the simulation, +the number of times they go through a cell membrane (in/out), as well as the time every particle +spends inside and outside cells can be recorded. These measurements are a simple example of +useful information that can be used to simulate an MR signal. + +The Particle Diffusion sample is intended to show the basic elements of the DPC++ programming +language as well as some basic optimizations as generating random numbers in the device (using +functionality from the oneAPI oneMKL library), as well as atomic functions to prevent memory +access inconsistencies. + + +## Key implementation details + +SYCL implementation explained. + +* DPC++ queues (including device selectors and exception handlers). +* DPC++ buffers and accessors. +* The ability to call a function inside a kernel definition and pass accessor arguments as pointers. +* Optimization using API-based programming and Atomic Functions. + + +## License + +This code sample is licensed under MIT license. + + +## Building the `Particle_Diffusion` Program for CPU and GPU + +### Include Files +The include folder is located at `%ONEAPI_ROOT%\dev-utilities\latest\include` on your +development system". + +### Running Samples In DevCloud +If running a sample in the Intel DevCloud, remember that you must specify the compute node (CPU, GPU, +FPGA) as well whether to run in batch or interactive mode. For more information see the Intel® oneAPI +Base Toolkit Get Started Guide (https://devcloud.intel.com/oneapi/get-started/base-toolkit/) + +### On a Linux* System +Perform the following steps: +1. Build the motionsim program using the following `cmake` commands. + ``` + cd Particle_Diffusion && + mkdir build && + cd build && + cmake ../. && + make VERBOSE=1 + ``` + +2. Run the program: + ``` + make run + ``` + +3. Clean the program using: + ``` + make clean + ``` + +### on Windows - Visual Studio 2017 or newer + * Build the program using VS2017 or VS2019 + Right click on the solution file and open using either VS2017 or VS2019 IDE + Right click on the project in Solution explorer and select Rebuild + From top menu select Debug -> Start without Debugging + +>If you see the following error message when compiling this sample: +> +``` +Error 'dpc_common.hpp' file not found +``` +>You need to add the following directory to the list of include folders, that are required by your project, in your project's Visual Studio project property panel. The missing include folder is located at `%ONEAPI_ROOT%\dev-utilities\latest\include` on your development system. + + * Build the program using MSBuild + Open "x64 Native Tools Command Prompt for VS2017" or "x64 Native Tools Command Prompt for VS2019" + Run - MSBuild Particle_Diffusion.sln /t:Rebuild /p:Configuration="Release" + + +## Running the Sample + +### Application Parameters + * You can execute the code with different parameters. For example, the following command will run the motionsim executable using 777 as the seed for the random number generator, and will iterate over 10000 time steps + + ``` + ./src/motionsim.exe 10000 777 + ``` + Usage: src/motionsim.exe + + +### Example of Output +The output grid show the cells indicating the number of times that particles have been inside the cell and how the number of particles diffuses from the center of the grid. + + ``` + type src/motionsim.exe 10000 777 (or type 'make run') + + $ src/motionsim.exe 10000 777 + + Running on:: Intel(R) Gen9 HD Graphics NEO + The Device Max Work Group Size is : 256 + The Device Max EUCount is : 72 + The number of iterations is : 10000 + The number of particles is : 20 + + Offload: Time: 138 + + + ********************** OUTPUT GRID: + + 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 + 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 + 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 + 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 + 0 0 0 0 0 0 0 0 0 0 0 367 27 16 0 0 0 0 0 0 0 + 0 0 0 0 0 0 0 0 0 0 84 750 98 84 0 0 0 0 0 0 0 + 0 0 0 0 0 0 0 0 0 0 669 0 116 55 0 0 0 0 0 0 0 + 0 0 0 0 0 0 0 130 211 250 170 35 30 261 0 0 0 0 0 0 0 + 0 0 0 0 0 0 0 10 353 539 243 809 61 878 174 0 0 0 0 0 0 + 0 0 0 0 0 0 1 118 1628 1050 1678 887 864 272 80 390 0 0 0 0 0 + 0 0 0 0 0 0 0 39 1173 1660 3549 1263 1155 2185 234 0 0 0 0 0 0 + 0 0 0 0 0 0 306 599 609 537 550 1134 1172 1261 13 0 0 0 0 0 0 + 0 0 0 0 0 0 283 120 92 282 851 512 658 872 40 0 0 0 0 0 0 + 0 0 0 0 0 157 284 133 817 151 175 271 147 286 57 0 0 0 0 0 0 + 0 0 0 0 0 0 294 428 0 0 0 0 0 17 0 0 0 0 0 0 0 + 0 0 0 0 0 0 0 364 0 0 0 0 0 0 0 0 0 0 0 0 0 + 0 0 0 0 0 0 0 182 0 0 0 0 0 0 0 0 0 0 0 0 0 + 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 + 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 + 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 + 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 + ``` +$ + diff --git a/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/draft b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/draft new file mode 100644 index 0000000000..8f3e1a6c57 --- /dev/null +++ b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/draft @@ -0,0 +1,73 @@ + +{ + "guid": "9483C0F0-7D63-4E99-86C5-37C40F77B2AE" , + "name": "iso2dfd_dpcpp", + "categories": [ "Toolkit/Intel® oneAPI HPC Toolkit" ], + "description": "ISO2DFD: Intel® oneAPI DPC++ Language Basics Using 2D Finite-Difference-Wave Propagation", + "toolchain": [ "dpcpp" ], + "targetDevice": [ "CPU", "GPU" ], + "languages": [ { "cpp": {} } ], + "os": [ "linux", "windows"], + "builder": [ "ide", "cmake" ], + "ciTests": { + "linux": [{ + "steps": [ + "mkdir build", + "cd build", + "cmake ..", + "make", + "make run" + ] + }], + "windows": [{ + "steps": [ + "MSBuild iso2dfd.sln /t:Rebuild /p:Configuration=\"Release\"", + "cd x64/Release", + "iso2dfd.exe 1000 1000 2000" + ] + }] + + } +} +-------------------------------------- + +linux: + + { + "path": "Compiler/Particle_Diffusion", + "configurations": [ + { + "build": [ + "mkdir build", + "cd build", + "cmake ..", + "make VERBOSE=1" + ], + "run": [ + "make run" + ] + } + ] + }, + + + + +windows: + + + { + "path": "Compiler/Particle_Diffusion", + "configurations": [ + { + "build": [ + "MSBuild Particle_Diffusion.sln /t:Rebuild /p:Configuration=\"Release\"" + ], + "run": [ + "cd x64/Release", + "Particle_Diffusion.exe 10000 777" + ] + } + ] + }, + diff --git a/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/sample.json b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/sample.json new file mode 100644 index 0000000000..fefc7bc9b5 --- /dev/null +++ b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/sample.json @@ -0,0 +1,30 @@ +{ + "guid": "FF17715B-10DD-43E2-A408-67146A93702B", + "name": "Particle-Diffusion", + "categories": ["Toolkit/Intel® oneAPI HPC Toolkit"], + "description": "This code sample shows a simple (non-optimized) implementation of a Monte Carlo simulation of the diffusion of water molecules in tissue.", + "toolchain": ["dpcpp"], + "languages": [{"cpp":{}}], + "targetDevice": ["CPU", "GPU"], + "os": ["linux", "windows"], + "builder": ["ide", "cmake"], + "ciTests": { + "linux": [{ + "steps": [ + "mkdir build", + "cd build", + "cmake ..", + "make", + "make run" + ] + }], + "windows": [{ + "steps": [ + "MSBuild Particle_Diffusion.sln /t:Rebuild /p:Configuration=\"Release\"", + "cd x64/Release", + "Particle_Diffusion.exe 10000 777" + ] + }] + + } +} diff --git a/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/src/CMakeLists.txt b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/src/CMakeLists.txt new file mode 100644 index 0000000000..0100583367 --- /dev/null +++ b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/src/CMakeLists.txt @@ -0,0 +1,11 @@ +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fsycl -std=c++17 -O3") +set(CMAKE_EXE_LINKER_FLAGS "${CMAKE_EXE_LINKER_FLAGS}") +set(CMAKE_EXE_LINKER_FLAGS " -L${MKLROOT}/lib/intel64 -lmkl_sycl -lmkl_intel_ilp64 -lmkl_sequential -lmkl_core -lsycl -lOpenCL -lpthread -lm -ldl") + +add_executable (motionsim.exe motionsim.cpp ) +target_link_libraries(motionsim.exe OpenCL sycl) +if(WIN32) + add_custom_target (run motionsim.exe 10000 777) +else() + add_custom_target (run ./motionsim.exe 10000 777) +endif() diff --git a/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/src/motionsim.cpp b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/src/motionsim.cpp new file mode 100644 index 0000000000..6332290484 --- /dev/null +++ b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/src/motionsim.cpp @@ -0,0 +1,264 @@ +//============================================================== +// Copyright © 2019 Intel Corporation +// +// SPDX-License-Identifier: MIT +// ============================================================= + +// motionsim: Intel® oneAPI DPC++ Language Basics Using a Monte Carlo +// Simulation +// +// This code sample will implement a simple example of a Monte Carlo +// simulation of the diffusion of water molecules in tissue. +// +// For comprehensive instructions regarding DPC++ Programming, go to +// https://software.intel.com/en-us/oneapi-programming-guide +// and search based on relevant terms noted in the comments. +// +// DPC++ material used in this code sample: +// +// Basic structures of DPC++: +// DPC++ Queues (including device selectors and exception handlers) +// DPC++ Buffers and accessors (communicate data between the host and the +// device) DPC++ Kernels (including parallel_for function and range<2> +// objects) API-based programming: Use oneMKL to generate random numbers +// (DPC++) DPC++ atomic operations for synchronization +// + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include "dpc_common.hpp" + +using namespace sycl; +using namespace std; + +// Helper functions + +// This function displays correct usage and parameters +void Usage(string program_name) { + cout << " Incorrect number of parameters \n Usage: "; + cout << program_name << " \n\n"; +} + +// This function prints a 1D vector as a matrix +template +void PrintVectorAsMatrix(T* vector, size_t size_X, size_t size_Y) { + cout << "\n"; + for (size_t j = 0; j < size_X; ++j) { + for (size_t i = 0; i < size_Y; ++i) { + cout << std::setw(3) << vector[j * size_Y + i] << " "; + } + cout << "\n"; + } +} + +// This function prints a 2D matrix +template +void PrintMatrix(T** matrix, size_t size_X, size_t size_Y) { + cout << "\n"; + for (size_t i = 0; i < size_X; ++i) { + for (size_t j = 0; j < size_Y; ++j) { + cout << std::setw(3) << matrix[i][j] << " "; + } + cout << "\n"; + } +} + +// This function prints a vector +template +void PrintVector(T* vector, size_t n) { + cout << "\n"; + for (size_t i = 0; i < n; ++i) { + cout << vector[i] << " "; + } + cout << "\n"; +} + +// This function distributes simulation work across workers +void ParticleMotion(queue& q, size_t seed, float* particle_X, float* particle_Y, + size_t* grid, size_t grid_size, size_t n_particles, + unsigned int n_iterations, float radius) { + auto device = q.get_device(); + auto maxBlockSize = device.get_info(); + auto maxEUCount = device.get_info(); + + // Total number of motion events + const size_t n_moves = n_particles * n_iterations; + + cout << " Running on:: " << device.get_info() << "\n"; + cout << " The Device Max Work Group Size is : " << maxBlockSize << "\n"; + cout << " The Device Max EUCount is : " << maxEUCount << "\n"; + cout << " The number of iterations is : " << n_iterations << "\n"; + cout << " The number of particles is : " << n_particles << "\n"; + + // Declare vectors to store random values for X and Y directions + float* random_X = new float[n_moves]; + float* random_Y = new float[n_moves]; + + // Declare RNG object to compute vectors of random values + // Basic random number generator object + mkl::rng::philox4x32x10 engine(q, seed); + // Distribution object + mkl::rng::gaussian distr(0.0, .03); + + // Begin scope for buffers + { + // Create buffers using DPC++ class buffer + buffer b_random_X(random_X, range(n_moves)); + buffer b_random_Y(random_Y, range(n_moves)); + buffer b_particle_X(particle_X, range(n_particles)); + buffer b_particle_Y(particle_Y, range(n_particles)); + buffer b_grid(grid, range(grid_size * grid_size)); + + // Compute vectors of random values for X and Y directions using RNG engine + // declared above + mkl::rng::generate(distr, engine, n_moves, b_random_X); + mkl::rng::generate(distr, engine, n_moves, b_random_Y); + + // Submit command group for execution + q.submit([&](handler& h) { + auto a_particle_X = b_particle_X.get_access(h); + auto a_particle_Y = b_particle_Y.get_access(h); + auto a_random_X = b_random_X.get_access(h); + auto a_random_Y = b_random_Y.get_access(h); + // Atomic accessors: Use DPC++ atomic access mode + auto a_grid = b_grid.get_access(h); + + // Send a DPC++ kernel (lambda) for parallel execution + h.parallel_for(range(n_particles), [=](id<1> index) { + int ii = index.get(0); + float displacement_X = 0.0f; + float displacement_Y = 0.0f; + + // Start iterations + // Each iteration: + // 1. Updates the position of all water molecules + // 2. Checks if water molecule is inside a cell or not. + // 3. Updates counter in cells array + // + for (size_t iter = 0; iter < n_iterations; iter++) { + // Computes random displacement for each molecule + // This example shows random distances between + // -0.05 units and 0.05 units in both X and Y directions + // Moves each water molecule by a random vector + + // Transform the random numbers into small displacements + displacement_X = a_random_X[iter * n_particles + ii]; + displacement_Y = a_random_Y[iter * n_particles + ii]; + + // Move particles using random displacements + a_particle_X[ii] += displacement_X; + a_particle_Y[ii] += displacement_Y; + + // Compute distances from particle position to grid point + float dX = a_particle_X[ii] - sycl::trunc(a_particle_X[ii]); + float dY = a_particle_Y[ii] - sycl::trunc(a_particle_Y[ii]); + + // Compute grid point indices + int iX = sycl::floor(a_particle_X[ii]); + int iY = sycl::floor(a_particle_Y[ii]); + + // Check if particle is still in computation grid + if ((a_particle_X[ii] < grid_size) && + (a_particle_Y[ii] < grid_size) && (a_particle_X[ii] >= 0) && + (a_particle_Y[ii] >= 0)) { + // Check if particle is (or remained) inside cell. + // Increment cell counter in map array if so + if ((dX * dX + dY * dY <= radius * radius)) { + // Use DPC++ atomic_fetch_add to add 1 to accessor using atomic + // mode + atomic_fetch_add(a_grid[iY * grid_size + iX], 1); + } + } + } // Next iteration + }); // End parallel for + }); // End queue submit. + + } // End scope for buffers + + delete[] random_X; + delete[] random_Y; +} // End of function ParticleMotion() + +int main(int argc, char* argv[]) { + // Cell and Particle parameters + const size_t grid_size = 21; // Size of square grid + const size_t n_particles = 20; // Number of particles + const float radius = 0.5f; // Cell radius = 0.5*(grid spacing) + + // Default number of operations + size_t n_iterations = 50; + // Default seed for RNG + size_t seed = 777; + + // Read command-line arguments + try { + n_iterations = stoi(argv[1]); + seed = stoi(argv[2]); + } catch (...) { + Usage(argv[0]); + return 1; + } + + // Allocate arrays + + // Stores a grid of cells + size_t* grid = new size_t[grid_size * grid_size]; + + // Stores X and Y position of particles in the cell grid + float* particle_X = new float[n_particles]; + float* particle_Y = new float[n_particles]; + + // Initialize arrays + for (size_t i = 0; i < n_particles; i++) { + // Initial position of particles in cell grid + particle_X[i] = 10.0f; + particle_Y[i] = 10.0f; + } + + for (size_t y = 0; y < grid_size; y++) { + for (size_t x = 0; x < grid_size; x++) { + grid[y * grid_size + x] = 0; + } + } + + // Create a device queue using default or host/gpu selectors + default_selector device_selector; + + // Create a device queue using DPC++ class queue + queue q(device_selector, dpc_common::exception_handler); + + // Start timers + dpc_common::TimeInterval t_offload; + + // Call simulation function + ParticleMotion(q, seed, particle_X, particle_Y, grid, grid_size, n_particles, + n_iterations, radius); + + q.wait_and_throw(); + + // End timers + auto time = t_offload.Elapsed(); + + cout << "\n Offload time: " << time << " s\n\n"; + + // Displays final grid only if grid small. + if (grid_size <= 45) { + cout << "\n ********************** OUTPUT GRID: \n "; + PrintVectorAsMatrix(grid, grid_size, grid_size); + } + + // Cleanup + delete[] grid; + delete[] particle_X; + delete[] particle_Y; + + return 0; +} From f3a6a8602bd2f4d7c73477fd1c3d4cfc929febf6 Mon Sep 17 00:00:00 2001 From: avillarr Date: Mon, 20 Jul 2020 16:49:40 -0700 Subject: [PATCH 09/17] Adding particle-diffusion directory Signed-off-by: avillarr --- .../DPC++/StructuredGrids/particle-diffusion/README.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/README.md b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/README.md index 3d138c967e..50c61fa567 100644 --- a/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/README.md +++ b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/README.md @@ -141,7 +141,7 @@ The output grid show the cells indicating the number of times that particles hav The number of iterations is : 10000 The number of particles is : 20 - Offload: Time: 138 + Offload: Time: 137 ********************** OUTPUT GRID: From a6e92d3e39a60ea3d853a4a32e0a6e06089e3321 Mon Sep 17 00:00:00 2001 From: avillarr Date: Mon, 20 Jul 2020 21:39:09 -0700 Subject: [PATCH 10/17] Fix copyright -> 2020 Signed-off-by: avillarr --- .../DPC++/StructuredGrids/particle-diffusion/License.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/License.txt b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/License.txt index da5f7c1888..148940418d 100644 --- a/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/License.txt +++ b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/License.txt @@ -1,4 +1,4 @@ -Copyright 2019 Intel Corporation +Copyright 2020 Intel Corporation Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights to use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies of the Software, and to permit persons to whom the Software is furnished to do so, subject to the following conditions: From fcf4462c3e8d16dae66ace355068e5cd35ac867b Mon Sep 17 00:00:00 2001 From: avillarr Date: Tue, 21 Jul 2020 08:47:04 -0700 Subject: [PATCH 11/17] Fixing MKL missing header file Signed-off-by: avillarr --- .../DPC++/StructuredGrids/particle-diffusion/src/motionsim.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/src/motionsim.cpp b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/src/motionsim.cpp index 6332290484..45b8853a33 100644 --- a/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/src/motionsim.cpp +++ b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/src/motionsim.cpp @@ -29,7 +29,6 @@ #include #include #include -#include #include #include #include From 74aae39007498c211d600cd5ae13827eee1e51dd Mon Sep 17 00:00:00 2001 From: avillarr Date: Tue, 21 Jul 2020 11:12:47 -0700 Subject: [PATCH 12/17] Fixing MKL missing header file - try 2 Signed-off-by: avillarr --- .../DPC++/StructuredGrids/particle-diffusion/src/motionsim.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/src/motionsim.cpp b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/src/motionsim.cpp index 45b8853a33..fda492d9e0 100644 --- a/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/src/motionsim.cpp +++ b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/src/motionsim.cpp @@ -30,7 +30,7 @@ #include #include #include -#include +#include #include #include #include "dpc_common.hpp" From c12134143b878ef32364714635839fb8bc4d4fb1 Mon Sep 17 00:00:00 2001 From: avillarr Date: Wed, 22 Jul 2020 23:20:29 -0700 Subject: [PATCH 13/17] Add the common includes directory to VS project file Signed-off-by: avillarr --- .../particle-diffusion/Particle_Diffusion.vcxproj | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/Particle_Diffusion.vcxproj b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/Particle_Diffusion.vcxproj index bd8c54ed7f..5033c25a5b 100644 --- a/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/Particle_Diffusion.vcxproj +++ b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/Particle_Diffusion.vcxproj @@ -99,7 +99,7 @@ true pch.h DisableAllWarnings - $(ONEAPI_ROOT)\mkl\latest\include;$(ONEAPI_ROOT)\mkl\latest\include\rng; %(AdditionalIncludeDirectories) + $(ONEAPI_ROOT)dev-utilities\latest\include;$(ONEAPI_ROOT)\mkl\latest\include;$(ONEAPI_ROOT)\mkl\latest\include\rng; %(AdditionalIncludeDirectories) Console @@ -138,7 +138,7 @@ true pch.h DisableAllWarnings - $(ONEAPI_ROOT)\mkl\latest\include;$(ONEAPI_ROOT)\mkl\latest\include\rng; %(AdditionalIncludeDirectories) + $(ONEAPI_ROOT)dev-utilities\latest\include;$(ONEAPI_ROOT)\mkl\latest\include;$(ONEAPI_ROOT)\mkl\latest\include\rng; %(AdditionalIncludeDirectories) Console @@ -162,4 +162,4 @@ - \ No newline at end of file + From bbf202e1b6fdec581d985bc37a0d2f9c60182963 Mon Sep 17 00:00:00 2001 From: Alberto Villarreal Date: Sun, 16 Aug 2020 15:11:50 -0700 Subject: [PATCH 14/17] Fix copyright and time units Signed-off-by: Alberto Villarreal --- .../DPC++/StructuredGrids/iso2dfd_dpcpp/License.txt | 2 +- .../DPC++/StructuredGrids/iso2dfd_dpcpp/src/iso2dfd.cpp | 4 ++-- .../DPC++/StructuredGrids/particle-diffusion/License.txt | 2 +- 3 files changed, 4 insertions(+), 4 deletions(-) diff --git a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/License.txt b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/License.txt index 6e9524bd74..415025cf03 100644 --- a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/License.txt +++ b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/License.txt @@ -1,4 +1,4 @@ -Copyright 2020 Intel Corporation +Copyright Intel Corporation Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights to use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies of the Software, and to permit persons to whom the Software is furnished to do so, subject to the following conditions: diff --git a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/src/iso2dfd.cpp b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/src/iso2dfd.cpp index e4638ba703..710d87051b 100644 --- a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/src/iso2dfd.cpp +++ b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/src/iso2dfd.cpp @@ -327,7 +327,7 @@ int main(int argc, char* argv[]) { // Compute and display time used by device auto time = t_offload.Elapsed(); - cout << "Offload time: " << time << " ms\n\n"; + cout << "Offload time: " << time << " s\n\n"; // Output final wavefield (computed by device) to binary file ofstream out_file; @@ -351,7 +351,7 @@ int main(int argc, char* argv[]) { // Compute and display time used by CPU time = t_cpu.Elapsed(); - cout << "CPU time: " << time << " ms\n\n"; + cout << "CPU time: " << time << " s\n\n"; // Compute error (difference between final wavefields computed in device and // CPU) diff --git a/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/License.txt b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/License.txt index 148940418d..e63c6e13dc 100644 --- a/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/License.txt +++ b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/License.txt @@ -1,4 +1,4 @@ -Copyright 2020 Intel Corporation +Copyright Intel Corporation Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights to use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies of the Software, and to permit persons to whom the Software is furnished to do so, subject to the following conditions: From d2840e391d2a3008c8d0b5b86633416d79da1686 Mon Sep 17 00:00:00 2001 From: Alberto Villarreal Date: Mon, 17 Aug 2020 13:51:45 -0700 Subject: [PATCH 15/17] Removing unused header file chrono Signed-off-by: Alberto Villarreal --- .../DPC++/StructuredGrids/iso2dfd_dpcpp/src/iso2dfd.cpp | 1 - .../DPC++/StructuredGrids/particle-diffusion/src/motionsim.cpp | 1 - 2 files changed, 2 deletions(-) diff --git a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/src/iso2dfd.cpp b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/src/iso2dfd.cpp index 710d87051b..62bd936ccf 100644 --- a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/src/iso2dfd.cpp +++ b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/src/iso2dfd.cpp @@ -31,7 +31,6 @@ #include #include #include -#include #include #include #include diff --git a/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/src/motionsim.cpp b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/src/motionsim.cpp index fda492d9e0..efbdb7c728 100644 --- a/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/src/motionsim.cpp +++ b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/src/motionsim.cpp @@ -25,7 +25,6 @@ // #include -#include #include #include #include From 4150f950e460add0afd4d12c94cfc5ad0fb42bea Mon Sep 17 00:00:00 2001 From: Alberto Villarreal Date: Mon, 17 Aug 2020 15:06:15 -0700 Subject: [PATCH 16/17] Adding path for dpc_common header file for Linux Signed-off-by: Alberto Villarreal --- .../DPC++/StructuredGrids/iso2dfd_dpcpp/README.md | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md index 604dd14b56..611574eb2f 100644 --- a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md +++ b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md @@ -53,9 +53,12 @@ global ID variable) for a single time step. This code sample is licensed under MIT license. - ## Building the `iso2dfd` Program for CPU and GPU +### Include Files + +The include folder is located at %ONEAPI_ROOT%\dev-utilities\latest\include on your development system. + ### Running Samples In DevCloud If running a sample in the Intel DevCloud, remember that you must specify the compute node (CPU, GPU, From 36ef52360d23702958ae8b49607c96a4990b395e Mon Sep 17 00:00:00 2001 From: Alberto Villarreal Date: Mon, 17 Aug 2020 15:10:56 -0700 Subject: [PATCH 17/17] Adding path for dpc_common header file for Linux and removing warning Signed-off-by: Alberto Villarreal --- .../DPC++/StructuredGrids/iso2dfd_dpcpp/README.md | 12 ------------ .../StructuredGrids/particle-diffusion/README.md | 12 ------------ 2 files changed, 24 deletions(-) diff --git a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md index 611574eb2f..43ca9e2d4e 100644 --- a/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md +++ b/DirectProgramming/DPC++/StructuredGrids/iso2dfd_dpcpp/README.md @@ -95,18 +95,6 @@ Perform the following steps: Right click on the project in Solution explorer and select Rebuild. From top menu select Debug -> Start without Debugging. ->If you see the following error message when compiling this sample: -> -``` -Error 'dpc_common.hpp' file not found -``` ->You need to add the following directory to the list of include folders, that are required by your project, in your project's Visual Studio project property panel. The missing include folder is located at `%ONEAPI_ROOT%\dev-utilities\latest\include` on your development system. - -* Build the program using MSBuild - Open "x64 Native Tools Command Prompt for VS2017" or "x64 Native Tools Command Prompt for VS2019" - Run - MSBuild iso2dfd.sln /t:Rebuild /p:Configuration="Release" - - ## Running the Sample ### Application Parameters diff --git a/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/README.md b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/README.md index 50c61fa567..fc1e4910a4 100644 --- a/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/README.md +++ b/DirectProgramming/DPC++/StructuredGrids/particle-diffusion/README.md @@ -104,18 +104,6 @@ Perform the following steps: Right click on the project in Solution explorer and select Rebuild From top menu select Debug -> Start without Debugging ->If you see the following error message when compiling this sample: -> -``` -Error 'dpc_common.hpp' file not found -``` ->You need to add the following directory to the list of include folders, that are required by your project, in your project's Visual Studio project property panel. The missing include folder is located at `%ONEAPI_ROOT%\dev-utilities\latest\include` on your development system. - - * Build the program using MSBuild - Open "x64 Native Tools Command Prompt for VS2017" or "x64 Native Tools Command Prompt for VS2019" - Run - MSBuild Particle_Diffusion.sln /t:Rebuild /p:Configuration="Release" - - ## Running the Sample ### Application Parameters