diff --git a/buildbot/configure.py b/buildbot/configure.py index 65b84fa2348a6..05a1e471ca7b5 100644 --- a/buildbot/configure.py +++ b/buildbot/configure.py @@ -181,6 +181,8 @@ def do_configure(args): "-DSYCL_ENABLE_KERNEL_FUSION={}".format(sycl_enable_fusion), "-DSYCL_ENABLE_MAJOR_RELEASE_PREVIEW_LIB={}".format(sycl_preview_lib), "-DBUG_REPORT_URL=https://github.com/intel/llvm/issues", + "-DSYCL_PI_UR_USE_FETCH_CONTENT=OFF", + "-DSYCL_PI_UR_SOURCE_DIR=/localdisk2/yzhao/work/sycl_workspace/unified-runtime" ] if args.l0_headers and args.l0_loader: diff --git a/libdevice/sanitizer_utils.cpp b/libdevice/sanitizer_utils.cpp index 29cfc23bb7f4a..4ebdcfbb126c7 100644 --- a/libdevice/sanitizer_utils.cpp +++ b/libdevice/sanitizer_utils.cpp @@ -34,6 +34,12 @@ DeviceGlobal __DeviceSanitizerReportMem; DeviceGlobal __DeviceType; +struct LaunchInfo { + uptr LocalShadowOffset = 0; + uptr LocalShadowOffsetEnd = 0; + DeviceSanitizerReport SPIR_DeviceSanitizerReportMem; +}; + #if defined(__SPIR__) #ifdef __SYCL_DEVICE_ONLY__ @@ -78,6 +84,8 @@ static const __SYCL_CONSTANT__ char __global_shadow_out_of_bound[] = static const __SYCL_CONSTANT__ char __local_shadow_out_of_bound[] = "ERROR: Local shadow memory out-of-bound (ptr: %p -> %p, wg: %d, base: " "%p)\n"; +static const __SYCL_CONSTANT__ char __private_shadow_out_of_bound[] = + "ERROR: Private shadow memory out-of-bound (ptr: %p -> %p, base: %p)\n"; static const __SYCL_CONSTANT__ char __unsupport_device_type[] = "ERROR: Unsupport device type: %d\n"; @@ -108,7 +116,7 @@ inline uptr MemToShadow_CPU(uptr addr, int32_t as) { return __AsanShadowMemoryGlobalStart + (addr >> 3); } -inline uptr MemToShadow_DG2(uptr addr, int32_t as) { +inline uptr MemToShadow_DG2(uptr addr, int32_t as, uptr launch_info) { uptr shadow_ptr = 0; if (addr & (~0xffffffffffff)) { shadow_ptr = @@ -125,7 +133,13 @@ inline uptr MemToShadow_DG2(uptr addr, int32_t as) { return shadow_ptr; } -inline uptr MemToShadow_PVC(uptr addr, int32_t as) { +static __SYCL_CONSTANT__ const char __mem_2_shadow_local[] = + "== wgid: %d (%d, %d, %d)\n"; + +static __SYCL_CONSTANT__ const char __mem_launch_info[] = + "== launch_info: %p (%p %p)\n"; + +inline uptr MemToShadow_PVC(uptr addr, int32_t as, uptr launch_info) { uptr shadow_ptr = 0; if (as == AS_GENERIC) { @@ -165,6 +179,13 @@ inline uptr MemToShadow_PVC(uptr addr, int32_t as) { __spirv_BuiltInWorkgroupId.y * __spirv_BuiltInNumWorkgroups.z + __spirv_BuiltInWorkgroupId.z; + auto plaunch = (__SYCL_GLOBAL__ const LaunchInfo *)launch_info; + const auto __AsanShadowMemoryLocalStart = plaunch->LocalShadowOffset; + const auto __AsanShadowMemoryLocalEnd = plaunch->LocalShadowOffsetEnd; + + __spirv_ocl_printf(__mem_launch_info, plaunch, plaunch->LocalShadowOffset, + plaunch->LocalShadowOffsetEnd); + shadow_ptr = __AsanShadowMemoryLocalStart + ((wg_lid * slm_size) >> 3) + ((addr & (slm_size - 1)) >> 3); @@ -178,18 +199,30 @@ inline uptr MemToShadow_PVC(uptr addr, int32_t as) { return shadow_ptr; } -inline uptr MemToShadow(uptr addr, int32_t as) { +static const __SYCL_CONSTANT__ char __mem_to_shadow_nonzero[] = + "__mem_to_shadow: %p(%d) -> %p : %p\n"; +static const __SYCL_CONSTANT__ char __mem_to_shadow_zero[] = + "__mem_to_shadow: %p(%d) -> %p : --\n"; + +inline uptr MemToShadow(uptr addr, int32_t as, uptr launch_info) { uptr shadow_ptr = 0; if (__DeviceType == DeviceType::CPU) { shadow_ptr = MemToShadow_CPU(addr, as); } else if (__DeviceType == DeviceType::GPU_PVC) { - shadow_ptr = MemToShadow_PVC(addr, as); + shadow_ptr = MemToShadow_PVC(addr, as, launch_info); } else { __spirv_ocl_printf(__unsupport_device_type, (int)__DeviceType); return shadow_ptr; } + // if (shadow_ptr) { + // __spirv_ocl_printf(__mem_to_shadow_nonzero, addr, as, shadow_ptr, + // *(char *)shadow_ptr); + // } else { + // __spirv_ocl_printf(__mem_to_shadow_zero, addr, as, shadow_ptr); + // } + return shadow_ptr; } @@ -220,8 +253,8 @@ bool MemIsZero(const char *beg, uptr size) { return all == 0; } -void print_shadow_memory(uptr addr, int32_t as) { - uptr shadow_address = MemToShadow(addr, as); +void print_shadow_memory(uptr addr, int32_t as, uptr launch_info) { + uptr shadow_address = MemToShadow(addr, as, launch_info); uptr p = shadow_address & (~0xf); __spirv_ocl_printf(__asan_shadow_value_start, addr, as, p); for (int i = 0; i < 0xf; ++i) { @@ -237,12 +270,12 @@ void print_shadow_memory(uptr addr, int32_t as) { } // namespace -bool __asan_region_is_value(uptr addr, int32_t as, std::size_t size, - char value) { +bool __asan_region_is_value(uptr addr, int32_t as, uptr launch_info, + std::size_t size, char value) { if (size == 0) return true; while (size--) { - char *shadow = (char *)MemToShadow(addr, as); + char *shadow = (char *)MemToShadow(addr, as, launch_info); if (*shadow != value) { return false; } @@ -252,16 +285,18 @@ bool __asan_region_is_value(uptr addr, int32_t as, std::size_t size, } static void __asan_internal_report_save( - uptr ptr, int32_t as, const char __SYCL_CONSTANT__ *file, int32_t line, - const char __SYCL_CONSTANT__ *func, bool is_write, uint32_t access_size, - DeviceSanitizerMemoryType memory_type, DeviceSanitizerErrorType error_type, - bool is_recover = false) { + uptr ptr, int32_t as, uptr launch_info, const char __SYCL_CONSTANT__ *file, + int32_t line, const char __SYCL_CONSTANT__ *func, bool is_write, + uint32_t access_size, DeviceSanitizerMemoryType memory_type, + DeviceSanitizerErrorType error_type, bool is_recover = false) { const int Expected = ASAN_REPORT_NONE; int Desired = ASAN_REPORT_START; if (atomicCompareAndSet(&__DeviceSanitizerReportMem.get().Flag, Desired, Expected) == Expected) { + // print_shadow_memory(ptr, as); + int FileLength = 0; int FuncLength = 0; @@ -311,14 +346,14 @@ static void __asan_internal_report_save( /// ASAN Error Reporters /// -void __asan_report_access_error(uptr addr, int32_t as, size_t size, - bool is_write, uptr poisoned_addr, +void __asan_report_access_error(uptr addr, int32_t as, uptr launch_info, + size_t size, bool is_write, uptr poisoned_addr, const char __SYCL_CONSTANT__ *file, int32_t line, const char __SYCL_CONSTANT__ *func, bool is_recover = false) { // Check Error Type - s8 *shadow_address = (s8 *)MemToShadow(poisoned_addr, as); + s8 *shadow_address = (s8 *)MemToShadow(poisoned_addr, as, launch_info); int shadow_value = *shadow_address; if (shadow_value > 0) { shadow_value = *(shadow_address + 1); @@ -376,8 +411,8 @@ void __asan_report_access_error(uptr addr, int32_t as, size_t size, error_type = DeviceSanitizerErrorType::UNKNOWN; } - __asan_internal_report_save(addr, as, file, line, func, is_write, size, - memory_type, error_type, is_recover); + __asan_internal_report_save(addr, as, launch_info, file, line, func, is_write, + size, memory_type, error_type, is_recover); } /// @@ -385,8 +420,9 @@ void __asan_report_access_error(uptr addr, int32_t as, size_t size, /// // NOTE: size < 8 -inline int __asan_address_is_poisoned(uptr a, int32_t as, size_t size) { - auto *shadow_address = (s8 *)MemToShadow(a, as); +inline int __asan_address_is_poisoned(uptr a, int32_t as, uptr launch_info, + size_t size) { + auto *shadow_address = (s8 *)MemToShadow(a, as, launch_info); if (shadow_address) { auto shadow_value = *shadow_address; if (shadow_value) { @@ -398,11 +434,12 @@ inline int __asan_address_is_poisoned(uptr a, int32_t as, size_t size) { } // NOTE: size = 1 -inline int __asan_address_is_poisoned(uptr a, int32_t as) { - return __asan_address_is_poisoned(a, as, 1); +inline int __asan_address_is_poisoned(uptr a, int32_t as, uptr launch_info) { + return __asan_address_is_poisoned(a, as, launch_info, 1); } -inline uptr __asan_region_is_poisoned(uptr beg, int32_t as, size_t size) { +inline uptr __asan_region_is_poisoned(uptr beg, int32_t as, uptr launch_info, + size_t size) { if (!size) return 0; @@ -410,11 +447,11 @@ inline uptr __asan_region_is_poisoned(uptr beg, int32_t as, size_t size) { uptr aligned_b = RoundUpTo(beg, ASAN_SHADOW_GRANULARITY); uptr aligned_e = RoundDownTo(end, ASAN_SHADOW_GRANULARITY); - uptr shadow_beg = MemToShadow(aligned_b, as); + uptr shadow_beg = MemToShadow(aligned_b, as, launch_info); if (!shadow_beg) { return 0; } - uptr shadow_end = MemToShadow(aligned_e, as); + uptr shadow_end = MemToShadow(aligned_e, as, launch_info); if (!shadow_end) { return 0; } @@ -422,8 +459,8 @@ inline uptr __asan_region_is_poisoned(uptr beg, int32_t as, size_t size) { // First check the first and the last application bytes, // then check the ASAN_SHADOW_GRANULARITY-aligned region by calling // MemIsZero on the corresponding shadow. - if (!__asan_address_is_poisoned(beg, as) && - !__asan_address_is_poisoned(end - 1, as) && + if (!__asan_address_is_poisoned(beg, as, launch_info) && + !__asan_address_is_poisoned(end - 1, as, launch_info) && (shadow_end <= shadow_beg || MemIsZero((const char *)shadow_beg, shadow_end - shadow_beg))) return 0; @@ -431,7 +468,7 @@ inline uptr __asan_region_is_poisoned(uptr beg, int32_t as, size_t size) { // The fast check failed, so we have a poisoned byte somewhere. // Find it slowly. for (; beg < end; beg++) - if (__asan_address_is_poisoned(beg, as)) + if (__asan_address_is_poisoned(beg, as, launch_info)) return beg; return 0; @@ -444,18 +481,18 @@ inline uptr __asan_region_is_poisoned(uptr beg, int32_t as, size_t size) { #define ASAN_REPORT_ERROR(type, is_write, size) \ DEVICE_EXTERN_C_NOINLINE void __asan_##type##size( \ uptr addr, int32_t as, const char __SYCL_CONSTANT__ *file, int32_t line, \ - const char __SYCL_CONSTANT__ *func) { \ - if (__asan_address_is_poisoned(addr, as, size)) { \ - __asan_report_access_error(addr, as, size, is_write, addr, file, line, \ - func); \ + const char __SYCL_CONSTANT__ *func, uptr launch_info) { \ + if (__asan_address_is_poisoned(addr, as, launch_info, size)) { \ + __asan_report_access_error(addr, as, launch_info, size, is_write, addr, \ + file, line, func); \ } \ } \ DEVICE_EXTERN_C_NOINLINE void __asan_##type##size##_noabort( \ uptr addr, int32_t as, const char __SYCL_CONSTANT__ *file, int32_t line, \ - const char __SYCL_CONSTANT__ *func) { \ - if (__asan_address_is_poisoned(addr, as, size)) { \ - __asan_report_access_error(addr, as, size, is_write, addr, file, line, \ - func, true); \ + const char __SYCL_CONSTANT__ *func, uptr launch_info) { \ + if (__asan_address_is_poisoned(addr, as, launch_info, size)) { \ + __asan_report_access_error(addr, as, launch_info, size, is_write, addr, \ + file, line, func, true); \ } \ } @@ -469,20 +506,20 @@ ASAN_REPORT_ERROR(store, true, 4) #define ASAN_REPORT_ERROR_BYTE(type, is_write, size) \ DEVICE_EXTERN_C_NOINLINE void __asan_##type##size( \ uptr addr, int32_t as, const char __SYCL_CONSTANT__ *file, int32_t line, \ - const char __SYCL_CONSTANT__ *func) { \ - u##size *shadow_address = (u##size *)MemToShadow(addr, as); \ + const char __SYCL_CONSTANT__ *func, uptr launch_info) { \ + u##size *shadow_address = (u##size *)MemToShadow(addr, as, launch_info); \ if (shadow_address && *shadow_address) { \ - __asan_report_access_error(addr, as, size, is_write, addr, file, line, \ - func); \ + __asan_report_access_error(addr, as, launch_info, size, is_write, addr, \ + file, line, func); \ } \ } \ DEVICE_EXTERN_C_NOINLINE void __asan_##type##size##_noabort( \ uptr addr, int32_t as, const char __SYCL_CONSTANT__ *file, int32_t line, \ - const char __SYCL_CONSTANT__ *func) { \ - u##size *shadow_address = (u##size *)MemToShadow(addr, as); \ + const char __SYCL_CONSTANT__ *func, uptr launch_info) { \ + u##size *shadow_address = (u##size *)MemToShadow(addr, as, launch_info); \ if (shadow_address && *shadow_address) { \ - __asan_report_access_error(addr, as, size, is_write, addr, file, line, \ - func, true); \ + __asan_report_access_error(addr, as, launch_info, size, is_write, addr, \ + file, line, func, true); \ } \ } @@ -494,31 +531,41 @@ ASAN_REPORT_ERROR_BYTE(store, true, 16) #define ASAN_REPORT_ERROR_N(type, is_write) \ DEVICE_EXTERN_C_NOINLINE void __asan_##type##N( \ uptr addr, size_t size, int32_t as, const char __SYCL_CONSTANT__ *file, \ - int32_t line, const char __SYCL_CONSTANT__ *func) { \ - if (auto poisoned_addr = __asan_region_is_poisoned(addr, as, size)) { \ - __asan_report_access_error(addr, as, size, is_write, poisoned_addr, \ - file, line, func); \ + int32_t line, const char __SYCL_CONSTANT__ *func, uptr launch_info) { \ + if (auto poisoned_addr = \ + __asan_region_is_poisoned(addr, as, launch_info, size)) { \ + __asan_report_access_error(addr, as, launch_info, size, is_write, \ + poisoned_addr, file, line, func); \ } \ } \ DEVICE_EXTERN_C_NOINLINE void __asan_##type##N_noabort( \ uptr addr, size_t size, int32_t as, const char __SYCL_CONSTANT__ *file, \ - int32_t line, const char __SYCL_CONSTANT__ *func) { \ - if (auto poisoned_addr = __asan_region_is_poisoned(addr, as, size)) { \ - __asan_report_access_error(addr, as, size, is_write, poisoned_addr, \ - file, line, func, true); \ + int32_t line, const char __SYCL_CONSTANT__ *func, uptr launch_info) { \ + if (auto poisoned_addr = \ + __asan_region_is_poisoned(addr, as, launch_info, size)) { \ + __asan_report_access_error(addr, as, launch_info, size, is_write, \ + poisoned_addr, file, line, func, true); \ } \ } ASAN_REPORT_ERROR_N(load, false) ASAN_REPORT_ERROR_N(store, true) +static const __SYCL_CONSTANT__ char __set_shadow_local0[] = + "LOCAL: __asan_set_shadow_local_memory(%p, %d, %d)\n"; +static const __SYCL_CONSTANT__ char __set_shadow_local1[] = + "LOCAL: memset(%p - %p, 0)\n"; +static const __SYCL_CONSTANT__ char __set_shadow_local2[] = + "LOCAL: memset(%p, %d)\n"; + DEVICE_EXTERN_C_NOINLINE void -__asan_set_shadow_local_memory(uptr ptr, size_t size, - size_t size_with_redzone) { +__asan_set_shadow_local_memory(uptr ptr, size_t size, size_t size_with_redzone, + uptr launch_info) { uptr aligned_size = RoundUpTo(size, ASAN_SHADOW_GRANULARITY); { - auto shadow_address = MemToShadow(ptr + aligned_size, AS_LOCAL); + auto shadow_address = + MemToShadow(ptr + aligned_size, AS_LOCAL, launch_info); auto count = (size_with_redzone - aligned_size) / ASAN_SHADOW_GRANULARITY; for (size_t i = 0; i < count; ++i) { ((u8 *)shadow_address)[i] = kSharedLocalRedzoneMagic; @@ -527,7 +574,7 @@ __asan_set_shadow_local_memory(uptr ptr, size_t size, if (size != aligned_size) { auto user_end = ptr + size - 1; - auto *shadow_end = (s8 *)MemToShadow(user_end, AS_LOCAL); + auto *shadow_end = (s8 *)MemToShadow(user_end, AS_LOCAL, launch_info); *shadow_end = user_end - RoundDownTo(user_end, ASAN_SHADOW_GRANULARITY); } } diff --git a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp index 22080d30e6c57..66d0592d65ca3 100644 --- a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp +++ b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp @@ -87,6 +87,7 @@ #include #include #include +#include using namespace llvm; @@ -1184,6 +1185,7 @@ AddressSanitizerPass::AddressSanitizerPass( PreservedAnalyses AddressSanitizerPass::run(Module &M, ModuleAnalysisManager &MAM) { + // M.dump(); ModuleAddressSanitizer ModuleSanitizer( M, Options.InsertVersionCheck, Options.CompileKernel, Options.Recover, UseGlobalGC, UseOdrIndicator, DestructorKind, ConstructorKind); @@ -1191,6 +1193,88 @@ PreservedAnalyses AddressSanitizerPass::run(Module &M, auto &FAM = MAM.getResult(M).getManager(); const StackSafetyGlobalInfo *const SSGI = ClUseStackSafety ? &MAM.getResult(M) : nullptr; + + SmallVector SpirKernels; + // SmallVector SpirFuncs; + for (Function &F : M) { + if (F.getCallingConv() == CallingConv::SPIR_KERNEL) { + SpirKernels.emplace_back(&F); + } + } + + int LongSize = M.getDataLayout().getPointerSizeInBits(); + auto* IntptrTy = Type::getIntNTy(M.getContext(), LongSize); + + for (auto* F : SpirKernels) { + SmallVector Types; + for (Function::const_arg_iterator I = F->arg_begin(), E = F->arg_end(); + I != E; ++I) { + Types.push_back(I->getType()); + } + + // New Argument + Types.push_back(IntptrTy); + + FunctionType *NewFTy = FunctionType::get(F->getReturnType(), Types, false); + + std::string OrigFuncName = F->getName().str(); + F->setName(OrigFuncName + "_del"); + + Function *NewF = + Function::Create(NewFTy, F->getLinkage(), OrigFuncName, F->getParent()); + NewF->copyAttributesFrom(F); + NewF->copyMetadata(F, 0); + NewF->setCallingConv(F->getCallingConv()); + NewF->setDSOLocal(F->isDSOLocal()); + + // Set original arguments' names. + Function::arg_iterator NewI = NewF->arg_begin(); + for (Function::const_arg_iterator I = F->arg_begin(), E = F->arg_end(); + I != E; ++I, ++NewI) { + NewI->setName(I->getName()); + } + + NewF->splice(NewF->begin(), F); + assert(F->isDeclaration() && + "splice does not work, original function body is not empty!"); + + NewF->setSubprogram(F->getSubprogram()); + + NewF->setComdat(F->getComdat()); + F->setComdat(nullptr); + + F->deleteBody(); + + for (Function::arg_iterator I = F->arg_begin(), E = F->arg_end(), + NI = NewF->arg_begin(); + I != E; ++I, ++NI) { + I->replaceAllUsesWith(&*NI); + } + + // Fixup metadata + IRBuilder<> Builder(M.getContext()); + + auto FixupMetadata = [&NewF](StringRef MDName, Constant* NewV) { + auto *Node = NewF->getMetadata(MDName); + if (!Node) + return; + SmallVector NewMD; + for (unsigned I = 0; I < Node->getNumOperands(); ++I) { + NewMD.emplace_back(Node->getOperand(I)); + } + NewMD.emplace_back(ConstantAsMetadata::get(NewV)); + NewF->setMetadata(MDName, llvm::MDNode::get(NewF->getContext(), NewMD)); + }; + + FixupMetadata("kernel_arg_buffer_location", Builder.getInt32(-1)); + FixupMetadata("kernel_arg_runtime_aligned", Builder.getFalse()); + FixupMetadata("kernel_arg_exclusive_ptr", Builder.getFalse()); + + F->removeFromParent(); + } + + M.dump(); + for (Function &F : M) { AddressSanitizer FunctionSanitizer( M, SSGI, Options.InstrumentationWithCallsThreshold, @@ -1306,6 +1390,10 @@ void AddressSanitizer::AppendDebugInfoToArgs(Instruction *InsertBefore, auto *FuncNameGV = GetOrCreateGlobalString(*M, "__asan_func", demangle(FuncName), ConstantAS); Args.push_back(ConstantExpr::getPointerCast(FuncNameGV, ConstASPtrTy)); + + // Launch Data + auto* F = InsertBefore->getFunction(); + Args.push_back(F->getArg(F->arg_size() - 1)); } Value *AddressSanitizer::memToShadow(Value *Shadow, IRBuilder<> &IRB) { @@ -1368,11 +1456,13 @@ void AddressSanitizer::instrumentSyclAllocateLocalMemory(CallInst *CI) { IRB.CreateCall(CI->getCalledFunction(), {SizeWithRedZone, ConstantInt::get(IntptrTy, Align)}); + auto* F = CI->getFunction(); + /// __asan_set_shadow_local_memory(uptr beg, size_t size, size_t - /// size_with_redzone) + /// size_with_redzone, launch_info) IRB.CreateCall( AsanSetShadowDeviceLocalFunc, - {IRB.CreatePointerCast(NewCI, IntptrTy), Size, SizeWithRedZone}); + {IRB.CreatePointerCast(NewCI, IntptrTy), Size, SizeWithRedZone, F->getArg(F->arg_size() - 1)}); CI->replaceAllUsesWith(NewCI); CI->eraseFromParent(); @@ -2830,20 +2920,22 @@ bool ModuleAddressSanitizer::instrumentModule(Module &M) { // Put the constructor and destructor in comdat if both // (1) global instrumentation is not TU-specific // (2) target is ELF. - if (UseCtorComdat && TargetTriple.isOSBinFormatELF() && CtorComdat) { - if (AsanCtorFunction) { - AsanCtorFunction->setComdat(M.getOrInsertComdat(kAsanModuleCtorName)); - appendToGlobalCtors(M, AsanCtorFunction, Priority, AsanCtorFunction); - } - if (AsanDtorFunction) { - AsanDtorFunction->setComdat(M.getOrInsertComdat(kAsanModuleDtorName)); - appendToGlobalDtors(M, AsanDtorFunction, Priority, AsanDtorFunction); + if (!TargetTriple.isSPIR()) { // SPIR kernel needn't AsanCtorFunction & AsanDtorFunction + if (UseCtorComdat && TargetTriple.isOSBinFormatELF() && CtorComdat) { + if (AsanCtorFunction) { + AsanCtorFunction->setComdat(M.getOrInsertComdat(kAsanModuleCtorName)); + appendToGlobalCtors(M, AsanCtorFunction, Priority, AsanCtorFunction); + } + if (AsanDtorFunction) { + AsanDtorFunction->setComdat(M.getOrInsertComdat(kAsanModuleDtorName)); + appendToGlobalDtors(M, AsanDtorFunction, Priority, AsanDtorFunction); + } + } else { + if (AsanCtorFunction) + appendToGlobalCtors(M, AsanCtorFunction, Priority); + if (AsanDtorFunction) + appendToGlobalDtors(M, AsanDtorFunction, Priority); } - } else { - if (AsanCtorFunction) - appendToGlobalCtors(M, AsanCtorFunction, Priority); - if (AsanDtorFunction) - appendToGlobalDtors(M, AsanDtorFunction, Priority); } return true; @@ -2873,8 +2965,7 @@ void AddressSanitizer::initializeCallbacks(Module &M, const TargetLibraryInfo *T } } - // Extend __asan_load/store arguments: unsigned int address_space, char* - // file, unsigned int line, char* func + // Extend __asan_load/store(unsigned int address_space, char* file, unsigned int line, char* func, void* launch_data) if (TargetTriple.isSPIR()) { constexpr unsigned ConstantAS = 2; auto *Int8PtrTy = Type::getInt8Ty(*C)->getPointerTo(ConstantAS); @@ -2883,11 +2974,13 @@ void AddressSanitizer::initializeCallbacks(Module &M, const TargetLibraryInfo *T Args1.push_back(Int8PtrTy); // file Args1.push_back(Type::getInt32Ty(*C)); // line Args1.push_back(Int8PtrTy); // func + Args1.push_back(IntptrTy); // launch_data Args2.push_back(Type::getInt32Ty(*C)); // address_space Args2.push_back(Int8PtrTy); // file Args2.push_back(Type::getInt32Ty(*C)); // line Args2.push_back(Int8PtrTy); // func + Args2.push_back(IntptrTy); // launch_data } AsanErrorCallbackSized[AccessIsWrite][Exp] = M.getOrInsertFunction( kAsanReportErrorTemplate + ExpStr + TypeStr + "_n" + EndingStr, @@ -3044,7 +3137,8 @@ bool AddressSanitizer::instrumentFunction(Function &F, return false; if (F.getLinkage() == GlobalValue::AvailableExternallyLinkage) return false; if (!ClDebugFunc.empty() && ClDebugFunc == F.getName()) return false; - if (F.getName().starts_with("__asan_")) return false; + if (F.getName().starts_with("__asan_")) + return false; if (F.getName().contains("__sycl_service_kernel__")) return false; diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/matrix_multiply.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/matrix_multiply.cpp new file mode 100644 index 0000000000000..07fb1130dc426 --- /dev/null +++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/matrix_multiply.cpp @@ -0,0 +1,88 @@ +// REQUIRES: linux +// RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O1 -g -o %t +// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s +// RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O2 -g -o %t +// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s +// RUN: %{build} %device_sanitizer_flags -DMALLOC_HOST -O2 -g -o %t +// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-HOST --input-file %t.txt %s +// RUN: %{build} %device_sanitizer_flags -DMALLOC_SHARED -O2 -g -o %t +// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-SHARED --input-file %t.txt %s +#include + +using namespace sycl; + +int main() { + sycl::queue Q; + constexpr std::size_t N = 16; + constexpr std::size_t M = 16; + constexpr std::size_t K = N / 4; +#if defined(MALLOC_HOST) + auto matrixA = (int(*)[N])sycl::malloc_host(N * M, Q); + auto matrixB = (int(*)[N])sycl::malloc_host(N * M, Q); + auto matrixC = (int(*)[N])sycl::malloc_host(N * M, Q); +#elif defined(MALLOC_SHARED) + auto matrixA = (int(*)[N])sycl::malloc_shared(N * M, Q); + auto matrixB = (int(*)[N])sycl::malloc_shared(N * M, Q); + auto matrixC = (int(*)[N])sycl::malloc_shared(N * M, Q); +#elif defined(MALLOC_DEVICE) + auto matrixA = (int(*)[N])sycl::malloc_device(N * M, Q); + auto matrixB = (int(*)[N])sycl::malloc_device(N * M, Q); + auto matrixC = (int(*)[N])sycl::malloc_device(N * M, Q); +#elif defined(MALLOC_SYSTEM) + auto matrixA = (int(*)[N])new int[N * M]; + auto matrixB = (int(*)[N])new int[N * M]; + auto matrixC = (int(*)[N])new int[N * M]; +#else +#error "Must provide malloc type to run the test" +#endif + + Q.single_task([=]() { + for (unsigned m = 0; m < M; ++m) { + for (unsigned n = 0; n < N; ++n) { + matrixA[m][n] = n; + matrixB[m][n] = n + m; + matrixC[m][n] = 0; + } + } + }); + Q.wait(); + + Q.submit([&](sycl::handler &h) { + // Local accessor, for one matrix tile: + constexpr unsigned int tile_size = 16; + local_accessor tileA{tile_size, h}; + h.parallel_for( + nd_range<2>{{M, N}, {1, tile_size}}, [=](nd_item<2> item) { + // Indices in the global index space: + int m = item.get_global_id()[0]; + int n = item.get_global_id()[1]; + // Index in the local index space: + int i = item.get_local_id()[1]; + int sum = 0; + for (unsigned int kk = 0; kk < K; kk += tile_size) { + // Load the matrix tile from matrix A, and synchronize + // to ensure all work-items have a consistent view + // of the matrix tile in local memory. + tileA[i] = matrixA[m][kk + i + 1]; // <== bug add "+1" intentionally + // CHECK-DEVICE: ERROR: DeviceSanitizer: out-of-bounds-access on USM Device Memory + // CHECK-HOST: ERROR: DeviceSanitizer: out-of-bounds-access on USM Host Memory + // CHECK-SHARED: ERROR: DeviceSanitizer: out-of-bounds-access on USM Shared Memory + // CHECK: {{READ of size 4 at kernel <.*MatMultiply> LID\(15, 0, 0\) GID\(15, 15, 0\)}} + // CHECK: {{ #0 .* .*matrix_multiply.cpp:}}[[@LINE-5]] + item.barrier(); + // Perform computation using the local memory tile, and + // matrix B in global memory. + for (unsigned int k = 0; k < tile_size; k++) + sum += tileA[k] * matrixB[kk + k][n]; + // After computation, synchronize again, to ensure all + // reads from the local memory tile are complete. + item.barrier(); + } + // Write the final result to global memory. + matrixC[m][n] = sum; + }); + }); + Q.wait(); + + return 0; +} diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_char.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_char.cpp index 7995db533a4f4..d443a32ae9f38 100644 --- a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_char.cpp +++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_char.cpp @@ -1,3 +1,14 @@ +<<<<<<< HEAD +// REQUIRES: linux +// RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O1 -g -o %t +// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s +// RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O2 -g -o %t +// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s +// RUN: %{build} %device_sanitizer_flags -DMALLOC_HOST -O2 -g -o %t +// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-HOST --input-file %t.txt %s +// RUN: %{build} %device_sanitizer_flags -DMALLOC_SHARED -O2 -g -o %t +// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-SHARED --input-file %t.txt %s +======= // REQUIRES: linux, cpu // RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O1 -g -o %t // RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s @@ -7,6 +18,7 @@ // RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-HOST --input-file %t.txt %s // RUN: %{build} %device_sanitizer_flags -DMALLOC_SHARED -O2 -g -o %t // RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-SHARED --input-file %t.txt %s +>>>>>>> sycl #include int main() { @@ -30,11 +42,19 @@ int main() { [=](sycl::nd_item<1> item) { ++array[item.get_global_id(0)]; }); }); Q.wait(); +<<<<<<< HEAD + // CHECK-DEVICE: ERROR: DeviceSanitizer: out-of-bounds-access on USM Device Memory + // CHECK-HOST: ERROR: DeviceSanitizer: out-of-bounds-access on USM Host Memory + // CHECK-SHARED: ERROR: DeviceSanitizer: out-of-bounds-access on USM Shared Memory + // CHECK: {{READ of size 1 at kernel <.*MyKernelR_4> LID\(0, 0, 0\) GID\(12345, 0, 0\)}} + // CHECK: {{ #0 .* .*parallel_for_char.cpp:}}[[@LINE-8]] +======= // CHECK-DEVICE: ERROR: DeviceSanitizer: out-of-bounds-access on Device USM // CHECK-HOST: ERROR: DeviceSanitizer: out-of-bounds-access on Host USM // CHECK-SHARED: ERROR: DeviceSanitizer: out-of-bounds-access on Shared USM // CHECK: {{READ of size 1 at kernel <.*MyKernelR_4> LID\(0, 0, 0\) GID\(12345, 0, 0\)}} // CHECK: {{ #0 .* .*parallel_for_char.cpp:}}[[@LINE-7]] +>>>>>>> sycl return 0; } diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_double.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_double.cpp index a3d29e3b3e5a5..6b3f26d20d4f2 100644 --- a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_double.cpp +++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_double.cpp @@ -1,3 +1,14 @@ +<<<<<<< HEAD +// REQUIRES: linux +// RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O1 -g -o %t +// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s +// RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O2 -g -o %t +// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s +// RUN: %{build} %device_sanitizer_flags -DMALLOC_HOST -O2 -g -o %t +// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-HOST --input-file %t.txt %s +// RUN: %{build} %device_sanitizer_flags -DMALLOC_SHARED -O2 -g -o %t +// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-SHARED --input-file %t.txt %s +======= // REQUIRES: linux, cpu, aspect-fp64 // RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O1 -g -o %t // RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s @@ -7,6 +18,7 @@ // RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-HOST --input-file %t.txt %s // RUN: %{build} %device_sanitizer_flags -DMALLOC_SHARED -O2 -g -o %t // RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-SHARED --input-file %t.txt %s +>>>>>>> sycl #include int main() { @@ -30,11 +42,19 @@ int main() { [=](sycl::nd_item<1> item) { ++array[item.get_global_id(0)]; }); }); Q.wait(); +<<<<<<< HEAD + // CHECK-DEVICE: ERROR: DeviceSanitizer: out-of-bounds-access on USM Device Memory + // CHECK-HOST: ERROR: DeviceSanitizer: out-of-bounds-access on USM Host Memory + // CHECK-SHARED: ERROR: DeviceSanitizer: out-of-bounds-access on USM Shared Memory + // CHECK: {{READ of size 8 at kernel <.*MyKernelR_4> LID\(0, 0, 0\) GID\(123456, 0, 0\)}} + // CHECK: {{ #0 .* .*parallel_for_double.cpp:}}[[@LINE-8]] +======= // CHECK-DEVICE: ERROR: DeviceSanitizer: out-of-bounds-access on Device USM // CHECK-HOST: ERROR: DeviceSanitizer: out-of-bounds-access on Host USM // CHECK-SHARED: ERROR: DeviceSanitizer: out-of-bounds-access on Shared USM // CHECK: {{READ of size 8 at kernel <.*MyKernelR_4> LID\(0, 0, 0\) GID\(123456, 0, 0\)}} // CHECK: {{ #0 .* .*parallel_for_double.cpp:}}[[@LINE-7]] +>>>>>>> sycl return 0; } diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_func.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_func.cpp index 44c822a22884f..ded72a6caaa65 100644 --- a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_func.cpp +++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_func.cpp @@ -1,3 +1,20 @@ +<<<<<<< HEAD +// REQUIRES: linux +// RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O1 -g -o %t +// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s +// RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O2 -g -o %t +// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s +// RUN: %{build} %device_sanitizer_flags -DMALLOC_HOST -O2 -g -o %t +// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-HOST --input-file %t.txt %s +// RUN: %{build} %device_sanitizer_flags -DMALLOC_SHARED -O2 -g -o %t +// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-SHARED --input-file %t.txt %s +#include + +__attribute__((noinline)) void foo(int *array, size_t i) { array[i] = 1; } +// CHECK-DEVICE: ERROR: DeviceSanitizer: out-of-bounds-access on USM Device Memory +// CHECK-HOST: ERROR: DeviceSanitizer: out-of-bounds-access on USM Host Memory +// CHECK-SHARED: ERROR: DeviceSanitizer: out-of-bounds-access on USM Shared Memory +======= // REQUIRES: linux, cpu // RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O1 -g -o %t // RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s @@ -13,6 +30,7 @@ __attribute__((noinline)) void foo(int *array, size_t i) { array[i] = 1; } // CHECK-DEVICE: ERROR: DeviceSanitizer: out-of-bounds-access on Device USM // CHECK-HOST: ERROR: DeviceSanitizer: out-of-bounds-access on Host USM // CHECK-SHARED: ERROR: DeviceSanitizer: out-of-bounds-access on Shared USM +>>>>>>> sycl // CHECK: {{WRITE of size 4 at kernel <.*MyKernel> LID\(0, 0, 0\) GID\(123, 0, 0\)}} // CHECK: {{ #0 foo\(int\*, unsigned long\) .*parallel_for_func.cpp:}}[[@LINE-5]] diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_int.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_int.cpp index fd0ec83ebdc14..e030884942441 100644 --- a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_int.cpp +++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_int.cpp @@ -1,3 +1,14 @@ +<<<<<<< HEAD +// REQUIRES: linux +// RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O1 -g -o %t +// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s +// RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O2 -g -o %t +// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s +// RUN: %{build} %device_sanitizer_flags -DMALLOC_HOST -O2 -g -o %t +// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-HOST --input-file %t.txt %s +// RUN: %{build} %device_sanitizer_flags -DMALLOC_SHARED -O2 -g -o %t +// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-SHARED --input-file %t.txt %s +======= // REQUIRES: linux, cpu // RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O1 -g -o %t // RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s @@ -7,6 +18,7 @@ // RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-HOST --input-file %t.txt %s // RUN: %{build} %device_sanitizer_flags -DMALLOC_SHARED -O2 -g -o %t // RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-SHARED --input-file %t.txt %s +>>>>>>> sycl #include int main() { @@ -30,11 +42,19 @@ int main() { [=](sycl::nd_item<1> item) { ++array[item.get_global_id(0)]; }); }); Q.wait(); +<<<<<<< HEAD + // CHECK-DEVICE: ERROR: DeviceSanitizer: out-of-bounds-access on USM Device Memory + // CHECK-HOST: ERROR: DeviceSanitizer: out-of-bounds-access on USM Host Memory + // CHECK-SHARED: ERROR: DeviceSanitizer: out-of-bounds-access on USM Shared Memory + // CHECK: {{READ of size 4 at kernel <.*MyKernelR_4> LID\(0, 0, 0\) GID\(1234567, 0, 0\)}} + // CHECK: {{ #0 .* .*parallel_for_int.cpp:}}[[@LINE-8]] +======= // CHECK-DEVICE: ERROR: DeviceSanitizer: out-of-bounds-access on Device USM // CHECK-HOST: ERROR: DeviceSanitizer: out-of-bounds-access on Host USM // CHECK-SHARED: ERROR: DeviceSanitizer: out-of-bounds-access on Shared USM // CHECK: {{READ of size 4 at kernel <.*MyKernelR_4> LID\(0, 0, 0\) GID\(1234567, 0, 0\)}} // CHECK: {{ #0 .* .*parallel_for_int.cpp:}}[[@LINE-7]] +>>>>>>> sycl return 0; } diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_short.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_short.cpp index d90003950419e..132245b48822a 100644 --- a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_short.cpp +++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_short.cpp @@ -1,3 +1,14 @@ +<<<<<<< HEAD +// REQUIRES: linux +// RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O1 -g -o %t +// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s +// RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O2 -g -o %t +// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s +// RUN: %{build} %device_sanitizer_flags -DMALLOC_HOST -O2 -g -o %t +// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-HOST --input-file %t.txt %s +// RUN: %{build} %device_sanitizer_flags -DMALLOC_SHARED -O2 -g -o %t +// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-SHARED --input-file %t.txt %s +======= // REQUIRES: linux, cpu // RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O1 -g -o %t // RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s @@ -7,6 +18,7 @@ // RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-HOST --input-file %t.txt %s // RUN: %{build} %device_sanitizer_flags -DMALLOC_SHARED -O2 -g -o %t // RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-SHARED --input-file %t.txt %s +>>>>>>> sycl #include int main() { @@ -30,11 +42,19 @@ int main() { [=](sycl::nd_item<1> item) { ++array[item.get_global_id(0)]; }); }); Q.wait(); +<<<<<<< HEAD + // CHECK-DEVICE: ERROR: DeviceSanitizer: out-of-bounds-access on USM Device Memory + // CHECK-HOST: ERROR: DeviceSanitizer: out-of-bounds-access on USM Host Memory + // CHECK-SHARED: ERROR: DeviceSanitizer: out-of-bounds-access on USM Shared Memory + // CHECK: {{READ of size 2 at kernel <.*MyKernelR_4> LID\(0, 0, 0\) GID\(123456789, 0, 0\)}} + // CHECK: {{ #0 .* .*parallel_for_short.cpp:}}[[@LINE-8]] +======= // CHECK-DEVICE: ERROR: DeviceSanitizer: out-of-bounds-access on Device USM // CHECK-HOST: ERROR: DeviceSanitizer: out-of-bounds-access on Host USM // CHECK-SHARED: ERROR: DeviceSanitizer: out-of-bounds-access on Shared USM // CHECK: {{READ of size 2 at kernel <.*MyKernelR_4> LID\(0, 0, 0\) GID\(123456789, 0, 0\)}} // CHECK: {{ #0 .* .*parallel_for_short.cpp:}}[[@LINE-7]] +>>>>>>> sycl return 0; } diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/single_task_int.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/single_task_int.cpp new file mode 100644 index 0000000000000..53b8b6aea9c62 --- /dev/null +++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/single_task_int.cpp @@ -0,0 +1,23 @@ +// REQUIRES: linux +// RUN: %{build} %device_sanitizer_flags -g -o %t.out +// RUN: %{run} not %t.out 2>&1 | FileCheck %s + +#include + +const int N = 1024; + +int main() { + sycl::queue Q; + + int *data = sycl::malloc_device(N, Q); + + // CHECK: DeviceSanitizer: out-of-bounds-access on USM Device Memory + Q.single_task([=]() { + for (int i = 0; i <= N; ++i) { + data[i] = i; // <== buffer-overflow here + } + }); + + Q.wait(); + return 0; +} diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/single_task_loop.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/single_task_loop.cpp new file mode 100644 index 0000000000000..6f18c88e11fd3 --- /dev/null +++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/single_task_loop.cpp @@ -0,0 +1,30 @@ +// REQUIRES: linux +// UNSUPPORTED: true +// TODO: rely on "printf_abort" + +// RUN: %{build} %device_sanitizer_flags -g -o %t.out +// RUN: %{run} not %t.out 2>&1 | FileCheck %s + +// XFAIL: * + +#include + +const int N = 1024; + +int main() { + sycl::queue Q; + + int *data = sycl::malloc_device(N, Q); + + // CHECK: DeviceSanitizer: out-of-bounds-access on USM Device Memory + Q.single_task([=]() { + int i = N; + // infinite loop + while (true) { + data[--i] = i; // <== buffer-underflow here + } + }); + + Q.wait(); + return 0; +} diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/single_task_underflow.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/single_task_underflow.cpp new file mode 100644 index 0000000000000..405079a3bd145 --- /dev/null +++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/single_task_underflow.cpp @@ -0,0 +1,25 @@ +// REQUIRES: linux +// RUN: %{build} %device_sanitizer_flags -g -o %t.out +// RUN: %{run} not %t.out 2>&1 | FileCheck %s + +#include + +const int N = 1024; + +int main() { + sycl::queue Q; + + int *data = sycl::malloc_device(N, Q); + + // CHECK: DeviceSanitizer: out-of-bounds-access on USM Device Memory + Q.single_task([=]() { + int i = N; + int n = N + 100; + while (n--) { + data[--i] = i; // <== buffer-underflow here + } + }); + + Q.wait(); + return 0; +} diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/local/local-overflow-1.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/local/local-overflow-1.cpp index b8278b209aa6a..0d59af9e41b7c 100644 --- a/sycl/test-e2e/AddressSanitizer/out-of-bounds/local/local-overflow-1.cpp +++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/local/local-overflow-1.cpp @@ -1,6 +1,11 @@ +<<<<<<< HEAD +// RUN: %{build} %device_sanitizer_flags -g -o %t.out +// RUN: %{run} not %t.out 2>&1 | FileCheck %s +======= // REQUIRES: linux, cpu // RUN: %{build} %device_sanitizer_flags -g -o %t.out // RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t.out 2>&1 | FileCheck %s +>>>>>>> sycl #include constexpr std::size_t N = 16; @@ -19,7 +24,11 @@ int main() { auto &ref = *ptr; ref[item.get_local_linear_id() * 2 + 4] = 42; // CHECK: ERROR: DeviceSanitizer: out-of-bounds-access on Local Memory +<<<<<<< HEAD + // CHECK: {{READ of size 1 at kernel <.*MyKernel> LID\(6, 0, 0\) GID\(6, 0, 0\)}} +======= // CHECK: {{WRITE of size 4 at kernel <.*MyKernel> LID\(6, 0, 0\) GID\(.*, 0, 0\)}} +>>>>>>> sycl // CHECK: {{ #0 .* .*local-overflow-1.cpp:}}[[@LINE-3]] }); }); diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/local/local-overflow-2.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/local/local-overflow-2.cpp new file mode 100644 index 0000000000000..39a964e3cbfad --- /dev/null +++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/local/local-overflow-2.cpp @@ -0,0 +1,26 @@ +// UNSUPPORTED: true +// TODO: support dynamic local + +// RUN: %{build} %device_sanitizer_flags -g -o %t.out +// RUN: %{run} not %t.out 2>&1 | FileCheck %s +#include "sycl/accessor.hpp" +#include +#include + +constexpr std::size_t N = 1024; +constexpr std::size_t group_size = 16; + +int main() { + sycl::queue Q; + + Q.submit([&](sycl::handler &h) { + h.parallel_for_work_group(N / group_size, group_size, [=]() { + int array[N]; + group.parallel_for_work_item( + [&](sycl::h_item item) { ++array[item.get_global_id()]; }); + }); + }); + Q.wait(); + + return 0; +} diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/local/local-overflow-3.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/local/local-overflow-3.cpp new file mode 100644 index 0000000000000..dcc5294f633a0 --- /dev/null +++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/local/local-overflow-3.cpp @@ -0,0 +1,36 @@ +// RUN: %{build} %device_sanitizer_flags -g -o %t.out +// RUN: %{run} not %t.out 2>&1 | FileCheck %s +#include + +constexpr std::size_t N = 16; +constexpr std::size_t group_size = 4; + +int main() { + sycl::queue q; + int *ptr = sycl::malloc_shared(N, q); + + q.parallel_for( + sycl::nd_range<1>{N, group_size}, + [=](sycl::nd_item<1> it) { + auto g = it.get_group(); + auto mem1 = + sycl::ext::oneapi::group_local_memory(g, 1, 2, 3, 4, 5); + auto mem2 = + sycl::ext::oneapi::group_local_memory(g, 1, 2, 3, 4, 5); + auto ref1 = *mem1, ref2 = *mem2; + for (int i = 0; i < N + 1; ++i) { + ptr[i] = ref1[i]; + // CHECK: ERROR: DeviceSanitizer: out-of-bounds-access on Local Memory + // CHECK: {{READ of size 1 at kernel <.*MyKernel> LID\(0, 0, 0\) GID\(.*, 0, 0\)}} + // CHECK: {{ #0 .* .*local-overflow-3.cpp:}}[[@LINE-3]] + } + for (int i = 0; i < N; ++i) { + ptr[i] += ref2[i]; + } + }) + .wait(); + + sycl::free(ptr, q); + + return 0; +} diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/local/local-overflow-4.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/local/local-overflow-4.cpp new file mode 100644 index 0000000000000..bf6915a605fc7 --- /dev/null +++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/local/local-overflow-4.cpp @@ -0,0 +1,51 @@ +// RUN: %{build} %device_sanitizer_flags -g -o %t.out +// RUN: %{run} not %t.out 2>&1 | FileCheck %s +#include +#include + +constexpr std::size_t N = 1024ULL; +constexpr std::size_t group_size = 4; + +int k_func(sycl::nd_item<2> item) { + auto mem1 = sycl::ext::oneapi::group_local_memory(item.get_group(), 1, + 2, 3, 4, 5); + auto mem2 = sycl::ext::oneapi::group_local_memory(item.get_group(), 1, + 2, 3, 4, 5); + auto mem3 = sycl::ext::oneapi::group_local_memory(item.get_group(), 1, + 2, 3, 4, 5); + auto ref1 = *mem1, ref2 = *mem2, ref3 = *mem3; + + int sum = 0; + sum += ref1[0]; + sum += ref2[3]; + for (int i = 0; i < 10; ++i) { + sum += ref3[i]; + // CHECK: ERROR: DeviceSanitizer: out-of-bounds-access on Local Memory + // CHECK: {{READ of size 1 at kernel <.*MyKernel> LID\(0, 1, 0\) GID\(.*, .*, 0\)}} + // CHECK: {{ #0 .* .*local-overflow-4.cpp:}}[[@LINE-3]] + } + + return sum; +} + +int main() { + sycl::queue Q; + int *ptr = sycl::malloc_shared(1, Q); + + Q.submit([&](sycl::handler &h) { + h.parallel_for( + sycl::nd_range<2>({N, N}, {group_size, group_size}), + [=](sycl::nd_item<2> item) { + auto wgid = item.get_group_linear_id(); + // *ptr = wgid; + // sycl::ext::oneapi::experimental::printf("wgid: %u\n", + // wgid); + auto i = item.get_local_id(0); + if (i % group_size == 1) { + *ptr += k_func(item); + } + }); + }).wait(); + + return 0; +} diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/local/local-overflow-5.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/local/local-overflow-5.cpp new file mode 100644 index 0000000000000..06e870e7764b1 --- /dev/null +++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/local/local-overflow-5.cpp @@ -0,0 +1,27 @@ +// UNSUPPORTED: true +// RUN: %{build} %device_sanitizer_flags -g -o %t.out +// RUN: %{run} not %t.out 2>&1 | FileCheck %s +#include + +using namespace syclex = sycl::ext::oneapi::experimental; +constexpr std::size_t N = 8ULL; +constexpr std::size_t group_size = 8; + +// optional: static const +syclex::work_group_local dynamic_program_scope_array; + +int main() { + sycl::queue Q; + + Q.submit([&](sycl::handler &h) { + Q.parallel_for( + sycl::nd_range<1>{N, group_size}, + syclex::properties{syclex::work_group_local_size(M * sizeof(int))}, + [=](sycl::nd_item<1> it) { + (*dynamic_program_scope_array)[it.get_local_id(0)] = 0; + }); + }); + + Q.wait(); + return 0; +} diff --git a/sycl/test-e2e/lit.cfg.py b/sycl/test-e2e/lit.cfg.py index 869dce39a5a13..a618c7198ed2b 100644 --- a/sycl/test-e2e/lit.cfg.py +++ b/sycl/test-e2e/lit.cfg.py @@ -68,6 +68,10 @@ del llvm_config.config.environment[name] # Propagate some variables from the host environment. +<<<<<<< HEAD +llvm_config.with_system_environment(['PATH', 'OCL_ICD_FILENAMES', + 'CL_CONFIG_DEVICES', 'SYCL_DEVICE_ALLOWLIST', 'SYCL_CONFIG_FILE_NAME', 'ASAN_OPTIONS']) +======= llvm_config.with_system_environment( [ "PATH", @@ -77,6 +81,7 @@ "SYCL_CONFIG_FILE_NAME", ] ) +>>>>>>> sycl llvm_config.with_environment("PATH", config.lit_tools_dir, append_path=True) @@ -693,6 +698,10 @@ except ImportError: pass +<<<<<<< HEAD +config.substitutions.append( ('%device_sanitizer_flags', "-fsanitize=address -fsanitize-target=device") ) +======= config.substitutions.append( ("%device_sanitizer_flags", "-Xsycl-target-frontend -fsanitize=address") ) +>>>>>>> sycl