diff --git a/clang/lib/CodeGen/BackendUtil.cpp b/clang/lib/CodeGen/BackendUtil.cpp
index d214f03eaa7be..ffae6a10c821d 100644
--- a/clang/lib/CodeGen/BackendUtil.cpp
+++ b/clang/lib/CodeGen/BackendUtil.cpp
@@ -42,6 +42,7 @@
 #include "llvm/Passes/PassBuilder.h"
 #include "llvm/Passes/PassPlugin.h"
 #include "llvm/Passes/StandardInstrumentations.h"
+#include "llvm/SYCLLowerIR/ESIMDVerifier.h"
 #include "llvm/SYCLLowerIR/LowerWGLocalMemory.h"
 #include "llvm/Support/BuryPointer.h"
 #include "llvm/Support/CommandLine.h"
@@ -849,6 +850,9 @@ void EmitAssemblyHelper::CreatePasses(legacy::PassManager &MPM,
     FPM.add(createVerifierPass());
 
   // Set up the per-module pass manager.
+  if (LangOpts.SYCLIsDevice)
+    MPM.add(createESIMDVerifierPass());
+
   if (!CodeGenOpts.RewriteMapFiles.empty())
     addSymbolRewriterPass(CodeGenOpts, &MPM);
 
diff --git a/llvm/include/llvm/InitializePasses.h b/llvm/include/llvm/InitializePasses.h
index 8cec011926259..1843009a799b1 100644
--- a/llvm/include/llvm/InitializePasses.h
+++ b/llvm/include/llvm/InitializePasses.h
@@ -438,6 +438,7 @@ void initializeSYCLLowerESIMDLegacyPassPass(PassRegistry &);
 void initializeSPIRITTAnnotationsLegacyPassPass(PassRegistry &);
 void initializeESIMDLowerLoadStorePass(PassRegistry &);
 void initializeESIMDLowerVecArgLegacyPassPass(PassRegistry &);
+void initializeESIMDVerifierPass(PassRegistry &);
 void initializeSYCLLowerWGLocalMemoryLegacyPass(PassRegistry &);
 void initializeTailCallElimPass(PassRegistry&);
 void initializeTailDuplicatePass(PassRegistry&);
diff --git a/llvm/include/llvm/LinkAllPasses.h b/llvm/include/llvm/LinkAllPasses.h
index 5707a4d8abe1d..e7151ff20d63a 100644
--- a/llvm/include/llvm/LinkAllPasses.h
+++ b/llvm/include/llvm/LinkAllPasses.h
@@ -37,6 +37,7 @@
 #include "llvm/CodeGen/Passes.h"
 #include "llvm/IR/Function.h"
 #include "llvm/IR/IRPrintingPasses.h"
+#include "llvm/SYCLLowerIR/ESIMDVerifier.h"
 #include "llvm/SYCLLowerIR/LowerESIMD.h"
 #include "llvm/SYCLLowerIR/LowerWGLocalMemory.h"
 #include "llvm/SYCLLowerIR/LowerWGScope.h"
@@ -208,6 +209,7 @@ namespace {
       (void)llvm::createSYCLLowerESIMDPass();
       (void)llvm::createESIMDLowerLoadStorePass();
       (void)llvm::createESIMDLowerVecArgPass();
+      (void)llvm::createESIMDVerifierPass();
       (void)llvm::createSPIRITTAnnotationsLegacyPass();
       (void)llvm::createSYCLLowerWGLocalMemoryLegacyPass();
       std::string buf;
diff --git a/llvm/include/llvm/SYCLLowerIR/ESIMDVerifier.h b/llvm/include/llvm/SYCLLowerIR/ESIMDVerifier.h
new file mode 100644
index 0000000000000..d5ef20b27e232
--- /dev/null
+++ b/llvm/include/llvm/SYCLLowerIR/ESIMDVerifier.h
@@ -0,0 +1,30 @@
+//===--------- ESIMDVerifier.h - ESIMD-specific IR verification -----------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// ESIMD verification pass.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_SYCLLOWERIR_ESIMDVERIFIER_H
+#define LLVM_SYCLLOWERIR_ESIMDVERIFIER_H
+
+#include "llvm/IR/PassManager.h"
+
+namespace llvm {
+
+struct ESIMDVerifierPass : public PassInfoMixin<ESIMDVerifierPass> {
+  ESIMDVerifierPass() {}
+  PreservedAnalyses run(Module &M, ModuleAnalysisManager &);
+  static bool isRequired() { return true; }
+};
+
+ModulePass *createESIMDVerifierPass();
+
+} // namespace llvm
+
+#endif // LLVM_SYCLLOWERIR_ESIMDVERIFIER_H
diff --git a/llvm/lib/Passes/PassBuilder.cpp b/llvm/lib/Passes/PassBuilder.cpp
index 074c1f354d5bd..4ae93b76eadd5 100644
--- a/llvm/lib/Passes/PassBuilder.cpp
+++ b/llvm/lib/Passes/PassBuilder.cpp
@@ -75,6 +75,7 @@
 #include "llvm/IR/PrintPasses.h"
 #include "llvm/IR/SafepointIRVerifier.h"
 #include "llvm/IR/Verifier.h"
+#include "llvm/SYCLLowerIR/ESIMDVerifier.h"
 #include "llvm/SYCLLowerIR/LowerESIMD.h"
 #include "llvm/SYCLLowerIR/LowerWGScope.h"
 #include "llvm/Support/CommandLine.h"
diff --git a/llvm/lib/Passes/PassRegistry.def b/llvm/lib/Passes/PassRegistry.def
index 28470c5b3f2d5..1331b75cf9113 100644
--- a/llvm/lib/Passes/PassRegistry.def
+++ b/llvm/lib/Passes/PassRegistry.def
@@ -117,6 +117,7 @@ MODULE_PASS("poison-checking", PoisonCheckingPass())
 MODULE_PASS("pseudo-probe-update", PseudoProbeUpdatePass())
 MODULE_PASS("LowerESIMD", SYCLLowerESIMDPass())
 MODULE_PASS("ESIMDLowerVecArg", ESIMDLowerVecArgPass())
+MODULE_PASS("esimd-verifier", ESIMDVerifierPass())
 MODULE_PASS("SPIRITTAnnotations", SPIRITTAnnotationsPass())
 MODULE_PASS("deadargelim-sycl", DeadArgumentEliminationSYCLPass())
 #undef MODULE_PASS
diff --git a/llvm/lib/SYCLLowerIR/CMakeLists.txt b/llvm/lib/SYCLLowerIR/CMakeLists.txt
index 900e62b9be5e1..c68f72b1b33fc 100644
--- a/llvm/lib/SYCLLowerIR/CMakeLists.txt
+++ b/llvm/lib/SYCLLowerIR/CMakeLists.txt
@@ -53,6 +53,7 @@ add_llvm_component_library(LLVMSYCLLowerIR
   LowerESIMDVLoadVStore.cpp
   LowerESIMDVecArg.cpp
   LowerWGLocalMemory.cpp
+  ESIMDVerifier.cpp
 
   ADDITIONAL_HEADER_DIRS
   ${LLVM_MAIN_INCLUDE_DIR}/llvm/SYCLLowerIR
diff --git a/llvm/lib/SYCLLowerIR/ESIMDVerifier.cpp b/llvm/lib/SYCLLowerIR/ESIMDVerifier.cpp
new file mode 100644
index 0000000000000..6535e7f20f652
--- /dev/null
+++ b/llvm/lib/SYCLLowerIR/ESIMDVerifier.cpp
@@ -0,0 +1,122 @@
+//===---------- ESIMDVerifier.cpp - ESIMD-specific IR verification --------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file implements ESIMD specific IR verification pass. So far it only
+// detects invalid API calls in ESIMD context.
+//
+//===----------------------------------------------------------------------===//
+
+#include "llvm/SYCLLowerIR/ESIMDVerifier.h"
+#include "llvm/Demangle/Demangle.h"
+#include "llvm/IR/InstIterator.h"
+#include "llvm/IR/Instructions.h"
+#include "llvm/IR/Module.h"
+#include "llvm/InitializePasses.h"
+#include "llvm/Pass.h"
+#include "llvm/Support/Debug.h"
+#include "llvm/Support/Regex.h"
+
+using namespace llvm;
+
+#define DEBUG_TYPE "esimd-verifier"
+
+// A list of unsupported functions in ESIMD context.
+static const char *IllegalFunctions[] = {
+    "^cl::sycl::multi_ptr<.+> cl::sycl::accessor<.+>::get_pointer<.+>\\(\\) "
+    "const"};
+
+namespace {
+
+class ESIMDVerifierImpl {
+  const Module &M;
+
+public:
+  ESIMDVerifierImpl(const Module &M) : M(M) {}
+
+  void verify() {
+    SmallPtrSet<const Function *, 8u> Visited;
+    SmallVector<const Function *, 8u> Worklist;
+
+    auto Add2Worklist = [&Worklist, &Visited](const Function *F) {
+      if (Visited.insert(F).second)
+        Worklist.push_back(F);
+    };
+
+    // Start with adding all ESIMD functions to the work list.
+    for (const Function &F : M)
+      if (F.hasMetadata("sycl_explicit_simd"))
+        Add2Worklist(&F);
+
+    // Then check ESIMD functions and all functions called from ESIMD context
+    // for invalid calls.
+    while (!Worklist.empty()) {
+      const Function *F = Worklist.pop_back_val();
+      for (const Instruction &I : instructions(F)) {
+        if (auto *CB = dyn_cast<CallBase>(&I)) {
+          Function *Callee = CB->getCalledFunction();
+          if (!Callee)
+            continue;
+
+          // Demangle called function name and check if it matches any illegal
+          // function name. Report an error if there is a match.
+          std::string DemangledName = demangle(Callee->getName().str());
+          for (const char *Name : IllegalFunctions) {
+            Regex NameRE(Name);
+            assert(NameRE.isValid() && "invalid function name regex");
+            if (NameRE.match(DemangledName)) {
+              std::string ErrorMsg = std::string("function '") + DemangledName +
+                                     "' is not supported in ESIMD context";
+              F->getContext().emitError(&I, ErrorMsg);
+            }
+          }
+
+          // Add callee to the list to be analyzed if it is not a declaration.
+          if (!Callee->isDeclaration())
+            Add2Worklist(Callee);
+        }
+      }
+    }
+  }
+};
+
+} // end anonymous namespace
+
+PreservedAnalyses ESIMDVerifierPass::run(Module &M, ModuleAnalysisManager &AM) {
+  ESIMDVerifierImpl(M).verify();
+  return PreservedAnalyses::all();
+}
+
+namespace {
+
+struct ESIMDVerifier : public ModulePass {
+  static char ID;
+
+  ESIMDVerifier() : ModulePass(ID) {
+    initializeESIMDVerifierPass(*PassRegistry::getPassRegistry());
+  }
+
+  void getAnalysisUsage(AnalysisUsage &AU) const override {
+    AU.setPreservesAll();
+  }
+
+  bool runOnModule(Module &M) override {
+    ESIMDVerifierImpl(M).verify();
+    return false;
+  }
+};
+
+} // end anonymous namespace
+
+char ESIMDVerifier::ID = 0;
+
+INITIALIZE_PASS_BEGIN(ESIMDVerifier, DEBUG_TYPE, "ESIMD-specific IR verifier",
+                      false, false)
+INITIALIZE_PASS_END(ESIMDVerifier, DEBUG_TYPE, "ESIMD-specific IR verifier",
+                    false, false)
+
+ModulePass *llvm::createESIMDVerifierPass() { return new ESIMDVerifier(); }
diff --git a/llvm/tools/opt/opt.cpp b/llvm/tools/opt/opt.cpp
index 165083b251797..e0fe6de4de12a 100644
--- a/llvm/tools/opt/opt.cpp
+++ b/llvm/tools/opt/opt.cpp
@@ -579,6 +579,7 @@ int main(int argc, char **argv) {
   initializeSPIRITTAnnotationsLegacyPassPass(Registry);
   initializeESIMDLowerLoadStorePass(Registry);
   initializeESIMDLowerVecArgLegacyPassPass(Registry);
+  initializeESIMDVerifierPass(Registry);
   initializeSYCLLowerWGLocalMemoryLegacyPass(Registry);
 
 #ifdef BUILD_EXAMPLES
diff --git a/sycl/test/esimd/esimd_verify.cpp b/sycl/test/esimd/esimd_verify.cpp
new file mode 100644
index 0000000000000..e9b91294187d2
--- /dev/null
+++ b/sycl/test/esimd/esimd_verify.cpp
@@ -0,0 +1,14 @@
+// RUN: not %clangxx -fsycl -fsycl-device-only -S %s -o %t 2>&1 | FileCheck %s
+
+#include <sycl/ext/intel/experimental/esimd.hpp>
+
+using namespace cl::sycl;
+using namespace sycl::ext::intel::experimental::esimd;
+
+// CHECK: error: function 'cl::sycl::multi_ptr<{{.+}}> cl::sycl::accessor<{{.+}}>::get_pointer<{{.+}}>() const' is not supported in ESIMD context
+
+SYCL_EXTERNAL auto
+test(accessor<int, 1, access::mode::read_write, access::target::device> &acc)
+    SYCL_ESIMD_FUNCTION {
+  return acc.get_pointer();
+}
diff --git a/sycl/test/esimd/spirv_intrins_trans.cpp b/sycl/test/esimd/spirv_intrins_trans.cpp
index 7f7da671e01fb..8aa6dc17d8d0a 100644
--- a/sycl/test/esimd/spirv_intrins_trans.cpp
+++ b/sycl/test/esimd/spirv_intrins_trans.cpp
@@ -16,10 +16,10 @@ __attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
 
 size_t caller() {
 
-  size_t DoNotOpt;
-  cl::sycl::buffer<size_t, 1> buf(&DoNotOpt, 1);
-  uint32_t DoNotOpt32;
-  cl::sycl::buffer<uint32_t, 1> buf32(&DoNotOpt32, 1);
+  size_t DoNotOpt[1];
+  cl::sycl::buffer<size_t, 1> buf(&DoNotOpt[0], 1);
+  uint32_t DoNotOpt32[1];
+  cl::sycl::buffer<uint32_t, 1> buf32(&DoNotOpt32[0], 1);
 
   size_t DoNotOptXYZ[3];
   cl::sycl::buffer<size_t, 1> bufXYZ(&DoNotOptXYZ[0], sycl::range<1>(3));
@@ -29,7 +29,7 @@ size_t caller() {
     auto DoNotOptimize32 = buf32.get_access<cl::sycl::access::mode::write>(cgh);
 
     kernel<class kernel_GlobalInvocationId_x>([=]() SYCL_ESIMD_KERNEL {
-      *DoNotOptimize.get_pointer() = __spirv_GlobalInvocationId_x();
+      DoNotOptimize[0] = __spirv_GlobalInvocationId_x();
     });
     // CHECK-LABEL: @{{.*}}kernel_GlobalInvocationId_x
     // CHECK: [[CALL_ESIMD1:%.*]] = call <3 x i32> @llvm.genx.local.id.v3i32()
@@ -39,7 +39,7 @@ size_t caller() {
     // CHECK: {{.*}} call i32 @llvm.genx.group.id.x()
 
     kernel<class kernel_GlobalInvocationId_y>([=]() SYCL_ESIMD_KERNEL {
-      *DoNotOptimize.get_pointer() = __spirv_GlobalInvocationId_y();
+      DoNotOptimize[0] = __spirv_GlobalInvocationId_y();
     });
     // CHECK-LABEL: @{{.*}}kernel_GlobalInvocationId_y
     // CHECK: [[CALL_ESIMD1:%.*]] = call <3 x i32> @llvm.genx.local.id.v3i32()
@@ -49,7 +49,7 @@ size_t caller() {
     // CHECK: {{.*}} call i32 @llvm.genx.group.id.y()
 
     kernel<class kernel_GlobalInvocationId_z>([=]() SYCL_ESIMD_KERNEL {
-      *DoNotOptimize.get_pointer() = __spirv_GlobalInvocationId_z();
+      DoNotOptimize[0] = __spirv_GlobalInvocationId_z();
     });
     // CHECK-LABEL: @{{.*}}kernel_GlobalInvocationId_z
     // CHECK: [[CALL_ESIMD1:%.*]] = call <3 x i32> @llvm.genx.local.id.v3i32()
@@ -58,27 +58,24 @@ size_t caller() {
     // CHECK: {{.*}} extractelement <3 x i32> [[CALL_ESIMD2]], i32 2
     // CHECK: {{.*}} call i32 @llvm.genx.group.id.z()
 
-    kernel<class kernel_GlobalSize_x>([=]() SYCL_ESIMD_KERNEL {
-      *DoNotOptimize.get_pointer() = __spirv_GlobalSize_x();
-    });
+    kernel<class kernel_GlobalSize_x>(
+        [=]() SYCL_ESIMD_KERNEL { DoNotOptimize[0] = __spirv_GlobalSize_x(); });
     // CHECK-LABEL: @{{.*}}kernel_GlobalSize_x
     // CHECK: [[CALL_ESIMD1:%.*]] = call <3 x i32> @llvm.genx.local.size.v3i32()
     // CHECK: {{.*}} extractelement <3 x i32> [[CALL_ESIMD1]], i32 0
     // CHECK: [[CALL_ESIMD2:%.*]] = call <3 x i32> @llvm.genx.group.count.v3i32()
     // CHECK: {{.*}} extractelement <3 x i32> [[CALL_ESIMD2]], i32 0
 
-    kernel<class kernel_GlobalSize_y>([=]() SYCL_ESIMD_KERNEL {
-      *DoNotOptimize.get_pointer() = __spirv_GlobalSize_y();
-    });
+    kernel<class kernel_GlobalSize_y>(
+        [=]() SYCL_ESIMD_KERNEL { DoNotOptimize[0] = __spirv_GlobalSize_y(); });
     // CHECK-LABEL: @{{.*}}kernel_GlobalSize_y
     // CHECK: [[CALL_ESIMD1:%.*]] = call <3 x i32> @llvm.genx.local.size.v3i32()
     // CHECK: {{.*}} extractelement <3 x i32> [[CALL_ESIMD1]], i32 1
     // CHECK: [[CALL_ESIMD2:%.*]] = call <3 x i32> @llvm.genx.group.count.v3i32()
     // CHECK: {{.*}} extractelement <3 x i32> [[CALL_ESIMD2]], i32 1
 
-    kernel<class kernel_GlobalSize_z>([=]() SYCL_ESIMD_KERNEL {
-      *DoNotOptimize.get_pointer() = __spirv_GlobalSize_z();
-    });
+    kernel<class kernel_GlobalSize_z>(
+        [=]() SYCL_ESIMD_KERNEL { DoNotOptimize[0] = __spirv_GlobalSize_z(); });
     // CHECK-LABEL: @{{.*}}kernel_GlobalSize_z
     // CHECK: [[CALL_ESIMD1:%.*]] = call <3 x i32> @llvm.genx.local.size.v3i32()
     // CHECK: {{.*}} extractelement <3 x i32> [[CALL_ESIMD1]], i32 2
@@ -86,99 +83,99 @@ size_t caller() {
     // CHECK: {{.*}} extractelement <3 x i32> [[CALL_ESIMD2]], i32 2
 
     kernel<class kernel_GlobalOffset_x>([=]() SYCL_ESIMD_KERNEL {
-      *DoNotOptimize.get_pointer() = __spirv_GlobalOffset_x();
+      DoNotOptimize[0] = __spirv_GlobalOffset_x();
     });
     // CHECK-LABEL: @{{.*}}kernel_GlobalOffset_x
     // CHECK: store i64 0
 
     kernel<class kernel_GlobalOffset_y>([=]() SYCL_ESIMD_KERNEL {
-      *DoNotOptimize.get_pointer() = __spirv_GlobalOffset_y();
+      DoNotOptimize[0] = __spirv_GlobalOffset_y();
     });
     // CHECK-LABEL: @{{.*}}kernel_GlobalOffset_y
     // CHECK: store i64 0
 
     kernel<class kernel_GlobalOffset_z>([=]() SYCL_ESIMD_KERNEL {
-      *DoNotOptimize.get_pointer() = __spirv_GlobalOffset_z();
+      DoNotOptimize[0] = __spirv_GlobalOffset_z();
     });
     // CHECK-LABEL: @{{.*}}kernel_GlobalOffset_z
     // CHECK: store i64 0
 
     kernel<class kernel_NumWorkgroups_x>([=]() SYCL_ESIMD_KERNEL {
-      *DoNotOptimize.get_pointer() = __spirv_NumWorkgroups_x();
+      DoNotOptimize[0] = __spirv_NumWorkgroups_x();
     });
     // CHECK-LABEL: @{{.*}}kernel_NumWorkgroups_x
     // CHECK: [[CALL_ESIMD:%.*]] = call <3 x i32> @llvm.genx.group.count.v3i32()
     // CHECK: {{.*}} extractelement <3 x i32> [[CALL_ESIMD]], i32 0
 
     kernel<class kernel_NumWorkgroups_y>([=]() SYCL_ESIMD_KERNEL {
-      *DoNotOptimize.get_pointer() = __spirv_NumWorkgroups_y();
+      DoNotOptimize[0] = __spirv_NumWorkgroups_y();
     });
     // CHECK-LABEL: @{{.*}}kernel_NumWorkgroups_y
     // CHECK: [[CALL_ESIMD:%.*]] = call <3 x i32> @llvm.genx.group.count.v3i32()
     // CHECK: {{.*}} extractelement <3 x i32> [[CALL_ESIMD]], i32 1
 
     kernel<class kernel_NumWorkgroups_z>([=]() SYCL_ESIMD_KERNEL {
-      *DoNotOptimize.get_pointer() = __spirv_NumWorkgroups_z();
+      DoNotOptimize[0] = __spirv_NumWorkgroups_z();
     });
     // CHECK-LABEL: @{{.*}}kernel_NumWorkgroups_z
     // CHECK: [[CALL_ESIMD:%.*]] = call <3 x i32> @llvm.genx.group.count.v3i32()
     // CHECK: {{.*}} extractelement <3 x i32> [[CALL_ESIMD]], i32 2
 
     kernel<class kernel_WorkgroupSize_x>([=]() SYCL_ESIMD_KERNEL {
-      *DoNotOptimize.get_pointer() = __spirv_WorkgroupSize_x();
+      DoNotOptimize[0] = __spirv_WorkgroupSize_x();
     });
     // CHECK-LABEL: @{{.*}}kernel_WorkgroupSize_x
     // CHECK: [[CALL_ESIMD:%.*]] = call <3 x i32> @llvm.genx.local.size.v3i32()
     // CHECK: {{.*}} extractelement <3 x i32> [[CALL_ESIMD]], i32 0
 
     kernel<class kernel_WorkgroupSize_y>([=]() SYCL_ESIMD_KERNEL {
-      *DoNotOptimize.get_pointer() = __spirv_WorkgroupSize_y();
+      DoNotOptimize[0] = __spirv_WorkgroupSize_y();
     });
     // CHECK-LABEL: @{{.*}}kernel_WorkgroupSize_y
     // CHECK: [[CALL_ESIMD:%.*]] = call <3 x i32> @llvm.genx.local.size.v3i32()
     // CHECK: {{.*}} extractelement <3 x i32> [[CALL_ESIMD]], i32 1
 
     kernel<class kernel_WorkgroupSize_z>([=]() SYCL_ESIMD_KERNEL {
-      *DoNotOptimize.get_pointer() = __spirv_WorkgroupSize_z();
+      DoNotOptimize[0] = __spirv_WorkgroupSize_z();
     });
     // CHECK-LABEL: @{{.*}}kernel_WorkgroupSize_z
     // CHECK: [[CALL_ESIMD:%.*]] = call <3 x i32> @llvm.genx.local.size.v3i32()
     // CHECK: {{.*}} extractelement <3 x i32> [[CALL_ESIMD]], i32 2
 
     kernel<class kernel_WorkgroupId_x>([=]() SYCL_ESIMD_KERNEL {
-      *DoNotOptimize.get_pointer() = __spirv_WorkgroupId_x();
+      DoNotOptimize[0] = __spirv_WorkgroupId_x();
     });
     // CHECK-LABEL: @{{.*}}kernel_WorkgroupId_x
     // CHECK: {{.*}} call i32 @llvm.genx.group.id.x()
 
     kernel<class kernel_WorkgroupId_y>([=]() SYCL_ESIMD_KERNEL {
-      *DoNotOptimize.get_pointer() = __spirv_WorkgroupId_y();
+      DoNotOptimize[0] = __spirv_WorkgroupId_y();
     });
     // CHECK-LABEL: @{{.*}}kernel_WorkgroupId_y
     // CHECK: {{.*}} call i32 @llvm.genx.group.id.y()
 
     kernel<class kernel_WorkgroupId_z>([=]() SYCL_ESIMD_KERNEL {
-      *DoNotOptimize.get_pointer() = __spirv_WorkgroupId_z();
+      DoNotOptimize[0] = __spirv_WorkgroupId_z();
     });
     // CHECK-LABEL: @{{.*}}kernel_WorkgroupId_z
     // CHECK: {{.*}} call i32 @llvm.genx.group.id.z()
 
     kernel<class kernel_LocalInvocationId_x>([=]() SYCL_ESIMD_KERNEL {
-      *DoNotOptimize.get_pointer() = __spirv_LocalInvocationId_x();
+      DoNotOptimize[0] = __spirv_LocalInvocationId_x();
     });
     // CHECK-LABEL: @{{.*}}kernel_LocalInvocationId_x
     // CHECK: [[CALL_ESIMD:%.*]] = call <3 x i32> @llvm.genx.local.id.v3i32()
     // CHECK: {{.*}} extractelement <3 x i32> [[CALL_ESIMD]], i32 0
 
     kernel<class kernel_LocalInvocationId_y>([=]() SYCL_ESIMD_KERNEL {
-      *DoNotOptimize.get_pointer() = __spirv_LocalInvocationId_y();
+      DoNotOptimize[0] = __spirv_LocalInvocationId_y();
     });
     // CHECK-LABEL: @{{.*}}kernel_LocalInvocationId_y
     // CHECK: [[CALL_ESIMD:%.*]] = call <3 x i32> @llvm.genx.local.id.v3i32()
     // CHECK: {{.*}} extractelement <3 x i32> [[CALL_ESIMD]], i32 1
 
     kernel<class kernel_LocalInvocationId_z>([=]() SYCL_ESIMD_KERNEL {
-      *DoNotOptimize.get_pointer() = __spirv_LocalInvocationId_z();
+      DoNotOptimize[0] = __spirv_LocalInvocationId_z();
     });
     // CHECK-LABEL: @{{.*}}kernel_LocalInvocationId_z
     // CHECK: [[CALL_ESIMD:%.*]] = call <3 x i32> @llvm.genx.local.id.v3i32()
@@ -218,8 +215,8 @@ size_t caller() {
     // CHECK: {{.*}} call i32 @llvm.genx.group.id.z()
 
     kernel<class kernel_SubgroupLocalInvocationId>([=]() SYCL_ESIMD_KERNEL {
-      *DoNotOptimize.get_pointer() = __spirv_SubgroupLocalInvocationId();
-      *DoNotOptimize32.get_pointer() = __spirv_SubgroupLocalInvocationId() + 3;
+      DoNotOptimize[0] = __spirv_SubgroupLocalInvocationId();
+      DoNotOptimize32[0] = __spirv_SubgroupLocalInvocationId() + 3;
     });
     // CHECK-LABEL: @{{.*}}kernel_SubgroupLocalInvocationId
     // CHECK: [[ZEXT0:%.*]] = zext i32 0 to i64
@@ -227,8 +224,8 @@ size_t caller() {
     // CHECK: add i32 0, 3
 
     kernel<class kernel_SubgroupSize>([=]() SYCL_ESIMD_KERNEL {
-      *DoNotOptimize.get_pointer() = __spirv_SubgroupSize();
-      *DoNotOptimize32.get_pointer() = __spirv_SubgroupSize() + 7;
+      DoNotOptimize[0] = __spirv_SubgroupSize();
+      DoNotOptimize32[0] = __spirv_SubgroupSize() + 7;
     });
     // CHECK-LABEL: @{{.*}}kernel_SubgroupSize
     // CHECK: [[ZEXT0:%.*]] = zext i32 1 to i64
@@ -236,13 +233,13 @@ size_t caller() {
     // CHECK: add i32 1, 7
 
     kernel<class kernel_SubgroupMaxSize>([=]() SYCL_ESIMD_KERNEL {
-      *DoNotOptimize.get_pointer() = __spirv_SubgroupMaxSize();
-      *DoNotOptimize32.get_pointer() = __spirv_SubgroupMaxSize() + 9;
+      DoNotOptimize[0] = __spirv_SubgroupMaxSize();
+      DoNotOptimize32[0] = __spirv_SubgroupMaxSize() + 9;
     });
     // CHECK-LABEL: @{{.*}}kernel_SubgroupMaxSize
     // CHECK: [[ZEXT0:%.*]] = zext i32 1 to i64
     // CHECK: store i64 [[ZEXT0]]
     // CHECK: add i32 1, 9
   });
-  return DoNotOpt;
+  return DoNotOpt[0];
 }
diff --git a/sycl/test/esimd/vadd.cpp b/sycl/test/esimd/vadd.cpp
index c1159fd29cf9d..90dcaf9ab3421 100644
--- a/sycl/test/esimd/vadd.cpp
+++ b/sycl/test/esimd/vadd.cpp
@@ -54,19 +54,33 @@ int main(void) {
   constexpr unsigned VL = 32;
   constexpr unsigned GroupSize = 2;
 
-  int A[Size];
-  int B[Size];
-  int C[Size] = {};
+  struct Deleter {
+    queue Q;
+    void operator()(int *Ptr) {
+      if (Ptr) {
+        sycl::free(Ptr, Q);
+      }
+    }
+  };
+
+  queue q(ESIMDSelector{}, exception_handler);
+
+  std::unique_ptr<int, Deleter> BufA(sycl::malloc_shared<int>(Size, q),
+                                     Deleter{q});
+  std::unique_ptr<int, Deleter> BufB(
+      sycl::aligned_alloc_shared<int>(16u, Size, q), Deleter{q});
+  std::unique_ptr<int, Deleter> BufC(
+      sycl::aligned_alloc_shared<int>(16u, Size, q), Deleter{q});
+
+  int *A = BufA.get();
+  int *B = BufB.get();
+  int *C = BufC.get();
 
   for (unsigned i = 0; i < Size; ++i) {
     A[i] = B[i] = i;
   }
 
   {
-    cl::sycl::buffer<int, 1> bufA(A, Size);
-    cl::sycl::buffer<int, 1> bufB(B, Size);
-    cl::sycl::buffer<int, 1> bufC(C, Size);
-
     // We need that many task groups
     cl::sycl::range<1> GroupRange{Size / VL};
 
@@ -75,28 +89,20 @@ int main(void) {
 
     cl::sycl::nd_range<1> Range{GroupRange, TaskRange};
 
-    queue q(ESIMDSelector{}, exception_handler);
     q.submit([&](cl::sycl::handler &cgh) {
-      auto accA = bufA.get_access<cl::sycl::access::mode::read>(cgh);
-      auto accB = bufB.get_access<cl::sycl::access::mode::read>(cgh);
-      auto accC = bufC.get_access<cl::sycl::access::mode::write>(cgh);
-
       cgh.parallel_for<class Test>(
           Range, [=](nd_item<1> ndi) SYCL_ESIMD_KERNEL {
             using namespace sycl::ext::intel::experimental::esimd;
-            auto pA = accA.get_pointer().get();
-            auto pB = accB.get_pointer().get();
-            auto pC = accC.get_pointer().get();
 
             int i = ndi.get_global_id(0);
             constexpr int ESIZE = sizeof(int);
             simd<uint32_t, VL> offsets(0, ESIZE);
 
-            simd<int, VL> va = gather<int, VL>(pA + i * VL, offsets);
-            simd<int, VL> vb = block_load<int, VL>(pB + i * VL);
+            simd<int, VL> va = gather<int, VL>(A + i * VL, offsets);
+            simd<int, VL> vb = block_load<int, VL>(B + i * VL);
             simd<int, VL> vc = va + vb;
 
-            block_store<int, VL>(pC + i * VL, vc);
+            block_store<int, VL>(C + i * VL, vc);
           });
     });
   }