diff --git a/libdevice/sanitizer/asan_rtl.cpp b/libdevice/sanitizer/asan_rtl.cpp index c4319c2c6f128..11e2f9e5ec7d7 100644 --- a/libdevice/sanitizer/asan_rtl.cpp +++ b/libdevice/sanitizer/asan_rtl.cpp @@ -935,4 +935,93 @@ __asan_set_private_base(__SYCL_PRIVATE__ void *ptr) { SubGroupBarrier(); } +// Intercept string functions +#define ASAN_MEMSET(as) \ + DEVICE_EXTERN_C_NOINLINE __attribute__((address_space(as))) void * \ + __asan_memset_p##as(__attribute__((address_space(as))) char *ptr, int val, \ + size_t size, const char __SYCL_CONSTANT__ *file, \ + uint32_t line, const char __SYCL_CONSTANT__ *func) { \ + if (__AsanLaunchInfo) { \ + DebugInfo debug{(uptr)ptr, as, size, true, file, func, line}; \ + if (auto poisoned_addr = \ + IsRegionPoisoned((uptr)ptr, as, size, &debug)) { \ + ReportAccessError(poisoned_addr, as, false, &debug); \ + } \ + } \ + return Memset(ptr, val, size); \ + } + +ASAN_MEMSET(0) +ASAN_MEMSET(1) +ASAN_MEMSET(3) +ASAN_MEMSET(4) + +#define ASAN_MEMCPY_BASE(dst_as, src_as) \ + DEVICE_EXTERN_C_NOINLINE __attribute__((address_space(dst_as))) void * \ + __asan_memcpy_p##dst_as##_p##src_as( \ + __attribute__((address_space(dst_as))) char *dst, \ + __attribute__((address_space(src_as))) char *src, size_t size, \ + const char __SYCL_CONSTANT__ *file, uint32_t line, \ + const char __SYCL_CONSTANT__ *func) { \ + if (__AsanLaunchInfo) { \ + DebugInfo debug_dst{(uptr)dst, dst_as, size, true, file, func, line}; \ + if (auto poisoned_addr = \ + IsRegionPoisoned((uptr)dst, dst_as, size, &debug_dst)) { \ + ReportAccessError(poisoned_addr, dst_as, false, &debug_dst); \ + } \ + DebugInfo debug_src{(uptr)src, src_as, size, false, file, func, line}; \ + if (auto poisoned_addr = \ + IsRegionPoisoned((uptr)src, src_as, size, &debug_src)) { \ + ReportAccessError(poisoned_addr, src_as, false, &debug_src); \ + } \ + } \ + return Memcpy(dst, src, size); \ + } + +#define ASAN_MEMCPY(dst_as) \ + ASAN_MEMCPY_BASE(dst_as, 0) \ + ASAN_MEMCPY_BASE(dst_as, 1) \ + ASAN_MEMCPY_BASE(dst_as, 2) \ + ASAN_MEMCPY_BASE(dst_as, 3) \ + ASAN_MEMCPY_BASE(dst_as, 4) + +ASAN_MEMCPY(0) +ASAN_MEMCPY(1) +ASAN_MEMCPY(3) +ASAN_MEMCPY(4) + +#define ASAN_MEMMOVE_BASE(dst_as, src_as) \ + DEVICE_EXTERN_C_NOINLINE __attribute__((address_space(dst_as))) void * \ + __asan_memmove_p##dst_as##_p##src_as( \ + __attribute__((address_space(dst_as))) char *dst, \ + __attribute__((address_space(src_as))) char *src, size_t size, \ + const char __SYCL_CONSTANT__ *file, uint32_t line, \ + const char __SYCL_CONSTANT__ *func) { \ + if (__AsanLaunchInfo) { \ + DebugInfo debug_dst{(uptr)dst, dst_as, size, true, file, func, line}; \ + if (auto poisoned_addr = \ + IsRegionPoisoned((uptr)dst, dst_as, size, &debug_dst)) { \ + ReportAccessError(poisoned_addr, dst_as, false, &debug_dst); \ + } \ + DebugInfo debug_src{(uptr)src, src_as, size, false, file, func, line}; \ + if (auto poisoned_addr = \ + IsRegionPoisoned((uptr)src, src_as, size, &debug_src)) { \ + ReportAccessError(poisoned_addr, src_as, false, &debug_src); \ + } \ + } \ + return Memmove(dst, src, size); \ + } + +#define ASAN_MEMMOVE(dst_as) \ + ASAN_MEMMOVE_BASE(dst_as, 0) \ + ASAN_MEMMOVE_BASE(dst_as, 1) \ + ASAN_MEMMOVE_BASE(dst_as, 2) \ + ASAN_MEMMOVE_BASE(dst_as, 3) \ + ASAN_MEMMOVE_BASE(dst_as, 4) + +ASAN_MEMMOVE(0) +ASAN_MEMMOVE(1) +ASAN_MEMMOVE(3) +ASAN_MEMMOVE(4) + #endif // __SPIR__ || __SPIRV__ diff --git a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp index 7b9aa86dc9bbc..456f1e3d3133b 100644 --- a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp +++ b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp @@ -903,7 +903,7 @@ struct AddressSanitizer { bool instrumentSyclDynamicLocalMemory(Function &F); void instrumentInitAsanLaunchInfo(Function &F, const TargetLibraryInfo *TLI); - void AppendDebugInfoToArgs(Instruction *InsertBefore, Value *Addr, + void AppendDebugInfoToArgs(Instruction *InsertBefore, SmallVectorImpl &Args); private: @@ -964,6 +964,9 @@ struct AddressSanitizer { FunctionCallee AsanMemoryAccessCallbackSizedAS[2][2][kNumberOfAddressSpace]; FunctionCallee AsanMemmove, AsanMemcpy, AsanMemset; + FunctionCallee AsanMemcpyAS[kNumberOfAddressSpace][kNumberOfAddressSpace], + AsanMemmoveAS[kNumberOfAddressSpace][kNumberOfAddressSpace], + AsanMemsetAS[kNumberOfAddressSpace]; Value *LocalDynamicShadow = nullptr; const StackSafetyGlobalInfo *SSGI; DenseMap ProcessedAllocas; @@ -1355,7 +1358,7 @@ class AddressSanitizerOnSpirv { void initializeCallbacks() { IRBuilder<> IRB(*C); - // __msan_set_private_base( + // __asan_set_private_base( // as(0) void * ptr // ) AsanSetPrivateBaseFunc = @@ -1803,7 +1806,6 @@ static bool isUnsupportedSPIRAccess(Value *Addr, Instruction *Inst) { } void AddressSanitizer::AppendDebugInfoToArgs(Instruction *InsertBefore, - Value *Addr, SmallVectorImpl &Args) { auto *M = InsertBefore->getModule(); auto &C = InsertBefore->getContext(); @@ -1933,17 +1935,48 @@ void AddressSanitizer::instrumentMemIntrinsic(MemIntrinsic *MI, RuntimeCallInserter &RTCI) { InstrumentationIRBuilder IRB(MI); if (isa(MI)) { - RTCI.createRuntimeCall( - IRB, isa(MI) ? AsanMemmove : AsanMemcpy, - {IRB.CreateAddrSpaceCast(MI->getOperand(0), PtrTy), - IRB.CreateAddrSpaceCast(MI->getOperand(1), PtrTy), - IRB.CreateIntCast(MI->getOperand(2), IntptrTy, false)}); + if (TargetTriple.isSPIROrSPIRV()) { + unsigned int DstAS = + cast(MI->getOperand(0)->getType()->getScalarType()) + ->getPointerAddressSpace(); + unsigned int SrcAS = + cast(MI->getOperand(1)->getType()->getScalarType()) + ->getPointerAddressSpace(); + SmallVector Args; + Args.push_back(MI->getOperand(0)); + Args.push_back(MI->getOperand(1)); + Args.push_back(IRB.CreateIntCast(MI->getOperand(2), IntptrTy, false)); + AppendDebugInfoToArgs(MI, Args); + RTCI.createRuntimeCall(IRB, + isa(MI) ? AsanMemmoveAS[DstAS][SrcAS] + : AsanMemcpyAS[DstAS][SrcAS], + Args); + } else { + RTCI.createRuntimeCall( + IRB, isa(MI) ? AsanMemmove : AsanMemcpy, + {IRB.CreateAddrSpaceCast(MI->getOperand(0), PtrTy), + IRB.CreateAddrSpaceCast(MI->getOperand(1), PtrTy), + IRB.CreateIntCast(MI->getOperand(2), IntptrTy, false)}); + } } else if (isa(MI)) { - RTCI.createRuntimeCall( - IRB, AsanMemset, - {IRB.CreateAddrSpaceCast(MI->getOperand(0), PtrTy), - IRB.CreateIntCast(MI->getOperand(1), IRB.getInt32Ty(), false), - IRB.CreateIntCast(MI->getOperand(2), IntptrTy, false)}); + if (TargetTriple.isSPIROrSPIRV()) { + unsigned int AS = + cast(MI->getOperand(0)->getType()->getScalarType()) + ->getPointerAddressSpace(); + SmallVector Args; + Args.push_back(MI->getOperand(0)); + Args.push_back( + IRB.CreateIntCast(MI->getOperand(1), IRB.getInt32Ty(), false)); + Args.push_back(IRB.CreateIntCast(MI->getOperand(2), IntptrTy, false)); + AppendDebugInfoToArgs(MI, Args); + RTCI.createRuntimeCall(IRB, AsanMemsetAS[AS], Args); + } else { + RTCI.createRuntimeCall( + IRB, AsanMemset, + {IRB.CreateAddrSpaceCast(MI->getOperand(0), PtrTy), + IRB.CreateIntCast(MI->getOperand(1), IRB.getInt32Ty(), false), + IRB.CreateIntCast(MI->getOperand(2), IntptrTy, false)}); + } } MI->eraseFromParent(); } @@ -2496,7 +2529,7 @@ void AddressSanitizer::instrumentAddress(Instruction *OrigIns, auto AS = cast(Addr->getType()->getScalarType()) ->getPointerAddressSpace(); Args.push_back(AddrLong); - AppendDebugInfoToArgs(InsertBefore, Addr, Args); + AppendDebugInfoToArgs(InsertBefore, Args); RTCI.createRuntimeCall( IRB, AsanMemoryAccessCallbackAS[IsWrite][0][AccessSizeIndex][AS], Args); @@ -2582,7 +2615,7 @@ void AddressSanitizer::instrumentUnusualSizeOrAlignment( ->getPointerAddressSpace(); Args.push_back(AddrLong); Args.push_back(Size); - AppendDebugInfoToArgs(InsertBefore, Addr, Args); + AppendDebugInfoToArgs(InsertBefore, Args); RTCI.createRuntimeCall( IRB, AsanMemoryAccessCallbackSizedAS[IsWrite][0][AS], Args); } else { @@ -3774,6 +3807,46 @@ void AddressSanitizer::initializeCallbacks(const TargetLibraryInfo *TLI) { } } + if (TargetTriple.isSPIROrSPIRV()) { + auto *Int8PtrTy = PointerType::get(*C, kSpirOffloadConstantAS); + for (size_t FirstArgAS = 0; FirstArgAS < kNumberOfAddressSpace; + FirstArgAS++) { + PointerType *FirstArgPtrTy = PointerType::get(*C, FirstArgAS); + // __asan_memset_pX ( + // ... + // char* file, + // unsigned int line, + // char* func + // ) + AsanMemsetAS[FirstArgAS] = M.getOrInsertFunction( + ClMemoryAccessCallbackPrefix + "memset_p" + itostr(FirstArgAS), + TLI->getAttrList(C, {1}, /*Signed=*/false), FirstArgPtrTy, + FirstArgPtrTy, IRB.getInt32Ty(), IntptrTy, Int8PtrTy, + IRB.getInt32Ty(), Int8PtrTy); + + for (size_t SecondArgAS = 0; SecondArgAS < kNumberOfAddressSpace; + SecondArgAS++) { + PointerType *SecondArgPtrTy = PointerType::get(*C, SecondArgAS); + // __asan_mem[cpy|move]_pX_pX ( + // ... + // char* file, + // unsigned int line, + // char* func + // ) + AsanMemcpyAS[FirstArgAS][SecondArgAS] = M.getOrInsertFunction( + ClMemoryAccessCallbackPrefix + "memcpy_p" + itostr(FirstArgAS) + + "_p" + itostr(SecondArgAS), + FirstArgPtrTy, FirstArgPtrTy, SecondArgPtrTy, IntptrTy, Int8PtrTy, + IRB.getInt32Ty(), Int8PtrTy); + AsanMemmoveAS[FirstArgAS][SecondArgAS] = M.getOrInsertFunction( + ClMemoryAccessCallbackPrefix + "memmove_p" + itostr(FirstArgAS) + + "_p" + itostr(SecondArgAS), + FirstArgPtrTy, FirstArgPtrTy, SecondArgPtrTy, IntptrTy, Int8PtrTy, + IRB.getInt32Ty(), Int8PtrTy); + } + } + } + const std::string MemIntrinCallbackPrefix = (CompileKernel && !ClKasanMemIntrinCallbackPrefix) ? std::string("") @@ -4048,12 +4121,10 @@ bool AddressSanitizer::instrumentFunction(Function &F, F.getDataLayout(), RTCI); FunctionModified = true; } - if (!TargetTriple.isSPIROrSPIRV()) { - for (auto *Inst : IntrinToInstrument) { - if (!suppressInstrumentationSiteForDebug(NumInstrumented)) - instrumentMemIntrinsic(Inst, RTCI); - FunctionModified = true; - } + for (auto *Inst : IntrinToInstrument) { + if (!suppressInstrumentationSiteForDebug(NumInstrumented)) + instrumentMemIntrinsic(Inst, RTCI); + FunctionModified = true; } FunctionStackPoisoner FSP(F, *this, RTCI); diff --git a/llvm/test/Instrumentation/AddressSanitizer/SPIRV/string_func.ll b/llvm/test/Instrumentation/AddressSanitizer/SPIRV/string_func.ll new file mode 100644 index 0000000000000..52edce251699e --- /dev/null +++ b/llvm/test/Instrumentation/AddressSanitizer/SPIRV/string_func.ll @@ -0,0 +1,162 @@ +; RUN: opt < %s -passes=asan -asan-instrumentation-with-call-threshold=0 -asan-stack=0 -asan-globals=0 -asan-constructor-kind=none -asan-use-after-return=never -S | FileCheck %s + +target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1" +target triple = "spir64-unknown-unknown" + +; Function Attrs: sanitize_address +define weak_odr dso_local spir_kernel void @test_memset_as0() #0 { +; CHECK-LABEL: define weak_odr dso_local spir_kernel void @test_memset_as0 +entry: + %p.i = alloca [4 x i8], align 4 + call void @llvm.memset.p0.i64(ptr %p.i, i8 1, i64 5, i1 false) + ; CHECK: [[MEMSET_PTR:%[0-9]+]] = getelementptr i8, ptr %MyAlloca + ; CHECK: call ptr @__asan_memset_p0(ptr [[MEMSET_PTR]], i32 1, i64 5 + ret void +} + +; Function Attrs: sanitize_address +define weak_odr dso_local spir_kernel void @test_memset_as1(ptr addrspace(1) %_arg_ptr) #0 { +; CHECK-LABEL: define weak_odr dso_local spir_kernel void @test_memset_as1 +entry: + call void @llvm.memset.p1.i64(ptr addrspace(1) %_arg_ptr, i8 1, i64 13, i1 false) + ; CHECK: call ptr addrspace(1) @__asan_memset_p1(ptr addrspace(1) %_arg_ptr, i32 1, i64 13 + ret void +} + +; Function Attrs: sanitize_address +define weak_odr dso_local spir_kernel void @test_memset_as3(ptr addrspace(3) %_arg_ptr) #0 { +; CHECK-LABEL: define weak_odr dso_local spir_kernel void @test_memset_as3 +entry: + call void @llvm.memset.p3.i64(ptr addrspace(3) %_arg_ptr, i8 1, i64 13, i1 false) + ; CHECK: call ptr addrspace(3) @__asan_memset_p3(ptr addrspace(3) %_arg_ptr, i32 1, i64 13 + ret void +} + +; Function Attrs: sanitize_address +define weak_odr dso_local spir_kernel void @test_memset_as4(ptr addrspace(4) %_arg_ptr) #0 { +; CHECK-LABEL: define weak_odr dso_local spir_kernel void @test_memset_as4 +entry: + call void @llvm.memset.p4.i64(ptr addrspace(4) %_arg_ptr, i8 1, i64 13, i1 false) + ; CHECK: call ptr addrspace(4) @__asan_memset_p4(ptr addrspace(4) %_arg_ptr, i32 1, i64 13 + ret void +} + +; Function Attrs: sanitize_address +define weak_odr dso_local spir_kernel void @test_memcpy_as0() #0 { +; CHECK-LABEL: define weak_odr dso_local spir_kernel void @test_memcpy_as0 +entry: + %dst = alloca [4 x i8], align 4 + %src = alloca [4 x i8], align 4 + call void @llvm.memcpy.p0.p0.i64(ptr %dst, ptr %src, i64 5, i1 false) + ; CHECK: [[MEMCPY_DST:%[0-9]+]] = getelementptr i8, ptr %MyAlloca + ; CHECK-NEXT: [[MEMCPY_SRC:%[0-9]+]] = getelementptr i8, ptr %MyAlloca + ; CHECK: call ptr @__asan_memcpy_p0_p0(ptr [[MEMCPY_DST]], ptr [[MEMCPY_SRC]], i64 5 + ret void +} + +; Function Attrs: sanitize_address +define weak_odr dso_local spir_kernel void @test_memcpy_as1(ptr addrspace(1) %_arg_dst, ptr addrspace(1) %_arg_src) #0 { +; CHECK-LABEL: define weak_odr dso_local spir_kernel void @test_memcpy_as1 +entry: + call void @llvm.memcpy.p1.p1.i64(ptr addrspace(1) %_arg_dst, ptr addrspace(1) %_arg_src, i64 12, i1 false) + ; CHECK: call ptr addrspace(1) @__asan_memcpy_p1_p1(ptr addrspace(1) %_arg_dst, ptr addrspace(1) %_arg_src, i64 12 + ret void +} + +; Function Attrs: sanitize_address +define weak_odr dso_local spir_kernel void @test_memcpy_as3(ptr addrspace(3) %_arg_dst, ptr addrspace(3) %_arg_src) #0 { +; CHECK-LABEL: define weak_odr dso_local spir_kernel void @test_memcpy_as3 +entry: + call void @llvm.memcpy.p3.p3.i64(ptr addrspace(3) %_arg_dst, ptr addrspace(3) %_arg_src, i64 12, i1 false) + ; CHECK: call ptr addrspace(3) @__asan_memcpy_p3_p3(ptr addrspace(3) %_arg_dst, ptr addrspace(3) %_arg_src, i64 12 + ret void +} + +; Function Attrs: sanitize_address +define weak_odr dso_local spir_kernel void @test_memcpy_as4(ptr addrspace(4) %_arg_dst, ptr addrspace(4) %_arg_src) #0 { +; CHECK-LABEL: define weak_odr dso_local spir_kernel void @test_memcpy_as4 +entry: + call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) %_arg_dst, ptr addrspace(4) %_arg_src, i64 12, i1 false) + ; CHECK: call ptr addrspace(4) @__asan_memcpy_p4_p4(ptr addrspace(4) %_arg_dst, ptr addrspace(4) %_arg_src, i64 12 + ret void +} + +; Function Attrs: sanitize_address +define weak_odr dso_local spir_kernel void @test_memmove_as0() #0 { +; CHECK-LABEL: define weak_odr dso_local spir_kernel void @test_memmove_as0 +entry: + %dst = alloca [4 x i8], align 4 + %src = alloca [4 x i8], align 4 + call void @llvm.memmove.p0.p0.i64(ptr %dst, ptr %src, i64 5, i1 false) + ; CHECK: [[MEMMOVE_DST:%[0-9]+]] = getelementptr i8, ptr %MyAlloca + ; CHECK-NEXT: [[MEMMOVE_SRC:%[0-9]+]] = getelementptr i8, ptr %MyAlloca + ; CHECK: call ptr @__asan_memmove_p0_p0(ptr [[MEMMOVE_DST]], ptr [[MEMMOVE_SRC]], i64 5 + ret void +} + +; Function Attrs: sanitize_address +define weak_odr dso_local spir_kernel void @test_memmove_as1(ptr addrspace(1) %_arg_dst, ptr addrspace(1) %_arg_src) #0 { +; CHECK-LABEL: define weak_odr dso_local spir_kernel void @test_memmove_as1 +entry: + call void @llvm.memmove.p1.p1.i64(ptr addrspace(1) %_arg_dst, ptr addrspace(1) %_arg_src, i64 12, i1 false) + ; CHECK: call ptr addrspace(1) @__asan_memmove_p1_p1(ptr addrspace(1) %_arg_dst, ptr addrspace(1) %_arg_src, i64 12 + ret void +} + +; Function Attrs: sanitize_address +define weak_odr dso_local spir_kernel void @test_memmove_as3(ptr addrspace(3) %_arg_dst, ptr addrspace(3) %_arg_src) #0 { +; CHECK-LABEL: define weak_odr dso_local spir_kernel void @test_memmove_as3 +entry: + call void @llvm.memmove.p3.p3.i64(ptr addrspace(3) %_arg_dst, ptr addrspace(3) %_arg_src, i64 12, i1 false) + ; CHECK: call ptr addrspace(3) @__asan_memmove_p3_p3(ptr addrspace(3) %_arg_dst, ptr addrspace(3) %_arg_src, i64 12 + ret void +} + +; Function Attrs: sanitize_address +define weak_odr dso_local spir_kernel void @test_memmove_as4(ptr addrspace(4) %_arg_dst, ptr addrspace(4) %_arg_src) #0 { +; CHECK-LABEL: define weak_odr dso_local spir_kernel void @test_memmove_as4 +entry: + call void @llvm.memmove.p4.p4.i64(ptr addrspace(4) %_arg_dst, ptr addrspace(4) %_arg_src, i64 12, i1 false) + ; CHECK: call ptr addrspace(4) @__asan_memmove_p4_p4(ptr addrspace(4) %_arg_dst, ptr addrspace(4) %_arg_src, i64 12 + ret void +} + +; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: write) +declare void @llvm.memset.p0.i64(ptr writeonly captures(none), i8, i64, i1 immarg) #1 + +; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: write) +declare void @llvm.memset.p1.i64(ptr addrspace(1) writeonly captures(none), i8, i64, i1 immarg) #1 + +; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: write) +declare void @llvm.memset.p3.i64(ptr addrspace(3) writeonly captures(none), i8, i64, i1 immarg) #1 + +; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: write) +declare void @llvm.memset.p4.i64(ptr addrspace(4) writeonly captures(none), i8, i64, i1 immarg) #1 + +; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: readwrite) +declare void @llvm.memcpy.p0.p0.i64(ptr noalias writeonly captures(none), ptr noalias readonly captures(none), i64, i1 immarg) #2 + +; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: readwrite) +declare void @llvm.memcpy.p1.p1.i64(ptr addrspace(1) noalias writeonly captures(none), ptr addrspace(1) noalias readonly captures(none), i64, i1 immarg) #2 + +; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: readwrite) +declare void @llvm.memcpy.p3.p3.i64(ptr addrspace(3) noalias writeonly captures(none), ptr addrspace(3) noalias readonly captures(none), i64, i1 immarg) #2 + +; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: readwrite) +declare void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) noalias writeonly captures(none), ptr addrspace(4) noalias readonly captures(none), i64, i1 immarg) #2 + +; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: readwrite) +declare void @llvm.memmove.p0.p0.i64(ptr writeonly captures(none), ptr readonly captures(none), i64, i1 immarg) #2 + +; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: readwrite) +declare void @llvm.memmove.p1.p1.i64(ptr addrspace(1) writeonly captures(none), ptr addrspace(1) readonly captures(none), i64, i1 immarg) #2 + +; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: readwrite) +declare void @llvm.memmove.p3.p3.i64(ptr addrspace(3) writeonly captures(none), ptr addrspace(3) readonly captures(none), i64, i1 immarg) #2 + +; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: readwrite) +declare void @llvm.memmove.p4.p4.i64(ptr addrspace(4) writeonly captures(none), ptr addrspace(4) readonly captures(none), i64, i1 immarg) #2 + +attributes #0 = { sanitize_address } +attributes #1 = { nocallback nofree nounwind willreturn memory(argmem: write) } +attributes #2 = { nocallback nofree nounwind willreturn memory(argmem: readwrite) } diff --git a/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/string_func.cpp b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/string_func.cpp new file mode 100644 index 0000000000000..c45ee54ee1c9c --- /dev/null +++ b/sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/string_func.cpp @@ -0,0 +1,104 @@ +// REQUIRES: linux, cpu || (gpu && level_zero) +// RUN: %{build} %device_asan_flags -O0 -g -o %t.out +// RUN: %{run} not %t.out memset 2>&1 | FileCheck --check-prefixes CHECK-MEMSET %s +// RUN: %{run} not %t.out memcpy src 2>&1 | FileCheck --check-prefixes CHECK-MEMCPY-SRC %s +// RUN: %{run} not %t.out memcpy dst 2>&1 | FileCheck --check-prefixes CHECK-MEMCPY-DST %s +// RUN: %{run} not %t.out memmove src 2>&1 | FileCheck --check-prefixes CHECK-MEMMOVE-SRC %s +// RUN: %{run} not %t.out memmove dst 2>&1 | FileCheck --check-prefixes CHECK-MEMMOVE-DST %s +#include + +#include + +constexpr size_t N = 12; + +void test_memset(sycl::queue &Q) { + auto *ptr = sycl::malloc_device(N, Q); + + Q.submit([&](sycl::handler &h) { + h.parallel_for( + sycl::nd_range<1>(1, 1), + [=](sycl::nd_item<1>) { memset(ptr, 1, N + 1); }); + }); + Q.wait(); + // CHECK-MEMSET: ERROR: DeviceSanitizer: out-of-bounds-access on Device USM + // CHECK-MEMSET: {{WRITE of size 13 at kernel <.*test_memset>}} + // CHECK-MEMSET: {{ #0 .* .*string_func.cpp:}}[[@LINE-5]] + + sycl::free(ptr, Q); +} + +void test_memcpy(sycl::queue &Q, bool is_src_oob) { + auto *dst = sycl::malloc_device(N, Q); + auto *src = sycl::malloc_device(N, Q); + + if (is_src_oob) + Q.submit([&](sycl::handler &h) { + h.parallel_for( + sycl::nd_range<1>(1, 1), + [=](sycl::nd_item<1>) { memcpy(dst, src + 1, N); }); + }); + else + Q.submit([&](sycl::handler &h) { + h.parallel_for( + sycl::nd_range<1>(1, 1), + [=](sycl::nd_item<1>) { memcpy(dst + 1, src, N); }); + }); + Q.wait(); + // CHECK-MEMCPY-SRC: ERROR: DeviceSanitizer: out-of-bounds-access on Device USM + // CHECK-MEMCPY-SRC: {{READ of size 12 at kernel <.*test_memcpy_src>}} + // CHECK-MEMCPY-SRC: {{ #0 .* .*string_func.cpp:}}[[@LINE-11]] + // CHECK-MEMCPY-DST: ERROR: DeviceSanitizer: out-of-bounds-access on Device USM + // CHECK-MEMCPY-DST: {{WRITE of size 12 at kernel <.*test_memcpy_dst>}} + // CHECK-MEMCPY-DST: {{ #0 .* .*string_func.cpp:}}[[@LINE-8]] + + sycl::free(dst, Q); + sycl::free(src, Q); +} + +void test_memmove(sycl::queue &Q, bool is_src_oob) { + auto *dst = sycl::malloc_device(N, Q); + auto *src = sycl::malloc_device(N, Q); + + if (is_src_oob) + Q.submit([&](sycl::handler &h) { + h.parallel_for( + sycl::nd_range<1>(1, 1), + [=](sycl::nd_item<1>) { memmove(dst, src + 1, N); }); + }); + else + Q.submit([&](sycl::handler &h) { + h.parallel_for( + sycl::nd_range<1>(1, 1), + [=](sycl::nd_item<1>) { memmove(dst + 1, src, N); }); + }); + Q.wait(); + // CHECK-MEMMOVE-SRC: ERROR: DeviceSanitizer: out-of-bounds-access on Device USM + // CHECK-MEMMOVE-SRC: {{READ of size 12 at kernel <.*test_memmove_src>}} + // CHECK-MEMMOVE-SRC: {{ #0 .* .*string_func.cpp:}}[[@LINE-11]] + // CHECK-MEMMOVE-DST: ERROR: DeviceSanitizer: out-of-bounds-access on Device USM + // CHECK-MEMMOVE-DST: {{WRITE of size 12 at kernel <.*test_memmove_dst>}} + // CHECK-MEMMOVE-DST: {{ #0 .* .*string_func.cpp:}}[[@LINE-8]] + + sycl::free(dst, Q); + sycl::free(src, Q); +} + +int main(int argc, char **argv) { + assert(argc > 1 && "test is not specified"); + sycl::queue Q; + + if (!strcmp(argv[1], "memset")) { + test_memset(Q); + } else if (!strcmp(argv[1], "memcpy")) { + if (!strcmp(argv[2], "src")) + test_memcpy(Q, true); + else + test_memcpy(Q, false); + } else if (!strcmp(argv[1], "memmove")) { + if (!strcmp(argv[2], "src")) + test_memmove(Q, true); + else + test_memmove(Q, false); + } + return 0; +}