[SYCL] Add a builtin for SYCL ptr annotation (#10267)

added a "__builtin_intel_sycl_ptr_annotation(ptr, PROPERTY1, VALUE1,
PROPERTY2, VALUE2 ...)" builtin function for FPGA SYCL ptr annotation
using SYCL extension properties.

Updated the "annotated_ptr.hpp" to use
"__builtin_intel_sycl_ptr_annotation" to pass annotation instead of
using "add_sycl_ir_member_annotation". This helps to insert the
llvm.ptr.annotation at the correct position

Update the compileTimeProperties pass to consume "sycl-alignment" and
thus do not propagate this annotation to SPIRV translator

---------------------------------------------------------------------

After using "__builtin_intel_sycl_ptr_annotation" , IR generated are
changed from:
```
%ptr.to.ptr = GEP
%ptr.to.ptr.annotated = llvm.ptr.annotation %ptr.to.ptr
%ptr = load %ptr.to.ptr.annotated
%4 = load %ptr
```
to
```
%ptr.to.ptr = GEP
%ptr = load %ptr.to.ptr
%ptr.annotated = llvm.ptr.annotation %ptr
%4 = load %ptr.annotated
``` 

Thus, SYCLLowerIR pass are expecting different IR forms and thus these
two changes are merged together.

Author: Brox Chen

clang/include/clang/Basic/Builtins.def

@@ -1754,6 +1754,9 @@ BUILTIN(__builtin_ms_va_copy, "vc*&c*&", "n")
 // Arithmetic Fence: to prevent FP reordering and reassociation optimizations
 LANGBUILTIN(__arithmetic_fence, "v.", "tE", ALL_LANGUAGES)
 
+// Builtins for Intel SYCL
+BUILTIN(__builtin_intel_sycl_ptr_annotation, "v.", "nt")
+
 // Builtins for Intel FPGA
 BUILTIN(__builtin_intel_fpga_reg, "v.", "nt")
 BUILTIN(__builtin_intel_fpga_mem, "v.", "nt")

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -167,6 +167,13 @@ def err_intel_fpga_mem_arg_mismatch
             "a pointer"
             "|a non-negative integer constant}0">;
 
+def err_intel_sycl_ptr_annotation_arg_number_mismatch
+    :Error<"number of parameters must be odd number">;
+def err_intel_sycl_ptr_annotation_mismatch
+    : Error<"builtin parameter must be %select{"
+            "a pointer"
+            "|a string literal or constexpr const char*}0">;
+
 // C99 variable-length arrays
 def ext_vla : Extension<"variable length arrays are a C99 feature">,
   InGroup<VLAExtension>;

clang/include/clang/Sema/Sema.h

@@ -14015,6 +14015,9 @@ class Sema final {
   bool CheckIntelFPGARegBuiltinFunctionCall(unsigned BuiltinID, CallExpr *Call);
   bool CheckIntelFPGAMemBuiltinFunctionCall(CallExpr *Call);
 
+  bool CheckIntelSYCLPtrAnnotationBuiltinFunctionCall(unsigned BuiltinID,
+                                                      CallExpr *Call);
+
   bool SemaBuiltinVAStart(unsigned BuiltinID, CallExpr *TheCall);
   bool SemaBuiltinVAStartARMMicrosoft(CallExpr *Call);
   bool SemaBuiltinUnorderedCompare(CallExpr *TheCall);

clang/lib/CodeGen/CGBuiltin.cpp

@@ -5573,7 +5573,8 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
     return EmitIntelFPGARegBuiltin(E, ReturnValue);
   case Builtin::BI__builtin_intel_fpga_mem:
     return EmitIntelFPGAMemBuiltin(E);
-
+  case Builtin::BI__builtin_intel_sycl_ptr_annotation:
+    return EmitIntelSYCLPtrAnnotationBuiltin(E);
   case Builtin::BI__builtin_get_device_side_mangled_name: {
     auto Name = CGM.getCUDARuntime().getDeviceSideName(
         cast<DeclRefExpr>(E->getArg(0)->IgnoreImpCasts())->getDecl());
@@ -22540,6 +22541,43 @@ llvm::CallInst *CodeGenFunction::EmitFPBuiltinIndirectCall(
   return nullptr;
 }
 
+RValue CodeGenFunction::EmitIntelSYCLPtrAnnotationBuiltin(const CallExpr *E) {
+  const Expr *PtrArg = E->getArg(0);
+  Value *PtrVal = EmitScalarExpr(PtrArg);
+  auto &Ctx = CGM.getContext();
+
+  // Create the pointer annotation.
+  Function *F = CGM.getIntrinsic(llvm::Intrinsic::ptr_annotation,
+                                 {PtrVal->getType(), CGM.ConstGlobalsPtrTy});
+
+  SmallString<256> AnnotStr;
+  llvm::raw_svector_ostream Out(AnnotStr);
+
+  SmallVector<std::pair<std::string, std::string>, 4> Properties;
+
+  for (unsigned I = 1, N = E->getNumArgs(); I <= N / 2; I++) {
+    auto Arg = E->getArg(I)->IgnoreParenCasts();
+    const StringLiteral *Str = dyn_cast<const StringLiteral>(Arg);
+    Expr::EvalResult Result;
+    if (!Str && Arg->EvaluateAsRValue(Result, Ctx) && Result.Val.isLValue()) {
+      const auto *LVE = Result.Val.getLValueBase().dyn_cast<const Expr *>();
+      Str = dyn_cast<const StringLiteral>(LVE);
+    }
+    assert(Str && "Constant parameter string is invalid?");
+
+    auto IntVal = E->getArg(I + N / 2)->getIntegerConstantExpr(Ctx);
+    assert(IntVal.has_value() &&
+           "Constant integer arg isn't actually constant?");
+
+    Properties.push_back(
+        std::make_pair(Str->getString().str(), toString(IntVal.value(), 10)));
+  }
+
+  llvm::Value *Ann =
+      EmitSYCLAnnotationCall(F, PtrVal, E->getExprLoc(), Properties);
+  return RValue::get(Ann);
+}
+
 Value *CodeGenFunction::EmitRISCVBuiltinExpr(unsigned BuiltinID,
                                              const CallExpr *E,
                                              ReturnValueSlot ReturnValue) {

clang/lib/CodeGen/CodeGenFunction.cpp

@@ -2836,13 +2836,25 @@ Address CodeGenFunction::EmitFieldAnnotations(const FieldDecl *D,
 llvm::Value *CodeGenFunction::EmitSYCLAnnotationCall(
     llvm::Function *AnnotationFn, llvm::Value *AnnotatedVal,
     SourceLocation Location, const SYCLAddIRAnnotationsMemberAttr *Attr) {
+
+  llvm::SmallVector<std::pair<std::string, std::string>, 4>
+      AnnotationNameValPairs =
+          Attr->getFilteredAttributeNameValuePairs(getContext());
+  return EmitSYCLAnnotationCall(AnnotationFn, AnnotatedVal, Location,
+                                AnnotationNameValPairs);
+}
+
+llvm::Value *CodeGenFunction::EmitSYCLAnnotationCall(
+    llvm::Function *AnnotationFn, llvm::Value *AnnotatedVal,
+    SourceLocation Location,
+    SmallVectorImpl<std::pair<std::string, std::string>> &Pair) {
   SmallVector<llvm::Value *, 5> Args = {
       AnnotatedVal,
       Builder.CreateBitCast(CGM.EmitAnnotationString("sycl-properties"),
                             ConstGlobalsPtrTy),
       Builder.CreateBitCast(CGM.EmitAnnotationUnit(Location),
                             ConstGlobalsPtrTy),
-      CGM.EmitAnnotationLineNo(Location), CGM.EmitSYCLAnnotationArgs(Attr)};
+      CGM.EmitAnnotationLineNo(Location), CGM.EmitSYCLAnnotationArgs(Pair)};
   return Builder.CreateCall(AnnotationFn, Args);
 }
 

clang/lib/CodeGen/CodeGenFunction.h

@@ -4321,6 +4321,8 @@ class CodeGenFunction : public CodeGenTypeCache {
                                  ReturnValueSlot ReturnValue);
   RValue EmitIntelFPGAMemBuiltin(const CallExpr *E);
 
+  RValue EmitIntelSYCLPtrAnnotationBuiltin(const CallExpr *E);
+
   llvm::CallInst *
   EmitFPBuiltinIndirectCall(llvm::FunctionType *IRFuncTy,
                             const SmallVectorImpl<llvm::Value *> &IRArgs,
@@ -4589,6 +4591,11 @@ class CodeGenFunction : public CodeGenTypeCache {
                          llvm::Value *AnnotatedVal, SourceLocation Location,
                          const SYCLAddIRAnnotationsMemberAttr *Attr);
 
+  llvm::Value *EmitSYCLAnnotationCall(
+      llvm::Function *AnnotationFn, llvm::Value *AnnotatedVal,
+      SourceLocation Location,
+      llvm::SmallVectorImpl<std::pair<std::string, std::string>> &Pair);
+
   /// Emit sycl field annotations for given field & value. Returns the
   /// annotation result.
   Address EmitFieldSYCLAnnotations(const FieldDecl *D, Address V);

clang/lib/CodeGen/CodeGenModule.cpp

@@ -3496,10 +3496,9 @@ void CodeGenModule::AddGlobalAnnotations(const ValueDecl *D,
 }
 
 llvm::Constant *CodeGenModule::EmitSYCLAnnotationArgs(
-    const SYCLAddIRAnnotationsMemberAttr *Attr) {
-  llvm::SmallVector<std::pair<std::string, std::string>, 4>
-      AnnotationNameValPairs =
-          Attr->getFilteredAttributeNameValuePairs(getContext());
+    llvm::SmallVectorImpl<std::pair<std::string, std::string>>
+        &AnnotationNameValPairs) {
+
   if (AnnotationNameValPairs.empty())
     return llvm::ConstantPointerNull::get(ConstGlobalsPtrTy);
 

clang/lib/CodeGen/CodeGenModule.h

@@ -1382,8 +1382,8 @@ class CodeGenModule : public CodeGenTypeCache {
   void AddGlobalAnnotations(const ValueDecl *D, llvm::GlobalValue *GV);
 
   /// Emit additional args of the annotation.
-  llvm::Constant *
-  EmitSYCLAnnotationArgs(const SYCLAddIRAnnotationsMemberAttr *Attr);
+  llvm::Constant *EmitSYCLAnnotationArgs(
+      SmallVectorImpl<std::pair<std::string, std::string>> &Pairs);
 
   /// Add attributes from add_ir_attributes_global_variable on TND to GV.
   void AddGlobalSYCLIRAttributes(llvm::GlobalVariable *GV,

clang/lib/Sema/SemaChecking.cpp

@@ -102,6 +102,9 @@
 using namespace clang;
 using namespace sema;
 
+static const Expr *maybeConstEvalStringLiteral(ASTContext &Context,
+                                               const Expr *E);
+
 SourceLocation Sema::getLocationOfStringLiteralByte(const StringLiteral *SL,
                                                     unsigned ByteNo) const {
   return SL->getLocationOfByte(ByteNo, getSourceManager(), LangOpts,
@@ -2590,6 +2593,16 @@ Sema::CheckBuiltinFunctionCall(FunctionDecl *FDecl, unsigned BuiltinID,
     if (CheckIntelFPGARegBuiltinFunctionCall(BuiltinID, TheCall))
       return ExprError();
     break;
+  case Builtin::BI__builtin_intel_sycl_ptr_annotation:
+    if (!Context.getLangOpts().SYCLIsDevice) {
+      Diag(TheCall->getBeginLoc(), diag::err_builtin_requires_language)
+          << "__builtin_intel_sycl_ptr_annotation"
+          << "SYCL device";
+      return ExprError();
+    }
+    if (CheckIntelSYCLPtrAnnotationBuiltinFunctionCall(BuiltinID, TheCall))
+      return ExprError();
+    break;
   case Builtin::BI__builtin_intel_fpga_mem:
     if (!Context.getLangOpts().SYCLIsDevice) {
       Diag(TheCall->getBeginLoc(), diag::err_builtin_requires_language)
@@ -6018,6 +6031,56 @@ bool Sema::CheckIntelFPGAMemBuiltinFunctionCall(CallExpr *TheCall) {
   return false;
 }
 
+bool Sema::CheckIntelSYCLPtrAnnotationBuiltinFunctionCall(unsigned BuiltinID,
+                                                          CallExpr *TheCall) {
+  unsigned NumArgs = TheCall->getNumArgs();
+  // Make sure we have the minimum number of provided arguments.
+  if (checkArgCountAtLeast(*this, TheCall, 1)) {
+    return true;
+  }
+
+  // Make sure we have odd number of arguments.
+  if (!(NumArgs & 0x1)) {
+    return Diag(TheCall->getEndLoc(),
+                diag::err_intel_sycl_ptr_annotation_arg_number_mismatch);
+  }
+
+  // First argument should be a pointer.
+  Expr *PointerArg = TheCall->getArg(0);
+  QualType PointerArgType = PointerArg->getType();
+
+  if (!isa<PointerType>(PointerArgType))
+    return Diag(PointerArg->getBeginLoc(),
+                diag::err_intel_sycl_ptr_annotation_mismatch)
+           << 0;
+
+  // Following arguments are paired in format ("String", integer).
+  unsigned I = 1;
+  for (; I <= NumArgs / 2; ++I) {
+    // must be string Literal/const char*
+    auto Arg = TheCall->getArg(I)->IgnoreParenImpCasts();
+    Expr::EvalResult Result;
+    if (!isa<StringLiteral>(Arg) &&
+        !maybeConstEvalStringLiteral(this->Context, Arg)) {
+      Diag(TheCall->getArg(I)->getBeginLoc(),
+           diag::err_intel_sycl_ptr_annotation_mismatch)
+          << 1;
+      return true;
+    }
+  }
+
+  llvm::APSInt Result;
+  for (; I != NumArgs; ++I) {
+    // must be integer
+    if (SemaBuiltinConstantArg(TheCall, I, Result))
+      return true;
+  }
+
+  // Set the return type to be the same as the type of the first argument
+  TheCall->setType(PointerArgType);
+  return false;
+}
+
 /// Given a FunctionDecl's FormatAttr, attempts to populate the FomatStringInfo
 /// parameter with the FormatAttr's correct format_idx and firstDataArg.
 /// Returns true when the format fits the function and the FormatStringInfo has
@@ -9160,9 +9223,6 @@ static void CheckFormatString(
     llvm::SmallBitVector &CheckedVarArgs, UncoveredArgHandler &UncoveredArg,
     bool IgnoreStringsWithoutSpecifiers);
 
-static const Expr *maybeConstEvalStringLiteral(ASTContext &Context,
-                                               const Expr *E);
-
 // Determine if an expression is a string literal or constant string.
 // If this function returns false on the arguments to a function expecting a
 // format string, we will usually need to emit a warning.

clang/test/CodeGenSYCL/intel-sycl-ptr-annotation-builtin.cpp

@@ -0,0 +1,138 @@
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-linux -disable-llvm-passes -emit-llvm %s -internal-isystem %S/Inputs -o - | FileCheck %s
+
+// This test checks that using of __builtin_intel_sycl_ptr_annotation results in correct
+// generation of annotations in LLVM IR.
+#include "sycl.hpp"
+class kernel;
+// CHECK: [[STRUCT:%.*]] = type { i32, float }
+struct State {
+  int x;
+  float y;
+};
+
+// CHECK-DAG: [[ANN1:@.str[\.]*[0-9]*]] = {{.*}}"testA\00"
+// CHECK-DAG: [[ANN2:@.str[\.]*[0-9]*]] = {{.*}}"testB\00"
+// CHECK-DAG: [[ANN3:@.str[\.]*[0-9]*]] = {{.*}}"testC\00"
+// CHECK-DAG: [[ANN4:@.str[\.]*[0-9]*]] = {{.*}}"testD\00"
+// CHECK-DAG: [[ANN5:@.str[\.]*[0-9]*]] = {{.*}}"testE\00"
+// CHECK-DAG: [[ANN6:@.str[\.]*[0-9]*]] = {{.*}}"testF\00"
+// CHECK-DAG: [[ANN7:@.str[\.]*[0-9]*]] = {{.*}}"0\00"
+// CHECK-DAG: [[ANN8:@.str[\.]*[0-9]*]] = {{.*}}"127\00"
+// CHECK-DAG: [[ANN9:@.str[\.]*[0-9]*]] = {{.*}}"7\00"
+// CHECK-DAG: [[ANN10:@.str[\.]*[0-9]*]] = {{.*}}"8\00"
+// CHECK-DAG: [[ANN11:@.str[\.]*[0-9]*]] = {{.*}}"testG\00"
+// CHECK-DAG: [[ANN12:@.str[\.]*[0-9]*]] = {{.*}}"testH\00"
+// CHECK-DAG: [[ARG1:@.args[\.]*[0-9]*]] = {{.*}}[[ANN1]]{{.*}}[[ANN7]]
+// CHECK-DAG: [[ARG2:@.args[\.]*[0-9]*]] = {{.*}}[[ANN2]]{{.*}}[[ANN8]]
+// CHECK-DAG: [[ARG3:@.args[\.]*[0-9]*]] = {{.*}}[[ANN3]]{{.*}}[[ANN7]]
+// CHECK-DAG: [[ARG4:@.args[\.]*[0-9]*]] = {{.*}}[[ANN4]]{{.*}}[[ANN7]]
+// CHECK-DAG: [[ARG5:@.args[\.]*[0-9]*]] = {{.*}}[[ANN5]]{{.*}}[[ANN7]]
+// CHECK-DAG: [[ARG6:@.args[\.]*[0-9]*]] = {{.*}}[[ANN6]]{{.*}}[[ANN7]]{{.*}}[[ANN9]]{{.*}}[[ANN10]]
+// CHECK-DAG: [[ARG7:@.args[\.]*[0-9]*]] = {{.*}}[[ANN11]]{{.*}}[[ANN7]]
+// CHECK-DAG: [[ARG8:@.args[\.]*[0-9]*]] = {{.*}}[[ANN12]]{{.*}}[[ANN7]]
+
+
+// CHECK: define {{.*}}spir_func void @{{.*}}(ptr addrspace(4) noundef %A, ptr addrspace(4) noundef %B, ptr addrspace(4) noundef %C, ptr addrspace(4){{.*}}%D)
+void foo(float *A, int *B, State *C, State &D) {
+  float *x;
+  int *y;
+  State *z;
+  double *f;
+
+  // CHECK-DAG: [[Aaddr:%.*]] = alloca ptr addrspace(4)
+  // CHECK-DAG: [[Baddr:%.*]] = alloca ptr addrspace(4)
+  // CHECK-DAG: [[Caddr:%.*]] = alloca ptr addrspace(4)
+  // CHECK-DAG: [[Daddr:%.*]] = alloca ptr addrspace(4)
+
+  // CHECK-DAG: [[A:%[0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[Aaddr]]
+  // CHECK-DAG: [[PTR1:%[0-9]+]] = call ptr addrspace(4) @llvm.ptr.annotation{{.*}}[[A]]{{.*}}[[ARG1]]{{.*}}
+  // CHECK-DAG: store ptr addrspace(4) [[PTR1]], ptr addrspace(4) %x
+  x = __builtin_intel_sycl_ptr_annotation(A, "testA", 0);
+
+  // CHECK-DAG: [[B:%[0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[Baddr]]
+  // CHECK-DAG: [[PTR2:%[0-9]+]] = call ptr addrspace(4) @llvm.ptr.annotation{{.*}}[[B]]{{.*}}[[ARG1]]{{.*}}
+  // CHECK-DAG: store ptr addrspace(4) [[PTR2]], ptr addrspace(4) %y
+  y = __builtin_intel_sycl_ptr_annotation(B, "testA", 0);
+
+  // CHECK-DAG: [[C:%[0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[Caddr]]
+  // CHECK-DAG: [[PTR3:%[0-9]+]] = call ptr addrspace(4) @llvm.ptr.annotation{{.*}}[[C]]{{.*}}[[ARG1]]{{.*}}
+  // CHECK-DAG: store ptr addrspace(4) [[PTR3]], ptr addrspace(4) %z
+  z = __builtin_intel_sycl_ptr_annotation(C, "testA", 0);
+
+  // CHECK-DAG: [[A2:%[0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[Aaddr]]
+  // CHECK-DAG: [[PTR4:%[0-9]+]] = call ptr addrspace(4) @llvm.ptr.annotation{{.*}}[[A2]]{{.*}}[[ARG2]]{{.*}}
+  // CHECK-DAG: store ptr addrspace(4) [[PTR4]], ptr addrspace(4) %x
+  x = __builtin_intel_sycl_ptr_annotation(A, "testB", 127);
+
+  // CHECK-DAG: [[B2:%[0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[Baddr]]
+  // CHECK-DAG: [[PTR5:%[0-9]+]] = call ptr addrspace(4) @llvm.ptr.annotation{{.*}}[[B2]]{{.*}}[[ARG2]]{{.*}}
+  // CHECK-DAG: store ptr addrspace(4) [[PTR5]], ptr addrspace(4) %y
+  y = __builtin_intel_sycl_ptr_annotation(B, "testB", 127);
+
+  // CHECK-DAG: [[C2:%[0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[Caddr]]
+  // CHECK-DAG: [[PTR6:%[0-9]+]] = call ptr addrspace(4) @llvm.ptr.annotation{{.*}}[[C2]]{{.*}}[[ARG2]]{{.*}}
+  // CHECK-DAG: store ptr addrspace(4) [[PTR6]], ptr addrspace(4) %z
+  z = __builtin_intel_sycl_ptr_annotation(C, "testB", 127);
+
+  // CHECK-DAG: [[D:%[0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[Daddr]]
+  // CHECK-DAG: [[PTR7:%[0-9]+]] = call ptr addrspace(4) @llvm.ptr.annotation{{.*}}[[D]]{{.*}}[[ARG2]]{{.*}}
+  // CHECK-DAG: store ptr addrspace(4) [[PTR7]], ptr addrspace(4) %z
+  z = __builtin_intel_sycl_ptr_annotation(&D, "testB", 127);
+
+  // CHECK-DAG: [[A3:%[0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[Aaddr]]
+  // CHECK-DAG: [[PTR9:%[0-9]+]] = call ptr addrspace(4) @llvm.ptr.annotation{{.*}}[[A3]]{{.*}}[[ARG3]]{{.*}}
+  // CHECK-DAG: store ptr addrspace(4) [[PTR9]], ptr addrspace(4) %x
+  x = __builtin_intel_sycl_ptr_annotation(A, "testC", "testC", "testC", 0, 0, 0);
+
+  // CHECK-DAG: [[A4:%[0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[Aaddr]]
+  // CHECK-DAG: [[PTR10:%[0-9]+]] = call ptr addrspace(4) @llvm.ptr.annotation{{.*}}[[A4]]{{.*}}[[ARG4]]{{.*}}
+  // CHECK-DAG: store ptr addrspace(4) [[PTR10]], ptr addrspace(4) %x
+  x = __builtin_intel_sycl_ptr_annotation(A, "testD", "testD", "testD", 0, 0, 0);
+
+  // CHECK-DAG: [[B3:%[0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[Baddr]]
+  // CHECK-DAG: [[PTR11:%[0-9]+]] = call ptr addrspace(4) @llvm.ptr.annotation{{.*}}[[B3]]{{.*}}[[ARG5]]{{.*}}
+  // CHECK-DAG: store ptr addrspace(4) [[PTR11]], ptr addrspace(4) %y
+  y = __builtin_intel_sycl_ptr_annotation(B, "testE", "testE", 0, 0);
+
+  constexpr int TestVal1 = 7;
+  constexpr int TestVal2 = 8;
+
+  // CHECK-DAG: [[D1:%[0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[Daddr]]
+  // CHECK-DAG: [[PTR12:%[0-9]+]] = call ptr addrspace(4) @llvm.ptr.annotation{{.*}}[[D1]]{{.*}}[[ARG6]]{{.*}}
+  // CHECK-DAG: store ptr addrspace(4) [[PTR12]], ptr addrspace(4) %z
+  z = __builtin_intel_sycl_ptr_annotation(&D, "testF", "testF", "testF", "testF", 0, 0, TestVal1, TestVal2);
+}
+
+// This check makes sure the generated LoadInst consumes the annotated ptr directly
+// CHECK: define {{.*}}spir_func noundef i32 @{{.*}}(ptr addrspace(4) noundef %g)
+int annotation_with_load(int* g) {
+  // CHECK: [[PTR13:%[0-9]+]] = call ptr addrspace(4) @llvm.ptr.annotation{{.*}}[[ARG7]]{{.*}}
+  // CHECK: load i32, ptr addrspace(4) [[PTR13]]
+  return *__builtin_intel_sycl_ptr_annotation(g, "testG", 0);
+}
+
+// This check makes sure the generated StoreInst consumes the annotated ptr directly
+// CHECK: define {{.*}}spir_func void @{{.*}}(ptr addrspace(4) noundef %h)
+void annotation_with_store(int* h) {
+  // CHECK: [[PTR14:%[0-9]+]] = call ptr addrspace(4) @llvm.ptr.annotation{{.*}}[[ARG8]]{{.*}}
+  // CHECK: store i32 1, ptr addrspace(4) [[PTR14]]
+  *__builtin_intel_sycl_ptr_annotation(h, "testH", 0) = 1;
+}
+
+int main() {
+  sycl::queue q;
+  q.submit([&](sycl::handler &h) {
+    h.single_task<class kernel>([=](){
+        float *A;
+        int *B;
+        State *C;
+        State D;
+        foo(A, B, C, D);
+
+        int *a;
+        annotation_with_load(a);
+        annotation_with_store(a);
+    });
+  });
+  return 0;
+}