add func test

Author: AllanZyne

llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp

@@ -1666,6 +1666,9 @@ static bool isUnsupportedDeviceGlobal(GlobalVariable *G) {
   if (G->getName().starts_with("__Asan"))
     return true;
 
+  if (G->getAddressSpace() == kSpirOffloadLocalAS)
+    return true;
+
   Attribute Attr = G->getAttribute("sycl-device-image-scope");
   return (!Attr.isStringAttribute() || Attr.getValueAsString() == "false");
 }
@@ -3062,7 +3065,8 @@ void ModuleAddressSanitizer::instrumentSyclStaticLocalMemory(IRBuilder<> &IRB) {
       // Get root spir_kernel of spir_func
       initializeKernelCallerMap(F);
       for (Function *Kernel : FuncToKernelCallerMap[F])
-        Instrument(G, Kernel);
+        if (!InstrumentedFunc.contains(Kernel))
+          Instrument(G, Kernel);
     }
   }
 }
@@ -3945,9 +3949,9 @@ bool AddressSanitizer::instrumentFunction(Function &F,
   if (ChangedStack || !NoReturnCalls.empty())
     FunctionModified = true;
 
-  // We need to instrument dynamic/static local arguments after stack poisoner
+  // We need to instrument dynamic local arguments after stack poisoner
   if (TargetTriple.isSPIROrSPIRV()) {
-    if (F.getCallingConv() == CallingConv::SPIR_KERNEL) {
+    if (ClSpirOffloadLocals && F.getCallingConv() == CallingConv::SPIR_KERNEL) {
       FunctionModified |= instrumentSyclDynamicLocalMemory(F, FSP.RetVec);
     }
   }

llvm/test/Instrumentation/AddressSanitizer/SPIRV/instrument_local_addess_space.ll

@@ -7,20 +7,20 @@ target triple = "spir64-unknown-unknown"
 %"class.sycl::_V1::detail::array" = type { [1 x i64] }
 %"class.sycl::_V1::id" = type { %"class.sycl::_V1::detail::array" }
 
-declare dso_local spir_func ptr addrspace(3) @__sycl_allocateLocalMemory(i64 noundef %0, i64 noundef %1) local_unnamed_addr #1
+@WGLocalMem = internal addrspace(3) global [64 x i8] poison, align 4
 
-define spir_kernel void @kernel_static_local() #0 {
+define spir_kernel void @kernel_static_local() sanitize_address {
 ; CHECK-LABEL: define spir_kernel void @kernel_static_local
 entry:
-  %1 = tail call spir_func ptr addrspace(3) @__sycl_allocateLocalMemory(i64 noundef 16, i64 noundef 4)
-  ; CHECK:      [[T0:%.*]] = call ptr addrspace(3) @__sycl_allocateLocalMemory(i64 32, i64 8)
-  ; CHECK-NEXT: [[T1:%.*]] = ptrtoint ptr addrspace(3) [[T0]] to i64
-  ; CHECK-NEXT: call void @__asan_set_shadow_static_local(i64 [[T1]], i64 16, i64 32)
-  ; CHECK-NEXT: call void @__asan_unpoison_shadow_static_local(i64 %1, i64 16, i64 32)
+  store i32 0, ptr addrspace(3) @WGLocalMem
+  ; CHECK: store ptr addrspace(1) %__asan_launch, ptr addrspace(3) @__AsanLaunchInfo, align 8
+  ; CHECK-NEXT: call void @__asan_set_shadow_static_local(i64 ptrtoint (ptr addrspace(3) @WGLocalMem to i64), i64 64, i64 96)
+  ; CHECK-NEXT: store i32 0, ptr addrspace(3) @WGLocalMem, align 4
+  ; CHECK-NEXT: call void @__asan_unpoison_shadow_static_local(i64 ptrtoint (ptr addrspace(3) @WGLocalMem to i64), i64 64, i64 96)
   ret void
 }
 
-define spir_kernel void @kernel_dynamic_local(ptr addrspace(3) noundef align 4 %_arg_acc, ptr noundef byval(%"class.sycl::_V1::range") align 8 %_arg_acc1, ptr noundef byval(%"class.sycl::_V1::range") align 8 %_arg_acc2, ptr noundef byval(%"class.sycl::_V1::id") align 8 %_arg_acc3) #0 {
+define spir_kernel void @kernel_dynamic_local(ptr addrspace(3) noundef align 4 %_arg_acc, ptr noundef byval(%"class.sycl::_V1::range") align 8 %_arg_acc1, ptr noundef byval(%"class.sycl::_V1::range") align 8 %_arg_acc2, ptr noundef byval(%"class.sycl::_V1::id") align 8 %_arg_acc3) sanitize_address {
 ; CHECK-LABEL: define spir_kernel void @kernel_dynamic_local
 entry:
   ; CHECK:      %local_args = alloca i64, align 8
@@ -32,6 +32,3 @@ entry:
   ; CHECK: call void @__asan_unpoison_shadow_dynamic_local(i64 %2, i32 1)
   ret void
 }
-
-attributes #0 = { sanitize_address }
-attributes #1 = { convergent nounwind }

sycl/test-e2e/AddressSanitizer/out-of-bounds/local/group_local_memory.cpp

@@ -30,6 +30,7 @@ int main() {
               ptr = sycl::ext::oneapi::group_local_memory<int[N]>(
                   item.get_group());
           auto &ref = *ptr;
+          // NOTE: direct access will be optimized out
           data[0] = check(ref, item.get_local_linear_id() * 2 + 4);
         });
   });

sycl/test-e2e/AddressSanitizer/out-of-bounds/local/group_local_memory_func.cpp

@@ -0,0 +1,42 @@
+// REQUIRES: linux, cpu || (gpu && level_zero)
+// RUN: %{build} %device_asan_flags -g -O0 -o %t1.out
+// RUN: %{run} not %t1.out 2>&1 | FileCheck %s
+// RUN: %{build} %device_asan_flags -g -O1 -o %t2.out
+// RUN: %{run} not %t2.out 2>&1 | FileCheck %s
+// RUN: %{build} %device_asan_flags -g -O2 -o %t3.out
+// RUN: %{run} not %t3.out 2>&1 | FileCheck %s
+
+#include <sycl/detail/core.hpp>
+
+#include <sycl/ext/oneapi/group_local_memory.hpp>
+#include <sycl/usm.hpp>
+
+constexpr std::size_t N = 16;
+constexpr std::size_t group_size = 8;
+
+__attribute__((noinline)) int check(int *ref, int index) { return ref[index]; }
+// CHECK: ERROR: DeviceSanitizer: out-of-bounds-access on Local Memory
+// CHECK: READ of size 4 at kernel {{<.*MyKernel>}} LID({{.*}}, 0, 0) GID({{.*}}, 0, 0)
+// CHECK:   #0 {{.*}} {{.*group_local_memory_func.cpp}}:[[@LINE-3]]
+
+__attribute__((noinline)) int test_local(sycl::nd_item<1> &item) {
+  auto local_mem = sycl::ext::oneapi::group_local_memory<int[group_size]>(item.get_group());
+  // NOTE: direct access will be optimized out
+  return check(*local_mem, group_size);
+}
+
+int main() {
+  sycl::queue Q;
+  auto data = sycl::malloc_device<int>(N, Q);
+
+  Q.submit([&](sycl::handler &h) {
+    h.parallel_for<class MyKernel>(
+        sycl::nd_range<1>(N, group_size), [=](sycl::nd_item<1> item) {
+          data[0] = test_local(item);
+        });
+  });
+  Q.wait();
+
+  sycl::free(data, Q);
+  return 0;
+}