fix sycl-host-access

Author: AllanZyne

llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp

@@ -1358,7 +1358,7 @@ static void ExtendSpirKernelArgs(Module &M, FunctionAnalysisManager &FAM,
   AsanSpirKernelMetadata->addAttribute(
       "sycl-device-global-size", std::to_string(DL.getTypeAllocSize(ArrayTy)));
   AsanSpirKernelMetadata->addAttribute("sycl-device-image-scope");
-  AsanSpirKernelMetadata->addAttribute("sycl-host-access", "2");
+  AsanSpirKernelMetadata->addAttribute("sycl-host-access", "0"); // read only
   AsanSpirKernelMetadata->addAttribute("sycl-unique-id",
                                        "_Z20__AsanKernelMetadata");
   AsanSpirKernelMetadata->setDSOLocal(true);

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

@@ -3,6 +3,7 @@
 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"
 target triple = "spir64-unknown-unknown"
 
+; CHECK: @__AsanKernelMetadata = appending dso_local local_unnamed_addr addrspace(1) global
 ; CHECK: @__AsanLaunchInfo = external addrspace(3) global ptr addrspace(1)
 
 define spir_kernel void @sycl_kernel1() #0 {
@@ -22,3 +23,6 @@ entry:
 }
 
 attributes #0 = { sanitize_address }
+;; sycl-device-global-size = 16 * 2
+;; sycl-host-access = 0 read-only
+; CHECK: attributes #{{.*}} = { "sycl-device-global-size"="32" "sycl-device-image-scope" "sycl-host-access"="0" "sycl-unique-id"="_Z20__AsanKernelMetadata" }

llvm/test/tools/sycl-post-link/device-sanitizer/asan.ll

@@ -1,10 +1,13 @@
 ; This test checks that the post-link tool properly generates "asanUsed=1"
-; in [SYCL/misc properties]
+; in prop file, and fixes the attributes and metadata of @__AsanKernelMetadata
 
 ; RUN: sycl-post-link -properties -split=kernel -symbols -S < %s -o %t.table
-; RUN: FileCheck %s -input-file=%t_0.prop
-; CHECK: [SYCL/misc properties]
-; CHECK: asanUsed=1
+
+; RUN: FileCheck %s -input-file=%t_0.prop --check-prefix CHECK-PROP
+; CHECK-PROP: [SYCL/misc properties]
+; CHECK-PROP: asanUsed=1
+
+; RUN: FileCheck %s -input-file=%t_0.ll --check-prefix CHECK-IR
 
 ; ModuleID = 'parallel_for_int.cpp'
 source_filename = "parallel_for_int.cpp"
@@ -15,6 +18,7 @@ $_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8MyKernel = comdat any
 
 @__asan_kernel = internal addrspace(1) constant [55 x i8] c"_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8MyKernel\00"
 @__AsanKernelMetadata = appending dso_local local_unnamed_addr addrspace(1) global [1 x { i64, i64 }] [{ i64, i64 } { i64 ptrtoint (ptr addrspace(1) @__asan_kernel to i64), i64 54 }] #2
+; CHECK-IR: @__AsanKernelMetadata {{.*}} !spirv.Decorations
 @__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
 @__asan_func = internal addrspace(2) constant [106 x i8] c"typeinfo name for main::'lambda'(sycl::_V1::handler&)::operator()(sycl::_V1::handler&) const::MyKernelR_4\00"
 
@@ -48,7 +52,7 @@ declare spir_func void @__itt_offload_wi_finish_wrapper()
 
 attributes #0 = { mustprogress nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: write) }
 attributes #1 = { mustprogress norecurse nounwind sanitize_address uwtable "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="parallel_for_int.cpp" "sycl-optlevel"="2" "uniform-work-group-size"="true" }
-attributes #2 = { "sycl-device-global-size"="16" "sycl-device-image-scope" "sycl-host-access"="2" "sycl-unique-id"="_Z20__AsanKernelMetadata" }
+attributes #2 = { "sycl-device-global-size"="16" "sycl-device-image-scope" "sycl-host-access"="0" "sycl-unique-id"="_Z20__AsanKernelMetadata" }
 
 !llvm.module.flags = !{!0, !1, !2}
 !opencl.spir.version = !{!3}