Skip to content

Commit 5f9b8b7

Browse files
jimmytweijimmytwei
authored
Merge pull request #2522 from IgorOchocki/asan
Address Sanitizer Sample
2 parents a694251 + eeb5a1d commit 5f9b8b7

File tree

14 files changed

+749
-0
lines changed

14 files changed

+749
-0
lines changed

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

Lines changed: 17 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,17 @@
1+
cmake_minimum_required (VERSION 3.5)
2+
3+
set(CMAKE_CXX_COMPILER icpx)
4+
5+
# Set default build type to RelWithDebInfo if not specified
6+
if (NOT CMAKE_BUILD_TYPE)
7+
message (STATUS "Default CMAKE_BUILD_TYPE not set using Release with Debug Info")
8+
set (CMAKE_BUILD_TYPE "RelWithDebInfo" CACHE
9+
STRING "Choose the type of build, options are: None Debug Release RelWithDebInfo MinSizeRel"
10+
FORCE)
11+
endif()
12+
project(address_sanitizer)
13+
14+
option(GPU_SELECTOR "Option description" OFF)
15+
option(CPU_SELECTOR "Option description" OFF)
16+
17+
add_subdirectory (src)

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

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,7 @@
1+
Copyright Intel Corporation
2+
3+
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:
4+
5+
The above copyright notice and this permission notice shall be included in all copies or substantial portions of the Software.
6+
7+
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

Lines changed: 135 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,135 @@
1+
# `Address Sanitizer` Sample
2+
3+
The `Address Sanitizer` sample demonstrates how to use the AddressSanitizer (ASan) memory detector with the SYCL library.
4+
5+
| Area | Description
6+
|:--- |:---
7+
| What you will learn | How to use the Address Sanitizer tool
8+
| Time to complete | 10 minutes
9+
| Category | Memory Management
10+
11+
12+
## Purpose
13+
14+
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.
15+
16+
All samples can be run on a CPU or a PVC GPU.
17+
18+
>**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
19+
20+
The sample includes nine different mini samples that showcase the usage of ASan.
21+
22+
## Prerequisites
23+
24+
| Optimized for | Description
25+
|:--- |:---
26+
| OS | Ubuntu* 20.04 (or newer)
27+
| Hardware | GEN9 (or newer)
28+
| Software | Intel® oneAPI DPC++/C++ Compiler
29+
30+
## Key Implementation Details
31+
32+
The basic SYCL implementation explained in the code includes:
33+
34+
- out-of-bounds
35+
- use-after-free
36+
- misalign-access
37+
- double-free
38+
- bad-free
39+
- bad-context
40+
41+
42+
## Set Environment Variables
43+
44+
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.
45+
46+
## Build the `Address Sanitizer` Sample
47+
48+
> **Note**: If you have not already done so, set up your CLI
49+
> environment by sourcing the `setvars` script in the root of your oneAPI installation.
50+
>
51+
> Linux*:
52+
> - For system wide installations: `. /opt/intel/oneapi/setvars.sh`
53+
> - For private installations: ` . ~/intel/oneapi/setvars.sh`
54+
> - For non-POSIX shells, like csh, use the following command: `bash -c 'source <install-dir>/setvars.sh ; exec csh'`
55+
>
56+
> 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)*.
57+
58+
### Using Visual Studio Code* (Optional)
59+
60+
You can use Visual Studio Code (VS Code) extensions to set your environment, create launch configurations, and browse and download samples.
61+
62+
The basic steps to build and run a sample using VS Code include:
63+
64+
1. Configure the oneAPI environment with the extension **Environment Configurator for Intel Software Developer Tools**.
65+
2. Download a sample using the extension **Code Sample Browser for Intel Software Developer Tools**.
66+
3. Open a terminal in VS Code (**Terminal > New Terminal**).
67+
4. Run the sample in the VS Code terminal using the instructions below.
68+
69+
To learn more about the extensions and how to configure the oneAPI environment, see the
70+
*[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)*.
71+
72+
### On Linux*
73+
74+
1. Change to the sample directory.
75+
2. Build the program.
76+
```
77+
mkdir build
78+
cd build
79+
cmake ..
80+
make
81+
```
82+
3. Run the program.
83+
```
84+
make run_array_reduction
85+
```
86+
6. Clean the project. (Optional)
87+
```
88+
make clean
89+
```
90+
91+
> **Note**: List of all samples and command to run them.
92+
> environment by sourcing the `setvars` script in the root of your oneAPI installation.
93+
>| File Name | Run command
94+
>|:--- |:---
95+
>|`array_reduction.cpp` | make run_array_reduction
96+
>|`bad_free.cpp` | make run_bad_free
97+
>|`device_global.cpp` | make run_device_global
98+
>|`group_local.cpp` | make run_group_local
99+
>|`local_stencil.cpp` | make run_local_stencil
100+
>|`map.cpp` | make run_map
101+
>|`matmul_broadcast.cpp` | make run_matmul_broadcast
102+
>|`misalign-long.cpp` | make run_misalign-long
103+
>|`nd_range_reduction.cpp` | make run_nd_range_reduction
104+
105+
## Example Output
106+
107+
The following output is for the arrayreduction.cpp sample.
108+
```
109+
Histogram:
110+
bin[0]: 4
111+
bin[1]: 4
112+
bin[2]: 4
113+
bin[3]: 4
114+
SUCCESS
115+
116+
====ERROR: DeviceSanitizer: bad-free on address 0x250d040
117+
#0 ./array_reduction() [0x405249]
118+
#1 /lib/x86_64-linux-gnu/libc.so.6(+0x29d90) [0x735d35229d90]
119+
#2 /lib/x86_64-linux-gnu/libc.so.6(__libc_start_main+0x80) [0x735d35229e40]
120+
#3 ./array_reduction() [0x4049b5]
121+
122+
0x250d040 may be allocated on Host Memory
123+
Segmentation fault (core dumped)
124+
make[3]: *** [src/CMakeFiles/run_array_reduction.dir/build.make:70: src/CMakeFiles/run_array_reduction] Error 139
125+
make[2]: *** [CMakeFiles/Makefile2:392: src/CMakeFiles/run_array_reduction.dir/all] Error 2
126+
make[1]: *** [CMakeFiles/Makefile2:399: src/CMakeFiles/run_array_reduction.dir/rule] Error 2
127+
make: *** [Makefile:254: run_array_reduction] Error 2
128+
```
129+
130+
## License
131+
132+
Code samples are licensed under the MIT license. See
133+
[License.txt](https://github.com/oneapi-src/oneAPI-samples/blob/master/License.txt) for details.
134+
135+
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

Lines changed: 26 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,26 @@
1+
{
2+
"guid": "198ADB7C-9E6E-4714-ACD3-E869C86CB174",
3+
"name": "Address Sanitizer",
4+
"categories": ["Toolkit/oneAPI Direct Programming/C++SYCL/Dense Linear Algebra"],
5+
"description": "Address Sanitizer sample provides examples of error memory detection in SYCL",
6+
"toolchain": ["dpcpp"],
7+
"os": ["linux"],
8+
"targetDevice": ["CPU", "GPU"],
9+
"gpuRequired": ["gen9","pvc"],
10+
"builder": ["cmake"],
11+
"languages": [{"cpp":{}}],
12+
"ciTests": {
13+
"linux": [
14+
{
15+
"steps": [
16+
"mkdir build",
17+
"cd build",
18+
"cmake ..",
19+
"make",
20+
"make run_all"
21+
]
22+
}
23+
]
24+
},
25+
"expertise": "Memory Management"
26+
}

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

Lines changed: 52 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,52 @@
1+
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fsycl -Xarch_device -fsanitize=address -g -O2")
2+
# Set default build type to RelWithDebInfo if not specified
3+
if (NOT CMAKE_BUILD_TYPE)
4+
message (STATUS "Default CMAKE_BUILD_TYPE not set using Release with Debug Info")
5+
set (CMAKE_BUILD_TYPE "RelWithDebInfo" CACHE
6+
STRING "Choose the type of build, options are: None Debug Release RelWithDebInfo MinSizeRel"
7+
FORCE)
8+
endif()
9+
10+
set(CMAKE_EXE_LINKER_FLAGS "${CMAKE_EXE_LINKER_FLAGS}")
11+
set(ENV{ONEAPI_DEVICE_SELECTOR} "level_zero:cpu")
12+
# set(ENV{ONEAPI_DEVICE_SELECTOR} "level_zero:gpu")
13+
14+
add_executable(array_reduction array_reduction.cpp)
15+
add_executable(bad_free bad_free.cpp)
16+
add_executable(device_global device_global.cpp)
17+
add_executable(group_local group_local.cpp)
18+
add_executable(local_stencil local_stencil.cpp)
19+
add_executable(map map.cpp)
20+
add_executable(matmul_broadcast matmul_broadcast.cpp)
21+
add_executable(misalign-long misalign-long.cpp)
22+
add_executable(nd_range_reduction nd_range_reduction.cpp)
23+
24+
target_link_libraries(array_reduction OpenCL sycl)
25+
target_link_libraries(bad_free OpenCL sycl)
26+
target_link_libraries(device_global OpenCL sycl)
27+
target_link_libraries(group_local OpenCL sycl)
28+
target_link_libraries(local_stencil OpenCL sycl)
29+
target_link_libraries(map OpenCL sycl)
30+
target_link_libraries(matmul_broadcast OpenCL sycl)
31+
target_link_libraries(misalign-long OpenCL sycl)
32+
target_link_libraries(nd_range_reduction OpenCL sycl)
33+
34+
add_custom_target(run_all array_reduction
35+
COMMAND bad_free
36+
COMMAND device_global
37+
COMMAND group_local
38+
COMMAND local_stencil
39+
COMMAND map
40+
COMMAND matmul_broadcast
41+
COMMAND misalign-long
42+
COMMAND nd_range_reduction)
43+
44+
add_custom_target(run_array_reduction array_reduction)
45+
add_custom_target(run_bad_free bad_free)
46+
add_custom_target(run_device_global device_global)
47+
add_custom_target(run_group_local group_local)
48+
add_custom_target(run_local_stencil local_stencil)
49+
add_custom_target(run_map map)
50+
add_custom_target(run_matmul_broadcast matmul_broadcast)
51+
add_custom_target(run_misalign-long misalign-long)
52+
add_custom_target(run_nd_range_reduction nd_range_reduction)

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

Lines changed: 45 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,45 @@
1+
// Copyright (C) 2023 Intel Corporation
2+
3+
// SPDX-License-Identifier: MIT
4+
5+
#include <iostream>
6+
#include <numeric>
7+
#include <sycl/sycl.hpp>
8+
9+
using namespace sycl;
10+
11+
int main() {
12+
constexpr size_t N = 16;
13+
constexpr size_t B = 4;
14+
15+
queue q;
16+
int* data = malloc_shared<int>(N, q);
17+
int* histogram = malloc_shared<int>(B, q);
18+
std::iota(data, data + N, 1);
19+
std::fill(histogram, histogram + B, 0);
20+
21+
q.submit([&](handler& h) {
22+
// BEGIN CODE SNIP
23+
h.parallel_for(
24+
range{N},
25+
reduction(span<int, B>(histogram, B), plus<>()),
26+
[=](id<1> i, auto& histogram) {
27+
histogram[data[i] % B]++;
28+
});
29+
// END CODE SNIP
30+
}).wait();
31+
32+
bool passed = true;
33+
std::cout << "Histogram:" << std::endl;
34+
for (int b = 0; b < B; ++b) {
35+
std::cout << "bin[" << b << "]: " << histogram[b]
36+
<< std::endl;
37+
passed &= (histogram[b] == N / B);
38+
}
39+
std::cout << ((passed) ? "SUCCESS" : "FAILURE") << "\n";
40+
41+
free(histogram, q); // first release
42+
free(histogram, q); // BUG: second release
43+
free(data, q);
44+
return (passed) ? 0 : 1;
45+
}

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

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,13 @@
1+
// REQUIRES: linux
2+
// RUN: %{build} %device_sanitizer_flags -O0 -g -o %t
3+
// RUN: env SYCL_PREFER_UR=1 UR_ENABLE_LAYERS=UR_LAYER_ASAN %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK %s
4+
#include <sycl/sycl.hpp>
5+
6+
constexpr size_t N = 64;
7+
8+
int main() {
9+
sycl::queue Q;
10+
auto *data = new int[N];
11+
sycl::free(data, Q);
12+
return 0;
13+
}

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

Lines changed: 24 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,24 @@
1+
#include <sycl/sycl.hpp>
2+
3+
using namespace sycl;
4+
using namespace sycl::ext::oneapi;
5+
using namespace sycl::ext::oneapi::experimental;
6+
7+
sycl::ext::oneapi::experimental::device_global<
8+
int[4], decltype(properties(device_image_scope, host_access_read_write))>
9+
dev_global;
10+
11+
int main() {
12+
sycl::queue Q;
13+
14+
Q.submit([&](sycl::handler &h) {
15+
h.single_task<class Test>([=]() {
16+
dev_global[4] = 42; // BUG: out-of-bounds
17+
});
18+
}).wait();
19+
20+
int val;
21+
Q.copy(dev_global, &val).wait();
22+
23+
return 0;
24+
}

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

Lines changed: 21 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,21 @@
1+
#include <sycl/sycl.hpp>
2+
3+
constexpr std::size_t N = 16;
4+
constexpr std::size_t group_size = 8;
5+
6+
int main() {
7+
sycl::queue Q;
8+
auto *data = sycl::malloc_host<int>(1, Q);
9+
10+
Q.submit([&](sycl::handler &h) {
11+
h.parallel_for<class MyKernel>(
12+
sycl::nd_range<1>(N, group_size), [=](sycl::nd_item<1> item) {
13+
auto &ref = *sycl::ext::oneapi::group_local_memory<int[N]>(
14+
item.get_group());
15+
ref[item.get_local_linear_id() * 2 + 4] = 42; // BUG: out-of-bounds on local memory
16+
});
17+
});
18+
19+
Q.wait();
20+
return 0;
21+
}

0 commit comments

Comments
 (0)