Skip to content

Commit 61905ef

Browse files
committed
Merge remote-tracking branch 'intel/sycl' into fix_aspect_ext_intel_free_memory_ignored
2 parents db1f972 + 3c274a8 commit 61905ef

File tree

201 files changed

+1002
-689
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

201 files changed

+1002
-689
lines changed

.github/CODEOWNERS

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -169,12 +169,12 @@ sycl/test-e2e/LLVMIntrinsicLowering/ @intel/dpcpp-spirv-reviewers
169169

170170
# Sanitizer
171171
clang/lib/Driver/SanitizerArgs.cpp @intel/dpcpp-sanitizers-review
172-
libdevice/sanitizer_utils.cpp @intel/dpcpp-sanitizers-review
173-
libdevice/include/asan_libdevice.hpp @intel/dpcpp-sanitizers-review
174-
libdevice/include/sanitizer_utils.hpp @intel/dpcpp-sanitizers-review
172+
libdevice/include/asan_rtl.hpp @intel/dpcpp-sanitizers-review
173+
libdevice/include/sanitizer_defs.hpp @intel/dpcpp-sanitizers-review
174+
libdevice/sanitizer/ @intel/dpcpp-sanitizers-review
175+
llvm/include/llvm/Transforms/Instrumentation/AddressSanitizer.h @intel/dpcpp-sanitizers-review
176+
llvm/include/llvm/Transforms/Instrumentation/AddressSanitizerCommon.h @intel/dpcpp-sanitizers-review
177+
llvm/include/llvm/Transforms/Instrumentation/AddressSanitizerOptions.h @intel/dpcpp-sanitizers-review
175178
llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp @intel/dpcpp-sanitizers-review
176-
sycl/test-e2e/AddressSanitizer/ @intel/dpcpp-sanitizers-review
177179
llvm/test/Instrumentation/AddressSanitizer/ @intel/dpcpp-sanitizers-review
178-
llvm/include/llvm/Transforms/Instrumentation/AddressSanitizerOptions.h @intel/dpcpp-sanitizers-review
179-
llvm/include/llvm/Transforms/Instrumentation/AddressSanitizerCommon.h @intel/dpcpp-sanitizers-review
180-
llvm/include/llvm/Transforms/Instrumentation/AddressSanitizer.h @intel/dpcpp-sanitizers-review
180+
sycl/test-e2e/AddressSanitizer/ @intel/dpcpp-sanitizers-review

clang/lib/Driver/ToolChains/SYCL.cpp

Lines changed: 27 additions & 26 deletions
Original file line numberDiff line numberDiff line change
@@ -658,43 +658,44 @@ SYCL::getDeviceLibraries(const Compilation &C, const llvm::Triple &TargetTriple,
658658
addLibraries(SYCLDeviceAnnotationLibs);
659659

660660
#if !defined(_WIN32)
661+
std::string SanitizeVal;
661662
size_t sanitizer_lib_idx = getSingleBuildTarget();
662663
if (Arg *A = Args.getLastArg(options::OPT_fsanitize_EQ,
663664
options::OPT_fno_sanitize_EQ)) {
664665
if (A->getOption().matches(options::OPT_fsanitize_EQ) &&
665-
A->getValues().size() == 1) {
666-
std::string SanitizeVal = A->getValue();
667-
if (SanitizeVal == "address")
668-
addSingleLibrary(SYCLDeviceAsanLibs[sanitizer_lib_idx]);
669-
}
666+
A->getValues().size() == 1)
667+
SanitizeVal = A->getValue();
670668
} else {
671669
// User can pass -fsanitize=address to device compiler via
672670
// -Xsycl-target-frontend, sanitize device library must be
673671
// linked with user's device image if so.
674-
bool IsDeviceAsanEnabled = false;
675-
auto SyclFEArg = Args.getAllArgValues(options::OPT_Xsycl_frontend);
676-
IsDeviceAsanEnabled = (std::count(SyclFEArg.begin(), SyclFEArg.end(),
677-
"-fsanitize=address") > 0);
678-
if (!IsDeviceAsanEnabled) {
679-
auto SyclFEArgEq = Args.getAllArgValues(options::OPT_Xsycl_frontend_EQ);
680-
IsDeviceAsanEnabled = (std::count(SyclFEArgEq.begin(), SyclFEArgEq.end(),
681-
"-fsanitize=address") > 0);
682-
}
683-
684-
// User can also enable asan for SYCL device via -Xarch_device option.
685-
if (!IsDeviceAsanEnabled) {
686-
auto DeviceArchVals = Args.getAllArgValues(options::OPT_Xarch_device);
687-
for (auto DArchVal : DeviceArchVals) {
688-
if (DArchVal.find("-fsanitize=address") != std::string::npos) {
689-
IsDeviceAsanEnabled = true;
690-
break;
691-
}
672+
std::vector<std::string> EnabledDeviceSanitizers;
673+
674+
// NOTE: "-fsanitize=" applies to all device targets
675+
auto SyclFEArgVals = Args.getAllArgValues(options::OPT_Xsycl_frontend);
676+
auto SyclFEEQArgVals = Args.getAllArgValues(options::OPT_Xsycl_frontend_EQ);
677+
auto ArchDeviceVals = Args.getAllArgValues(options::OPT_Xarch_device);
678+
679+
std::vector<std::string> ArgVals(
680+
SyclFEArgVals.size() + SyclFEEQArgVals.size() + ArchDeviceVals.size());
681+
ArgVals.insert(ArgVals.end(), SyclFEArgVals.begin(), SyclFEArgVals.end());
682+
ArgVals.insert(ArgVals.end(), SyclFEEQArgVals.begin(),
683+
SyclFEEQArgVals.end());
684+
ArgVals.insert(ArgVals.end(), ArchDeviceVals.begin(), ArchDeviceVals.end());
685+
686+
// Driver will report error if address sanitizer and memory sanitizer are
687+
// both enabled, so we only need to check first one here.
688+
for (const std::string &Arg : ArgVals) {
689+
if (Arg.find("-fsanitize=address") != std::string::npos) {
690+
SanitizeVal = "address";
691+
break;
692692
}
693693
}
694-
695-
if (IsDeviceAsanEnabled)
696-
addSingleLibrary(SYCLDeviceAsanLibs[sanitizer_lib_idx]);
697694
}
695+
696+
if (SanitizeVal == "address")
697+
addSingleLibrary(SYCLDeviceAsanLibs[sanitizer_lib_idx]);
698+
698699
#endif
699700

700701
if (isNativeCPU)

clang/lib/Sema/SemaSYCL.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -6969,6 +6969,7 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
69696969
Policy.adjustForCPlusPlusFwdDecl();
69706970
Policy.SuppressTypedefs = true;
69716971
Policy.SuppressUnwrittenScope = true;
6972+
Policy.PrintCanonicalTypes = true;
69726973

69736974
llvm::SmallSet<const VarDecl *, 8> Visited;
69746975
bool EmittedFirstSpecConstant = false;
Lines changed: 32 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,32 @@
1+
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -fsycl-int-footer=%t.footer.h -emit-llvm %s -o -
2+
// RUN: FileCheck -input-file=%t.footer.h %s --check-prefix=CHECK-FOOTER
3+
4+
// This test checks that integration footer is emitted correctly when a
5+
// device_global has an explicit template specialization in template arguments.
6+
7+
#include "sycl.hpp"
8+
9+
namespace sycl {
10+
template <typename T> struct X {};
11+
template <> struct X<int> {};
12+
namespace detail {
13+
struct Y {};
14+
} // namespace detail
15+
template <> struct X<detail::Y> {};
16+
} // namespace sycl
17+
18+
using namespace sycl;
19+
template <typename T, typename = X<detail::Y>> struct Arg1 { T val; };
20+
21+
using namespace sycl::ext::oneapi;
22+
template <typename properties_t>
23+
device_global<properties_t> dev_global;
24+
25+
SYCL_EXTERNAL auto foo() {
26+
(void)dev_global<Arg1<int>>;
27+
}
28+
29+
// CHECK-FOOTER: __sycl_device_global_registration::__sycl_device_global_registration() noexcept {
30+
// CHECK-FOOTER-NEXT: device_global_map::add((void *)&::dev_global<Arg1<int, sycl::X<sycl::detail::Y>>>, "_Z10dev_globalI4Arg1IiN4sycl1XINS1_6detail1YEEEEE");
31+
// CHECK-FOOTER-NEXT: }
32+
// CHECK-FOOTER-NEXT: } // namespace (unnamed)

devops/dependencies-igc-dev.json

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,10 +1,10 @@
11
{
22
"linux": {
33
"igc_dev": {
4-
"github_tag": "igc-dev-0b4b682",
5-
"version": "0b4b682",
6-
"updated_at": "2024-11-17T01:09:50Z",
7-
"url": "https://api.github.com/repos/intel/intel-graphics-compiler/actions/artifacts/2197388704/zip",
4+
"github_tag": "igc-dev-ac93a93",
5+
"version": "ac93a93",
6+
"updated_at": "2024-11-21T02:09:35Z",
7+
"url": "https://api.github.com/repos/intel/intel-graphics-compiler/actions/artifacts/2216471673/zip",
88
"root": "{DEPS_ROOT}/opencl/runtime/linux/oclgpu"
99
}
1010
}

libdevice/cmake/modules/SYCLLibdevice.cmake

Lines changed: 24 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -197,7 +197,7 @@ function(add_devicelibs filename)
197197
cmake_parse_arguments(ARG
198198
""
199199
""
200-
"SRC;EXTRA_OPTS;DEPENDENCIES"
200+
"SRC;EXTRA_OPTS;DEPENDENCIES;SKIP_ARCHS"
201201
${ARGN})
202202

203203
foreach(filetype IN LISTS filetypes)
@@ -209,6 +209,9 @@ function(add_devicelibs filename)
209209
endforeach()
210210

211211
foreach(arch IN LISTS devicelib_arch)
212+
if(arch IN_LIST ARG_SKIP_ARCHS)
213+
continue()
214+
endif()
212215
compile_lib(${filename}-${arch}
213216
FILETYPE bc
214217
SRC ${ARG_SRC}
@@ -229,16 +232,17 @@ set(imf_obj_deps device_imf.hpp imf_half.hpp imf_bf16.hpp imf_rounding_op.hpp im
229232
set(itt_obj_deps device_itt.h spirv_vars.h device.h sycl-compiler)
230233
set(bfloat16_obj_deps sycl-headers sycl-compiler)
231234
if (NOT MSVC AND UR_SANITIZER_INCLUDE_DIR)
232-
set(sanitizer_obj_deps
235+
set(asan_obj_deps
233236
device.h atomic.hpp spirv_vars.h
234-
${UR_SANITIZER_INCLUDE_DIR}/asan_libdevice.hpp
235-
include/sanitizer_utils.hpp
237+
${UR_SANITIZER_INCLUDE_DIR}/asan/asan_libdevice.hpp
238+
include/asan_rtl.hpp
236239
include/spir_global_var.hpp
237240
sycl-compiler)
238241

239242
set(sanitizer_generic_compile_opts ${compile_opts}
240243
-fno-sycl-instrument-device-code
241-
-I${UR_SANITIZER_INCLUDE_DIR})
244+
-I${UR_SANITIZER_INCLUDE_DIR}
245+
-I${CMAKE_CURRENT_SOURCE_DIR})
242246

243247
set(asan_pvc_compile_opts_obj -fsycl -c
244248
${sanitizer_generic_compile_opts}
@@ -346,19 +350,27 @@ if(MSVC)
346350
DEPENDENCIES ${cmath_obj_deps})
347351
else()
348352
if(UR_SANITIZER_INCLUDE_DIR)
353+
# asan jit
349354
add_devicelibs(libsycl-asan
350-
SRC sanitizer_utils.cpp
351-
DEPENDENCIES ${sanitizer_obj_deps}
352-
EXTRA_OPTS -fno-sycl-instrument-device-code -I${UR_SANITIZER_INCLUDE_DIR})
355+
SRC sanitizer/asan_rtl.cpp
356+
DEPENDENCIES ${asan_obj_deps}
357+
SKIP_ARCHS nvptx64-nvidia-cuda
358+
amdgcn-amd-amdhsa
359+
EXTRA_OPTS -fno-sycl-instrument-device-code
360+
-I${UR_SANITIZER_INCLUDE_DIR}
361+
-I${CMAKE_CURRENT_SOURCE_DIR})
362+
363+
# asan aot
353364
set(asan_filetypes obj obj-new-offload bc)
354365
set(asan_devicetypes pvc cpu dg2)
366+
355367
foreach(asan_ft IN LISTS asan_filetypes)
356368
foreach(asan_device IN LISTS asan_devicetypes)
357369
compile_lib_ext(libsycl-asan-${asan_device}
358-
SRC sanitizer_utils.cpp
359-
FILETYPE ${asan_ft}
360-
DEPENDENCIES ${sanitizer_obj_deps}
361-
OPTS ${asan_${asan_device}_compile_opts_${asan_ft}})
370+
SRC sanitizer/asan_rtl.cpp
371+
FILETYPE ${asan_ft}
372+
DEPENDENCIES ${asan_obj_deps}
373+
OPTS ${asan_${asan_device}_compile_opts_${asan_ft}})
362374
endforeach()
363375
endforeach()
364376
endif()

libdevice/include/sanitizer_utils.hpp renamed to libdevice/include/asan_rtl.hpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
//==-- sanitizer_device_utils.hpp - Declaration for sanitizer global var ---==//
1+
//==-- asan_rtl.hpp - Declaration for sanitizer global var ---==//
22
//
33
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
44
// See https://llvm.org/LICENSE.txt for license information.
@@ -7,6 +7,7 @@
77
//===----------------------------------------------------------------------===//
88
#pragma once
99

10+
#include "sanitizer_defs.hpp"
1011
#include "spir_global_var.hpp"
1112
#include <cstdint>
1213

Lines changed: 24 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,24 @@
1+
//==-- sanitizer_defs.hpp - common macros shared by sanitizers ---==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
#pragma once
9+
10+
#include <cstdint>
11+
12+
using uptr = uintptr_t;
13+
using u8 = uint8_t;
14+
using u16 = uint16_t;
15+
using u32 = uint32_t;
16+
using u64 = uint64_t;
17+
using s8 = int8_t;
18+
using s16 = int16_t;
19+
using s32 = int32_t;
20+
using s64 = int64_t;
21+
22+
#define LIKELY(x) __builtin_expect(!!(x), 1)
23+
#define UNLIKELY(x) __builtin_expect(!!(x), 0)
24+
#define NORETURN __declspec(noreturn)

0 commit comments

Comments
 (0)