Merge pull request #2522 from IgorOchocki/asan

Address Sanitizer Sample

Author: jimmytwei

DirectProgramming/C++SYCL/DenseLinearAlgebra/address_sanitizer/CMakeLists.txt

@@ -0,0 +1,17 @@
+cmake_minimum_required (VERSION 3.5)
+
+set(CMAKE_CXX_COMPILER icpx)
+
+# 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(address_sanitizer)
+
+option(GPU_SELECTOR "Option description" OFF)
+option(CPU_SELECTOR "Option description" OFF)
+
+add_subdirectory (src)

DirectProgramming/C++SYCL/DenseLinearAlgebra/address_sanitizer/License.txt

@@ -0,0 +1,7 @@
+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:
+
+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

DirectProgramming/C++SYCL/DenseLinearAlgebra/address_sanitizer/README.md

@@ -0,0 +1,135 @@
+# `Address Sanitizer` Sample
+
+The `Address Sanitizer` sample demonstrates how to use the AddressSanitizer (ASan) memory detector with the SYCL library.
+
+| Area                   | Description
+|:---                    |:---
+| What you will learn    | How to use the Address Sanitizer tool
+| Time to complete       | 10 minutes
+| Category               | Memory Management
+
+
+## Purpose
+
+The `Address Sanitizer` sample illustrates how to use Address Sanitizer to manage memory errors with the SYCL library. Each of the examples shows a different error and how to initialize it.
+
+All samples can be run on a CPU or a PVC GPU.
+
+>**Note**: The gfx-driver needed to run ASan is version 1.3.28986 for a level zero GPU, 2024.18.3.0.head.prerelease for a CPU and 24.11.028986 for a GPU
+
+The sample includes nine different mini samples that showcase the usage of ASan.
+
+## Prerequisites
+
+| Optimized for          | Description
+|:---                    |:---
+| OS                     | Ubuntu* 20.04 (or newer)
+| Hardware               | GEN9 (or newer)
+| Software               | Intel® oneAPI DPC++/C++ Compiler
+
+## Key Implementation Details
+
+The basic SYCL implementation explained in the code includes:
+
+- out-of-bounds
+- use-after-free
+- misalign-access
+- double-free
+- bad-free
+- bad-context
+
+
+## Set Environment Variables
+
+When working with the command-line interface (CLI), you should configure the oneAPI toolkits using environment variables. Set up your CLI environment by sourcing the `setvars` script every time you open a new terminal window. This practice ensures that your compiler, libraries, and tools are ready for development.
+
+## Build the `Address Sanitizer` Sample
+
+> **Note**: If you have not already done so, set up your CLI
+> environment by sourcing  the `setvars` script in the root of your oneAPI installation.
+>
+> Linux*:
+> - For system wide installations: `. /opt/intel/oneapi/setvars.sh`
+> - For private installations: ` . ~/intel/oneapi/setvars.sh`
+> - For non-POSIX shells, like csh, use the following command: `bash -c 'source <install-dir>/setvars.sh ; exec csh'`
+>
+> For more information on configuring environment variables, see *[Use the setvars Script with Linux* or macOS*](https://www.intel.com/content/www/us/en/develop/documentation/oneapi-programming-guide/top/oneapi-development-environment-setup/use-the-setvars-script-with-linux-or-macos.html)*.
+
+### Using Visual Studio Code*  (Optional)
+
+You can use Visual Studio Code (VS Code) extensions to set your environment, create launch configurations, and browse and download samples.
+
+The basic steps to build and run a sample using VS Code include:
+
+1. Configure the oneAPI environment with the extension **Environment Configurator for Intel Software Developer Tools**.
+2. Download a sample using the extension **Code Sample Browser for Intel Software Developer Tools**.
+3. Open a terminal in VS Code (**Terminal > New Terminal**).
+4. Run the sample in the VS Code terminal using the instructions below.
+
+To learn more about the extensions and how to configure the oneAPI environment, see the 
+*[Using Visual Studio Code with Intel® oneAPI Toolkits User Guide](https://software.intel.com/content/www/us/en/develop/documentation/using-vs-code-with-intel-oneapi/top.html)*.
+
+### On Linux*
+
+1. Change to the sample directory.
+2. Build the program.
+   ```
+   mkdir build
+   cd build
+   cmake ..
+   make
+   ```
+3. Run the program.
+   ```
+   make run_array_reduction
+   ```   
+6. Clean the project. (Optional)
+   ```
+   make clean
+   ```
+
+> **Note**: List of all samples and command to run them.
+> environment by sourcing  the `setvars` script in the root of your oneAPI installation.
+>| File Name                                      | Run command
+>|:---                                            |:---
+>|`array_reduction.cpp`  			  | make run_array_reduction
+>|`bad_free.cpp`          			  | make run_bad_free
+>|`device_global.cpp`				  | make run_device_global
+>|`group_local.cpp`				  | make run_group_local
+>|`local_stencil.cpp`				  | make run_local_stencil
+>|`map.cpp`					  | make run_map
+>|`matmul_broadcast.cpp`		 	  | make run_matmul_broadcast
+>|`misalign-long.cpp`			 	  | make run_misalign-long
+>|`nd_range_reduction.cpp`			  | make run_nd_range_reduction
+
+## Example Output
+
+The following output is for the arrayreduction.cpp sample.
+```
+Histogram:
+bin[0]: 4
+bin[1]: 4
+bin[2]: 4
+bin[3]: 4
+SUCCESS
+
+====ERROR: DeviceSanitizer: bad-free on address 0x250d040
+  #0 ./array_reduction() [0x405249]
+  #1 /lib/x86_64-linux-gnu/libc.so.6(+0x29d90) [0x735d35229d90]
+  #2 /lib/x86_64-linux-gnu/libc.so.6(__libc_start_main+0x80) [0x735d35229e40]
+  #3 ./array_reduction() [0x4049b5]
+
+0x250d040 may be allocated on Host Memory
+Segmentation fault (core dumped)
+make[3]: *** [src/CMakeFiles/run_array_reduction.dir/build.make:70: src/CMakeFiles/run_array_reduction] Error 139
+make[2]: *** [CMakeFiles/Makefile2:392: src/CMakeFiles/run_array_reduction.dir/all] Error 2
+make[1]: *** [CMakeFiles/Makefile2:399: src/CMakeFiles/run_array_reduction.dir/rule] Error 2
+make: *** [Makefile:254: run_array_reduction] Error 2
+```
+
+## License
+
+Code samples are licensed under the MIT license. See
+[License.txt](https://github.com/oneapi-src/oneAPI-samples/blob/master/License.txt) for details.
+
+Third party program Licenses can be found here: [third-party-programs.txt](https://github.com/oneapi-src/oneAPI-samples/blob/master/third-party-programs.txt).

DirectProgramming/C++SYCL/DenseLinearAlgebra/address_sanitizer/sample.json

@@ -0,0 +1,26 @@
+{
+  "guid": "198ADB7C-9E6E-4714-ACD3-E869C86CB174",
+  "name": "Address Sanitizer",
+  "categories": ["Toolkit/oneAPI Direct Programming/C++SYCL/Dense Linear Algebra"],
+  "description": "Address Sanitizer sample provides examples of error memory detection in SYCL",
+  "toolchain": ["dpcpp"],
+  "os": ["linux"],
+  "targetDevice": ["CPU", "GPU"],
+  "gpuRequired": ["gen9","pvc"],
+  "builder": ["cmake"],
+  "languages": [{"cpp":{}}],
+  "ciTests": {
+    "linux": [
+      {
+        "steps": [
+          "mkdir build",
+          "cd build",
+          "cmake ..",
+          "make",
+          "make run_all"
+        ]
+      }
+    ]
+  },
+  "expertise": "Memory Management"
+}

DirectProgramming/C++SYCL/DenseLinearAlgebra/address_sanitizer/src/CMakeLists.txt

@@ -0,0 +1,52 @@
+set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fsycl -Xarch_device -fsanitize=address -g -O2")
+# 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_EXE_LINKER_FLAGS "${CMAKE_EXE_LINKER_FLAGS}")
+set(ENV{ONEAPI_DEVICE_SELECTOR} "level_zero:cpu")
+# set(ENV{ONEAPI_DEVICE_SELECTOR} "level_zero:gpu")
+
+add_executable(array_reduction array_reduction.cpp)
+add_executable(bad_free bad_free.cpp)
+add_executable(device_global device_global.cpp)
+add_executable(group_local group_local.cpp)
+add_executable(local_stencil local_stencil.cpp)
+add_executable(map map.cpp)
+add_executable(matmul_broadcast matmul_broadcast.cpp)
+add_executable(misalign-long misalign-long.cpp)
+add_executable(nd_range_reduction nd_range_reduction.cpp)
+
+target_link_libraries(array_reduction OpenCL sycl)
+target_link_libraries(bad_free OpenCL sycl)
+target_link_libraries(device_global OpenCL sycl)
+target_link_libraries(group_local OpenCL sycl)
+target_link_libraries(local_stencil OpenCL sycl)
+target_link_libraries(map OpenCL sycl)
+target_link_libraries(matmul_broadcast OpenCL sycl)
+target_link_libraries(misalign-long OpenCL sycl)
+target_link_libraries(nd_range_reduction OpenCL sycl)
+
+add_custom_target(run_all array_reduction
+                    COMMAND bad_free
+                    COMMAND device_global
+                    COMMAND group_local
+                    COMMAND local_stencil
+                    COMMAND map
+                    COMMAND matmul_broadcast
+                    COMMAND misalign-long
+                    COMMAND nd_range_reduction)
+
+add_custom_target(run_array_reduction array_reduction)
+add_custom_target(run_bad_free bad_free)
+add_custom_target(run_device_global device_global)
+add_custom_target(run_group_local group_local)
+add_custom_target(run_local_stencil local_stencil)
+add_custom_target(run_map map)
+add_custom_target(run_matmul_broadcast matmul_broadcast)
+add_custom_target(run_misalign-long misalign-long)
+add_custom_target(run_nd_range_reduction nd_range_reduction)

DirectProgramming/C++SYCL/DenseLinearAlgebra/address_sanitizer/src/array_reduction.cpp

@@ -0,0 +1,45 @@
+// Copyright (C) 2023 Intel Corporation
+
+// SPDX-License-Identifier: MIT
+
+#include <iostream>
+#include <numeric>
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+
+int main() {
+  constexpr size_t N = 16;
+  constexpr size_t B = 4;
+
+  queue q;
+  int* data = malloc_shared<int>(N, q);
+  int* histogram = malloc_shared<int>(B, q);
+  std::iota(data, data + N, 1);
+  std::fill(histogram, histogram + B, 0);
+
+  q.submit([&](handler& h) {
+     // BEGIN CODE SNIP
+     h.parallel_for(
+         range{N},
+         reduction(span<int, B>(histogram, B), plus<>()),
+         [=](id<1> i, auto& histogram) {
+           histogram[data[i] % B]++;
+         });
+     // END CODE SNIP
+   }).wait();
+
+  bool passed = true;
+  std::cout << "Histogram:" << std::endl;
+  for (int b = 0; b < B; ++b) {
+    std::cout << "bin[" << b << "]: " << histogram[b]
+              << std::endl;
+    passed &= (histogram[b] == N / B);
+  }
+  std::cout << ((passed) ? "SUCCESS" : "FAILURE") << "\n";
+
+  free(histogram, q); // first release
+  free(histogram, q); // BUG: second release
+  free(data, q);
+  return (passed) ? 0 : 1;
+}

DirectProgramming/C++SYCL/DenseLinearAlgebra/address_sanitizer/src/bad_free.cpp

@@ -0,0 +1,13 @@
+// REQUIRES: linux
+// RUN: %{build} %device_sanitizer_flags -O0 -g -o %t
+// RUN: env SYCL_PREFER_UR=1 UR_ENABLE_LAYERS=UR_LAYER_ASAN %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK %s
+#include <sycl/sycl.hpp>
+
+constexpr size_t N = 64;
+
+int main() {
+  sycl::queue Q;
+  auto *data = new int[N];
+  sycl::free(data, Q);
+  return 0;
+}

DirectProgramming/C++SYCL/DenseLinearAlgebra/address_sanitizer/src/device_global.cpp

@@ -0,0 +1,24 @@
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+using namespace sycl::ext::oneapi;
+using namespace sycl::ext::oneapi::experimental;
+
+sycl::ext::oneapi::experimental::device_global<
+    int[4], decltype(properties(device_image_scope, host_access_read_write))>
+    dev_global;
+
+int main() {
+  sycl::queue Q;
+
+  Q.submit([&](sycl::handler &h) {
+    h.single_task<class Test>([=]() {
+      dev_global[4] = 42; // BUG: out-of-bounds
+    });
+  }).wait();
+
+  int val;
+  Q.copy(dev_global, &val).wait();
+
+  return 0;
+}

DirectProgramming/C++SYCL/DenseLinearAlgebra/address_sanitizer/src/group_local.cpp

@@ -0,0 +1,21 @@
+#include <sycl/sycl.hpp>
+
+constexpr std::size_t N = 16;
+constexpr std::size_t group_size = 8;
+
+int main() {
+  sycl::queue Q;
+  auto *data = sycl::malloc_host<int>(1, Q);
+
+  Q.submit([&](sycl::handler &h) {
+    h.parallel_for<class MyKernel>(
+        sycl::nd_range<1>(N, group_size), [=](sycl::nd_item<1> item) {
+          auto &ref = *sycl::ext::oneapi::group_local_memory<int[N]>(
+                  item.get_group());
+          ref[item.get_local_linear_id() * 2 + 4] = 42; // BUG: out-of-bounds on local memory
+        });
+  });
+
+  Q.wait();
+  return 0;
+}