diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
index 695bf13a5a762..6c9bdc364e0df 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -10548,7 +10548,8 @@ static void getSPIRVBackendOpts(const llvm::opt::ArgList &TCArgs,
                             ",+SPV_INTEL_subgroups"
                             ",+SPV_INTEL_tensor_float32_conversion"
                             ",+SPV_INTEL_variable_length_array"
-                            ",+SPV_INTEL_memory_access_aliasing";
+                            ",+SPV_INTEL_memory_access_aliasing"
+                            ",+SPV_INTEL_global_variable_host_access";
   std::string KHRExtArg = ",+SPV_KHR_16bit_storage"
                           ",+SPV_KHR_cooperative_matrix"
                           ",+SPV_KHR_expect_assume"
@@ -10644,7 +10645,8 @@ static void getTripleBasedSPIRVTransOpts(Compilation &C,
       ",+SPV_KHR_cooperative_matrix"
       ",+SPV_EXT_shader_atomic_float16_add"
       ",+SPV_INTEL_fp_max_error"
-      ",+SPV_INTEL_memory_access_aliasing";
+      ",+SPV_INTEL_memory_access_aliasing"
+      ",+SPV_INTEL_global_variable_host_access";
 
   TranslatorArgs.push_back(TCArgs.MakeArgString(ExtArg));
 }
diff --git a/clang/test/Driver/sycl-spirv-backend.cpp b/clang/test/Driver/sycl-spirv-backend.cpp
index 148f24548b3cd..d0b1270d46e62 100644
--- a/clang/test/Driver/sycl-spirv-backend.cpp
+++ b/clang/test/Driver/sycl-spirv-backend.cpp
@@ -23,6 +23,7 @@
 // CHECK-SAME:,+SPV_INTEL_tensor_float32_conversion
 // CHECK-SAME:,+SPV_INTEL_variable_length_array
 // CHECK-SAME:,+SPV_INTEL_memory_access_aliasing
+// CHECK-SAME:,+SPV_INTEL_global_variable_host_access
 // CHECK-SAME:,+SPV_KHR_16bit_storage
 // CHECK-SAME:,+SPV_KHR_cooperative_matrix
 // CHECK-SAME:,+SPV_KHR_expect_assume
diff --git a/clang/test/Driver/sycl-spirv-ext-old-model.cpp b/clang/test/Driver/sycl-spirv-ext-old-model.cpp
index f3c920979841b..a594f1eabd527 100644
--- a/clang/test/Driver/sycl-spirv-ext-old-model.cpp
+++ b/clang/test/Driver/sycl-spirv-ext-old-model.cpp
@@ -83,3 +83,4 @@
 // CHECK-CPU-SAME:,+SPV_KHR_non_semantic_info
 // CHECK-CPU-SAME:,+SPV_KHR_cooperative_matrix
 // CHECK-CPU-SAME:,+SPV_INTEL_fp_max_error
+// CHECK-CPU-SAME:,+SPV_INTEL_global_variable_host_access
diff --git a/clang/test/Driver/sycl-spirv-ext.cpp b/clang/test/Driver/sycl-spirv-ext.cpp
index a8394d1ece837..26b484e7fd2f0 100644
--- a/clang/test/Driver/sycl-spirv-ext.cpp
+++ b/clang/test/Driver/sycl-spirv-ext.cpp
@@ -100,3 +100,4 @@
 // CHECK-CPU-SAME:,+SPV_KHR_non_semantic_info
 // CHECK-CPU-SAME:,+SPV_KHR_cooperative_matrix
 // CHECK-CPU-SAME:,+SPV_INTEL_fp_max_error
+// CHECK-CPU-SAME:,+SPV_INTEL_global_variable_host_access
diff --git a/clang/test/Driver/sycl-spirv-metadata-old-model.cpp b/clang/test/Driver/sycl-spirv-metadata-old-model.cpp
index 2e19dd9ed1dc4..935e11e9fa6c2 100644
--- a/clang/test/Driver/sycl-spirv-metadata-old-model.cpp
+++ b/clang/test/Driver/sycl-spirv-metadata-old-model.cpp
@@ -9,7 +9,7 @@
 // RUN:  FileCheck -check-prefix CHECK-WITHOUT %s
 
 // CHECK-WITH: llvm-spirv{{.*}} "--spirv-preserve-auxdata"
-// CHECK-WITH-SAME: "-spirv-ext=-all,{{.*}},+SPV_INTEL_memory_access_aliasing"
+// CHECK-WITH-SAME: "-spirv-ext=-all,{{.*}},+SPV_INTEL_global_variable_host_access"
 
 // CHECK-WITHOUT: "{{.*}}llvm-spirv"
 // CHECK-WITHOUT-NOT: --spirv-preserve-auxdata
diff --git a/clang/test/Driver/sycl-spirv-obj-old-model.cpp b/clang/test/Driver/sycl-spirv-obj-old-model.cpp
index 7a5f6215e31ae..6e9352b1cd4e3 100644
--- a/clang/test/Driver/sycl-spirv-obj-old-model.cpp
+++ b/clang/test/Driver/sycl-spirv-obj-old-model.cpp
@@ -11,7 +11,7 @@
 // SPIRV_DEVICE_OBJ-SAME: "-o" "[[DEVICE_BC:.+\.bc]]"
 // SPIRV_DEVICE_OBJ: llvm-spirv{{.*}} "-o" "[[DEVICE_SPV:.+\.spv]]"
 // SPIRV_DEVICE_OBJ-SAME: "--spirv-preserve-auxdata"
-// SPIRV_DEVICE_OBJ-SAME: "-spirv-ext=-all,{{.*}},+SPV_INTEL_memory_access_aliasing"
+// SPIRV_DEVICE_OBJ-SAME: "-spirv-ext=-all,{{.*}},+SPV_INTEL_global_variable_host_access"
 // SPIRV_DEVICE_OBJ-SAME: "[[DEVICE_BC]]"
 // SPIRV_DEVICE_OBJ: clang{{.*}} "-cc1" "-triple" "x86_64-unknown-linux-gnu"
 // SPIRV_DEVICE_OBJ-SAME: "-fsycl-is-host"
diff --git a/clang/test/Driver/sycl-spirv-obj.cpp b/clang/test/Driver/sycl-spirv-obj.cpp
index 975099db36413..a8afd57c55c10 100644
--- a/clang/test/Driver/sycl-spirv-obj.cpp
+++ b/clang/test/Driver/sycl-spirv-obj.cpp
@@ -11,7 +11,7 @@
 // SPIRV_DEVICE_OBJ-SAME: "-o" "[[DEVICE_BC:.+\.bc]]"
 // SPIRV_DEVICE_OBJ: llvm-spirv{{.*}} "-o" "[[DEVICE_SPV:.+\.spv]]"
 // SPIRV_DEVICE_OBJ-SAME: "--spirv-preserve-auxdata"
-// SPIRV_DEVICE_OBJ-SAME: "-spirv-ext=-all,{{.*}},+SPV_INTEL_memory_access_aliasing"
+// SPIRV_DEVICE_OBJ-SAME: "-spirv-ext=-all,{{.*}},+SPV_INTEL_global_variable_host_access"
 // SPIRV_DEVICE_OBJ-SAME: "[[DEVICE_BC]]"
 // SPIRV_DEVICE_OBJ: llvm-offload-binary{{.*}} "--image=file=[[DEVICE_SPV]]{{.*}}"
 // SPIRV_DEVICE_OBJ: clang{{.*}} "-cc1" "-triple" "x86_64-unknown-linux-gnu"
diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
index 4a41e1a610a1f..f094e70b642c6 100644
--- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
+++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
@@ -909,7 +909,8 @@ getTripleBasedSPIRVTransOpts(const ArgList &Args,
       ",+SPV_KHR_cooperative_matrix"
       ",+SPV_EXT_shader_atomic_float16_add"
       ",+SPV_INTEL_fp_max_error"
-      ",+SPV_INTEL_memory_access_aliasing";
+      ",+SPV_INTEL_memory_access_aliasing"
+      ",+SPV_INTEL_global_variable_host_access";
   TranslatorArgs.push_back(Args.MakeArgString(ExtArg));
 }
 
diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
index 48900f9db80d6..e438785965e7f 100644
--- a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
+++ b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
@@ -34,11 +34,9 @@ constexpr StringRef SpirvDecorMdKind = "spirv.Decorations";
 constexpr StringRef SpirvDecorCacheControlMdKind =
     "spirv.DecorationCacheControlINTEL";
 constexpr StringRef SpirvParamDecorMdKind = "spirv.ParameterDecorations";
-// The corresponding SPIR-V OpCode for the host_access property is documented
-// in the SPV_INTEL_global_variable_decorations design document:
-// https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/DeviceGlobal/SPV_INTEL_global_variable_decorations.asciidoc#decoration
-constexpr uint32_t SpirvHostAccessDecor = 6147;
-constexpr uint32_t SpirvHostAccessDefaultValue = 2; // Read/Write
+
+constexpr uint32_t SpirvHostAccessDecor = 6188;
+constexpr uint32_t SpirvHostAccessDefaultValue = 3; // Read/Write
 
 constexpr uint32_t SpirvInitiationIntervalDecor = 5917;
 constexpr uint32_t SpirvPipelineEnableDecor = 5919;
diff --git a/llvm/lib/SYCLLowerIR/SanitizerPostOptimizer.cpp b/llvm/lib/SYCLLowerIR/SanitizerPostOptimizer.cpp
index 5500f83dea8a7..601698bf0e618 100644
--- a/llvm/lib/SYCLLowerIR/SanitizerPostOptimizer.cpp
+++ b/llvm/lib/SYCLLowerIR/SanitizerPostOptimizer.cpp
@@ -24,7 +24,7 @@ using namespace llvm;
 namespace llvm {
 
 constexpr StringRef SPIRV_DECOR_MD_KIND = "spirv.Decorations";
-constexpr uint32_t SPIRV_HOST_ACCESS_DECOR = 6147;
+constexpr uint32_t SPIRV_HOST_ACCESS_DECOR = 6188;
 
 struct EliminateDeadCheck : public InstVisitor<EliminateDeadCheck> {
   void visitCallInst(CallInst &CI) {
diff --git a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp
index 7b9aa86dc9bbc..4c862f489bca1 100644
--- a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp
+++ b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp
@@ -1480,7 +1480,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", "0"); // read only
+  AsanSpirKernelMetadata->addAttribute("sycl-host-access", "1"); // read only
   AsanSpirKernelMetadata->addAttribute(
       "sycl-unique-id",
       computeKernelMetadataUniqueId("__AsanKernelMetadata", KernelNamesBytes));
diff --git a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp
index 3fa125d3c8ca9..018e9eb193f67 100644
--- a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp
+++ b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp
@@ -1296,7 +1296,7 @@ void MemorySanitizerOnSpirv::instrumentKernelsMetadata(int TrackOrigins) {
       "sycl-device-global-size", std::to_string(DL.getTypeAllocSize(ArrayTy)));
   MsanSpirKernelMetadata->addAttribute("sycl-device-image-scope");
   MsanSpirKernelMetadata->addAttribute("sycl-host-access",
-                                       "0"); // read only
+                                       "1"); // read only
   MsanSpirKernelMetadata->addAttribute(
       "sycl-unique-id",
       computeKernelMetadataUniqueId("__MsanKernelMetadata", KernelNamesBytes));
diff --git a/llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp b/llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp
index 9a4ddb1ef57a5..c1d4fee396405 100644
--- a/llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp
+++ b/llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp
@@ -732,7 +732,7 @@ void ThreadSanitizerOnSpirv::instrumentKernelsMetadata() {
   TsanSpirKernelMetadata->addAttribute(
       "sycl-device-global-size", std::to_string(DL.getTypeAllocSize(ArrayTy)));
   TsanSpirKernelMetadata->addAttribute("sycl-device-image-scope");
-  TsanSpirKernelMetadata->addAttribute("sycl-host-access", "0"); // read only
+  TsanSpirKernelMetadata->addAttribute("sycl-host-access", "1"); // read only
   TsanSpirKernelMetadata->addAttribute(
       "sycl-unique-id",
       computeKernelMetadataUniqueId("__TsanKernelMetadata", KernelNamesBytes));
diff --git a/llvm/test/Instrumentation/AddressSanitizer/SPIRV/extend_launch_info_arg.ll b/llvm/test/Instrumentation/AddressSanitizer/SPIRV/extend_launch_info_arg.ll
index 7f844a10ed49f..02a052c04e7fb 100644
--- a/llvm/test/Instrumentation/AddressSanitizer/SPIRV/extend_launch_info_arg.ll
+++ b/llvm/test/Instrumentation/AddressSanitizer/SPIRV/extend_launch_info_arg.ll
@@ -25,5 +25,5 @@ entry:
 
 attributes #0 = { sanitize_address }
 ;; sycl-device-global-size = 16 * 2
-;; sycl-host-access = 0 read-only
-; CHECK: attributes [[ATTR0]] = { "sycl-device-global-size"="48" "sycl-device-image-scope" "sycl-host-access"="0" "sycl-unique-id"="__AsanKernelMetadata833c47834a0b74946e370c23c39607cc" }
+;; sycl-host-access = 1 read-only
+; CHECK: attributes [[ATTR0]] = { "sycl-device-global-size"="48" "sycl-device-image-scope" "sycl-host-access"="1" "sycl-unique-id"="__AsanKernelMetadata833c47834a0b74946e370c23c39607cc" }
diff --git a/llvm/test/Instrumentation/MemorySanitizer/SPIRV/instrument_global_address_space.ll b/llvm/test/Instrumentation/MemorySanitizer/SPIRV/instrument_global_address_space.ll
index bf3496d6667ba..b50f50c24afa4 100644
--- a/llvm/test/Instrumentation/MemorySanitizer/SPIRV/instrument_global_address_space.ll
+++ b/llvm/test/Instrumentation/MemorySanitizer/SPIRV/instrument_global_address_space.ll
@@ -31,4 +31,4 @@ entry:
 }
 
 ; CHECK: attributes [[ATTR0]]
-; CHECK-SAME: "sycl-device-global-size"="24" "sycl-device-image-scope" "sycl-host-access"="0" "sycl-unique-id"="__MsanKernelMetadata3ff767e9a7a43f1f3968062dbb4ee3b4"
+; CHECK-SAME: "sycl-device-global-size"="24" "sycl-device-image-scope" "sycl-host-access"="1" "sycl-unique-id"="__MsanKernelMetadata3ff767e9a7a43f1f3968062dbb4ee3b4"
diff --git a/llvm/test/Instrumentation/MemorySanitizer/SPIRV/track_origins.ll b/llvm/test/Instrumentation/MemorySanitizer/SPIRV/track_origins.ll
index d20b133efff2a..9dcfd9ce3ad34 100644
--- a/llvm/test/Instrumentation/MemorySanitizer/SPIRV/track_origins.ll
+++ b/llvm/test/Instrumentation/MemorySanitizer/SPIRV/track_origins.ll
@@ -34,4 +34,4 @@ entry:
 }
 
 ; CHECK: attributes [[ATTR0]]
-; CHECK-SAME: "sycl-device-global-size"="24" "sycl-device-image-scope" "sycl-host-access"="0" "sycl-unique-id"="__MsanKernelMetadata3ff767e9a7a43f1f3968062dbb4ee3b4"
+; CHECK-SAME: "sycl-device-global-size"="24" "sycl-device-image-scope" "sycl-host-access"="1" "sycl-unique-id"="__MsanKernelMetadata3ff767e9a7a43f1f3968062dbb4ee3b4"
diff --git a/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/device-globals/test_global_variable.ll b/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/device-globals/test_global_variable.ll
index aaa5681ab83f1..988fe1f9e5c57 100644
--- a/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/device-globals/test_global_variable.ll
+++ b/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/device-globals/test_global_variable.ll
@@ -43,10 +43,10 @@ declare spir_func align 4 dereferenceable(4) ptr addrspace(4) @_ZNK2cl4sycl3ext6
 ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
 declare spir_func align 1 dereferenceable(1) ptr addrspace(4) @_ZNK2cl4sycl3ext6oneapi13device_globalIbJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(ptr addrspace(4) align 1 dereferenceable_or_null(1) %this) #4 align 2
 
-attributes #0 = { "sycl-unique-id"="6da74a122db9f35d____ZL7dg_int1" "sycl-device-image-scope"="false" "sycl-host-access"="1" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-device-global-size"="4" }
+attributes #0 = { "sycl-unique-id"="6da74a122db9f35d____ZL7dg_int1" "sycl-device-image-scope"="false" "sycl-host-access"="2" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-device-global-size"="4" }
 attributes #1 = { "sycl-unique-id"="7da74a1187b9f35d____ZL7dg_int2" "sycl-implement-in-csr"="false" "sycl-init-mode"="1" "sycl-device-global-size"="4" }
-attributes #2 = { "sycl-unique-id"="9d329ad59055e972____ZL8dg_bool3" "sycl-device-image-scope"="true" "sycl-host-access"="0" "sycl-implement-in-csr" "sycl-init-mode"="0" "sycl-device-global-size"="1" }
-attributes #3 = { "sycl-unique-id"="dda2bad52c45c432____ZL8dg_bool4" "sycl-device-image-scope" "sycl-host-access"="2" "sycl-device-global-size"="1" }
+attributes #2 = { "sycl-unique-id"="9d329ad59055e972____ZL8dg_bool3" "sycl-device-image-scope"="true" "sycl-host-access"="1" "sycl-implement-in-csr" "sycl-init-mode"="0" "sycl-device-global-size"="1" }
+attributes #3 = { "sycl-unique-id"="dda2bad52c45c432____ZL8dg_bool4" "sycl-device-image-scope" "sycl-host-access"="3" "sycl-device-global-size"="1" }
 attributes #4 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
 attributes #5 = { convergent nounwind }
 ; no sycl-device-global-size attribute, this is not a device global variable but it contains compile-time properties,
@@ -70,18 +70,18 @@ attributes #6 = { "sycl-unique-id"="6da74a122db9f35d____ZL7no_dg_int1" "sycl-dev
 ; CHECK-IR-DAG: ![[#MN0]] = !{![[#MN1:]], ![[#MN2:]], ![[#MN3:]]}
 ; CHECK-IR-DAG: ![[#MN1]] = !{i32 6149, i32 1}
 ; CHECK-IR-DAG: ![[#MN2]] = !{i32 6148, i32 0}
-; CHECK-IR-DAG: ![[#MN3]] = !{i32 6147, i32 1, !"6da74a122db9f35d____ZL7dg_int1"}
+; CHECK-IR-DAG: ![[#MN3]] = !{i32 6188, i32 2, !"6da74a122db9f35d____ZL7dg_int1"}
 
 ; CHECK-IR-DAG: ![[#MN4]] = !{![[#MN5:]], ![[#MN6:]], ![[#MN7:]]}
 ; CHECK-IR-DAG: ![[#MN5]] = !{i32 6149, i32 0}
 ; CHECK-IR-DAG: ![[#MN6]] = !{i32 6148, i32 1}
-; CHECK-IR-DAG: ![[#MN7]] = !{i32 6147, i32 2, !"7da74a1187b9f35d____ZL7dg_int2"}
+; CHECK-IR-DAG: ![[#MN7]] = !{i32 6188, i32 3, !"7da74a1187b9f35d____ZL7dg_int2"}
 
 ; CHECK-IR-DAG: ![[#MN8]] = !{![[#MN1]], ![[#MN2]], ![[#MN9:]]}
-; CHECK-IR-DAG: ![[#MN9]] = !{i32 6147, i32 0, !"9d329ad59055e972____ZL8dg_bool3"}
+; CHECK-IR-DAG: ![[#MN9]] = !{i32 6188, i32 1, !"9d329ad59055e972____ZL8dg_bool3"}
 
 ; CHECK-IR-DAG: ![[#MN10]] = !{![[#MN11:]]}
-; CHECK-IR-DAG: ![[#MN11]] = !{i32 6147, i32 2, !"dda2bad52c45c432____ZL8dg_bool4"}
+; CHECK-IR-DAG: ![[#MN11]] = !{i32 6188, i32 3, !"dda2bad52c45c432____ZL8dg_bool4"}
 
 ; For not a device global variable, only actually present compile-time
 ; properties are handled
diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable.ll
index cd07c948fb91d..05a1f277af94a 100644
--- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable.ll
+++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable.ll
@@ -34,11 +34,11 @@ declare spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext
 ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
 declare spir_func align 1 dereferenceable(1) i8 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIbJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global.1" addrspace(4)* align 1 dereferenceable_or_null(1)) #5 align 2
 
-attributes #0 = { "sycl-device-global-size"="4" "sycl-device-image-scope"="false" "sycl-host-access"="1" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-unique-id"="6da74a122db9f35d____ZL7dg_int1" }
-attributes #1 = { "sycl-device-global-size"="4" "sycl-implement-in-csr"="false" "sycl-init-mode"="1" "sycl-unique-id"="7da74a1187b9f35d____ZL7dg_int2" }
-attributes #2 = { "sycl-device-global-size"="1" "sycl-device-image-scope"="true" "sycl-host-access"="0" "sycl-implement-in-csr" "sycl-init-mode"="0" "sycl-unique-id"="9d329ad59055e972____ZL8dg_bool3" }
-attributes #3 = { "sycl-device-global-size"="1" "sycl-device-image-scope" "sycl-host-access"="2" "sycl-unique-id"="dda2bad52c45c432____ZL8dg_bool4" }
-attributes #4 = { "sycl-device-image-scope"="false" "sycl-host-access"="1" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-unique-id"="6da74a122db9f35d____ZL7no_dg_int1" }
+attributes #0 = { "sycl-device-global-size"="4" "sycl-device-image-scope"="false" "sycl-host-access"="2" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-unique-id"="6da74a122db9f35d____ZL7dg_int1" }
+attributes #1 = { "sycl-device-global-size"="4" "sycl-implement-in-csr"="false" "sycl-init-mode"="2" "sycl-unique-id"="7da74a1187b9f35d____ZL7dg_int2" }
+attributes #2 = { "sycl-device-global-size"="1" "sycl-device-image-scope"="true" "sycl-host-access"="1" "sycl-implement-in-csr" "sycl-init-mode"="0" "sycl-unique-id"="9d329ad59055e972____ZL8dg_bool3" }
+attributes #3 = { "sycl-device-global-size"="1" "sycl-device-image-scope" "sycl-host-access"="3" "sycl-unique-id"="dda2bad52c45c432____ZL8dg_bool4" }
+attributes #4 = { "sycl-device-image-scope"="false" "sycl-host-access"="2" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-unique-id"="6da74a122db9f35d____ZL7no_dg_int1" }
 attributes #5 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
 attributes #6 = { convergent nounwind }
 
@@ -51,15 +51,15 @@ attributes #6 = { convergent nounwind }
 !0 = !{!1, !2, !3}
 !1 = !{i32 6149, i32 1}
 !2 = !{i32 6148, i32 0}
-!3 = !{i32 6147, i32 1, !"6da74a122db9f35d____ZL7dg_int1"}
+!3 = !{i32 6188, i32 2, !"6da74a122db9f35d____ZL7dg_int1"}
 !4 = !{!5, !6, !7}
 !5 = !{i32 6149, i32 0}
 !6 = !{i32 6148, i32 1}
-!7 = !{i32 6147, i32 2, !"7da74a1187b9f35d____ZL7dg_int2"}
+!7 = !{i32 6188, i32 3, !"7da74a1187b9f35d____ZL7dg_int2"}
 !8 = !{!1, !2, !9}
-!9 = !{i32 6147, i32 0, !"9d329ad59055e972____ZL8dg_bool3"}
+!9 = !{i32 6188, i32 1, !"9d329ad59055e972____ZL8dg_bool3"}
 !10 = !{!11}
-!11 = !{i32 6147, i32 2, !"dda2bad52c45c432____ZL8dg_bool4"}
+!11 = !{i32 6188, i32 3, !"dda2bad52c45c432____ZL8dg_bool4"}
 !12 = !{!1, !2}
 !13 = !{!"libcpmt"}
 !14 = !{i32 1, !"wchar_size", i32 2}
diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used.ll
index 2e2c6127083e1..60a89af39b726 100644
--- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used.ll
+++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used.ll
@@ -48,11 +48,11 @@ declare spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext
 ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
 declare spir_func align 1 dereferenceable(1) i8 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIbJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global.1" addrspace(4)* align 1 dereferenceable_or_null(1)) #5 align 2
 
-attributes #0 = { "sycl-device-global-size"="4" "sycl-device-image-scope"="false" "sycl-host-access"="1" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-unique-id"="6da74a122db9f35d____ZL7dg_int1" }
-attributes #1 = { "sycl-device-global-size"="4" "sycl-implement-in-csr"="false" "sycl-init-mode"="1" "sycl-unique-id"="7da74a1187b9f35d____ZL7dg_int2" }
-attributes #2 = { "sycl-device-global-size"="1" "sycl-device-image-scope"="true" "sycl-host-access"="0" "sycl-implement-in-csr" "sycl-init-mode"="0" "sycl-unique-id"="9d329ad59055e972____ZL8dg_bool3" }
-attributes #3 = { "sycl-device-global-size"="1" "sycl-device-image-scope" "sycl-host-access"="2" "sycl-unique-id"="dda2bad52c45c432____ZL8dg_bool4" }
-attributes #4 = { "sycl-device-global-size"="4" "sycl-device-image-scope"="false" "sycl-host-access"="1" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-unique-id"="6da74a122db9f35d____ZL7no_dg_int1" }
+attributes #0 = { "sycl-device-global-size"="4" "sycl-device-image-scope"="false" "sycl-host-access"="2" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-unique-id"="6da74a122db9f35d____ZL7dg_int1" }
+attributes #1 = { "sycl-device-global-size"="4" "sycl-implement-in-csr"="false" "sycl-init-mode"="2" "sycl-unique-id"="7da74a1187b9f35d____ZL7dg_int2" }
+attributes #2 = { "sycl-device-global-size"="1" "sycl-device-image-scope"="true" "sycl-host-access"="1" "sycl-implement-in-csr" "sycl-init-mode"="0" "sycl-unique-id"="9d329ad59055e972____ZL8dg_bool3" }
+attributes #3 = { "sycl-device-global-size"="1" "sycl-device-image-scope" "sycl-host-access"="3" "sycl-unique-id"="dda2bad52c45c432____ZL8dg_bool4" }
+attributes #4 = { "sycl-device-global-size"="4" "sycl-device-image-scope"="false" "sycl-host-access"="2" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-unique-id"="6da74a122db9f35d____ZL7no_dg_int1" }
 attributes #5 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
 attributes #6 = { convergent nounwind }
 
@@ -65,16 +65,16 @@ attributes #6 = { convergent nounwind }
 !0 = !{!1, !2, !3}
 !1 = !{i32 6149, i32 1}
 !2 = !{i32 6148, i32 0}
-!3 = !{i32 6147, i32 1, !"6da74a122db9f35d____ZL7dg_int1"}
+!3 = !{i32 6188, i32 2, !"6da74a122db9f35d____ZL7dg_int1"}
 !4 = !{!5, !6, !7}
 !5 = !{i32 6149, i32 0}
 !6 = !{i32 6148, i32 1}
-!7 = !{i32 6147, i32 2, !"7da74a1187b9f35d____ZL7dg_int2"}
+!7 = !{i32 6188, i32 3, !"7da74a1187b9f35d____ZL7dg_int2"}
 !8 = !{!1, !2, !9}
-!9 = !{i32 6147, i32 0, !"9d329ad59055e972____ZL8dg_bool3"}
+!9 = !{i32 6188, i32 1, !"9d329ad59055e972____ZL8dg_bool3"}
 !10 = !{!11}
-!11 = !{i32 6147, i32 2, !"dda2bad52c45c432____ZL8dg_bool4"}
-!12 = !{i32 6147, i32 1, !"6da74a122db9f35d____ZL7no_dg_int1"}
+!11 = !{i32 6188, i32 3, !"dda2bad52c45c432____ZL8dg_bool4"}
+!12 = !{i32 6188, i32 2, !"6da74a122db9f35d____ZL7no_dg_int1"}
 !13 = !{!"libcpmt"}
 !14 = !{i32 1, !"wchar_size", i32 2}
 !15 = !{i32 7, !"frame-pointer", i32 2}
diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used_opaque_ptr.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used_opaque_ptr.ll
index 80905a39b9518..7728e394554f0 100644
--- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used_opaque_ptr.ll
+++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used_opaque_ptr.ll
@@ -48,11 +48,11 @@ declare spir_func align 4 dereferenceable(4) ptr addrspace(4) @_ZNK2cl4sycl3ext6
 ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
 declare spir_func align 1 dereferenceable(1) ptr addrspace(4) @_ZNK2cl4sycl3ext6oneapi13device_globalIbJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(ptr addrspace(4) align 1 dereferenceable_or_null(1)) #5 align 2
 
-attributes #0 = { "sycl-device-global-size"="4" "sycl-device-image-scope"="false" "sycl-host-access"="1" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-unique-id"="6da74a122db9f35d____ZL7dg_int1" }
-attributes #1 = { "sycl-device-global-size"="4" "sycl-implement-in-csr"="false" "sycl-init-mode"="1" "sycl-unique-id"="7da74a1187b9f35d____ZL7dg_int2" }
-attributes #2 = { "sycl-device-global-size"="1" "sycl-device-image-scope"="true" "sycl-host-access"="0" "sycl-implement-in-csr" "sycl-init-mode"="0" "sycl-unique-id"="9d329ad59055e972____ZL8dg_bool3" }
-attributes #3 = { "sycl-device-global-size"="1" "sycl-device-image-scope" "sycl-host-access"="2" "sycl-unique-id"="dda2bad52c45c432____ZL8dg_bool4" }
-attributes #4 = { "sycl-device-global-size"="4" "sycl-device-image-scope"="false" "sycl-host-access"="1" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-unique-id"="6da74a122db9f35d____ZL7no_dg_int1" }
+attributes #0 = { "sycl-device-global-size"="4" "sycl-device-image-scope"="false" "sycl-host-access"="2" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-unique-id"="6da74a122db9f35d____ZL7dg_int1" }
+attributes #1 = { "sycl-device-global-size"="4" "sycl-implement-in-csr"="false" "sycl-init-mode"="2" "sycl-unique-id"="7da74a1187b9f35d____ZL7dg_int2" }
+attributes #2 = { "sycl-device-global-size"="1" "sycl-device-image-scope"="true" "sycl-host-access"="1" "sycl-implement-in-csr" "sycl-init-mode"="0" "sycl-unique-id"="9d329ad59055e972____ZL8dg_bool3" }
+attributes #3 = { "sycl-device-global-size"="1" "sycl-device-image-scope" "sycl-host-access"="3" "sycl-unique-id"="dda2bad52c45c432____ZL8dg_bool4" }
+attributes #4 = { "sycl-device-global-size"="4" "sycl-device-image-scope"="false" "sycl-host-access"="2" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-unique-id"="6da74a122db9f35d____ZL7no_dg_int1" }
 attributes #5 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
 attributes #6 = { convergent nounwind }
 
@@ -65,16 +65,16 @@ attributes #6 = { convergent nounwind }
 !0 = !{!1, !2, !3}
 !1 = !{i32 6149, i32 1}
 !2 = !{i32 6148, i32 0}
-!3 = !{i32 6147, i32 1, !"6da74a122db9f35d____ZL7dg_int1"}
+!3 = !{i32 6188, i32 2, !"6da74a122db9f35d____ZL7dg_int1"}
 !4 = !{!5, !6, !7}
 !5 = !{i32 6149, i32 0}
 !6 = !{i32 6148, i32 1}
-!7 = !{i32 6147, i32 2, !"7da74a1187b9f35d____ZL7dg_int2"}
+!7 = !{i32 6188, i32 3, !"7da74a1187b9f35d____ZL7dg_int2"}
 !8 = !{!1, !2, !9}
-!9 = !{i32 6147, i32 0, !"9d329ad59055e972____ZL8dg_bool3"}
+!9 = !{i32 6188, i32 1, !"9d329ad59055e972____ZL8dg_bool3"}
 !10 = !{!11}
-!11 = !{i32 6147, i32 2, !"dda2bad52c45c432____ZL8dg_bool4"}
-!12 = !{i32 6147, i32 1, !"6da74a122db9f35d____ZL7no_dg_int1"}
+!11 = !{i32 6188, i32 3, !"dda2bad52c45c432____ZL8dg_bool4"}
+!12 = !{i32 6188, i32 2, !"6da74a122db9f35d____ZL7no_dg_int1"}
 !13 = !{!"libcpmt"}
 !14 = !{i32 1, !"wchar_size", i32 2}
 !15 = !{i32 7, !"frame-pointer", i32 2}
diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_img_scope.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_img_scope.ll
index 9934318e77793..86345299be632 100644
--- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_img_scope.ll
+++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_img_scope.ll
@@ -103,7 +103,7 @@ entry:
   ret i32 addrspace(4)* %val
 }
 
-attributes #0 = { "sycl-device-global-size"="4" "sycl-host-access"="1" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-unique-id"="dg_int2" }
+attributes #0 = { "sycl-device-global-size"="4" "sycl-host-access"="2" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-unique-id"="dg_int2" }
 attributes #1 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_main.cpp" "uniform-work-group-size"="true" }
 attributes #2 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_1.cpp" "uniform-work-group-size"="true" }
 attributes #3 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
@@ -120,7 +120,7 @@ attributes #6 = { convergent nounwind }
 !0 = !{!1, !2, !3}
 !1 = !{i32 6149, i32 1}
 !2 = !{i32 6148, i32 0}
-!3 = !{i32 6147, i32 1, !"dg_int2"}
+!3 = !{i32 6188, i32 2, !"dg_int2"}
 !4 = !{!"libcpmt"}
 !5 = !{i32 1, i32 2}
 !6 = !{i32 4, i32 100000}
@@ -128,6 +128,6 @@ attributes #6 = { convergent nounwind }
 !8 = !{i32 1, !"wchar_size", i32 2}
 !9 = !{i32 7, !"frame-pointer", i32 2}
 !10 = !{}
-; CHECK-MOD0: !{i32 6147, i32 1, !"dg_int2"}
-; CHECK-MOD1: !{i32 6147, i32 1, !"dg_int2"}
-; CHECK-MOD2-NOT: !{i32 6147, i32 1, !"dg_int2"}
+; CHECK-MOD0: !{i32 6188, i32 2, !"dg_int2"}
+; CHECK-MOD1: !{i32 6188, i32 2, !"dg_int2"}
+; CHECK-MOD2-NOT: !{i32 6188, i32 2, !"dg_int2"}
diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_one_var_error.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_one_var_error.ll
index a5de3dc6f4925..40d71563d5958 100644
--- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_one_var_error.ll
+++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_one_var_error.ll
@@ -112,7 +112,7 @@ entry:
   ret i32 addrspace(4)* %val
 }
 
-attributes #0 = { "sycl-unique-id"="dg_int2" "sycl-device-image-scope"="true" "sycl-host-access"="1" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-device-global-size"="4" }
+attributes #0 = { "sycl-unique-id"="dg_int2" "sycl-device-image-scope"="true" "sycl-host-access"="2" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-device-global-size"="4" }
 attributes #1 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
 attributes #2 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_1.cpp" "uniform-work-group-size"="true" }
 attributes #3 = { convergent mustprogress noinline norecurse optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="test_global_variable_2.cpp" "uniform-work-group-size"="true" }
diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_name_mapping_metadata.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_name_mapping_metadata.ll
index f3d8c6e0bff07..9ee57cf2e2337 100644
--- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_name_mapping_metadata.ll
+++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_name_mapping_metadata.ll
@@ -37,11 +37,11 @@ declare spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext
 ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
 declare spir_func align 1 dereferenceable(1) i8 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIbJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global.1" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #4 align 2
 
-attributes #0 = { "sycl-unique-id"="6da74a122db9f35d____ZL7dg_int1" "sycl-device-image-scope"="false" "sycl-host-access"="1" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-device-global-size"="4" }
-attributes #1 = { "sycl-unique-id"="7da74a1187b9f35d____ZL7dg_int2" "sycl-implement-in-csr"="false" "sycl-init-mode"="1" "sycl-device-global-size"="4" }
-attributes #2 = { "sycl-unique-id"="9d329ad59055e972____ZL8dg_bool3" "sycl-device-image-scope"="true" "sycl-host-access"="0" "sycl-implement-in-csr" "sycl-init-mode"="0" "sycl-device-global-size"="1" }
-attributes #3 = { "sycl-unique-id"="dda2bad52c45c432____ZL8dg_bool4" "sycl-device-image-scope" "sycl-host-access"="2" "sycl-device-global-size"="1" }
-attributes #4 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+attributes #0 = { "sycl-unique-id"="6da74a122db9f35d____ZL7dg_int1" "sycl-device-image-scope"="false" "sycl-host-access"="2" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-device-global-size"="4" }
+attributes #1 = { "sycl-unique-id"="7da74a1187b9f35d____ZL7dg_int2" "sycl-implement-in-csr"="false" "sycl-init-mode"="2" "sycl-device-global-size"="4" }
+attributes #2 = { "sycl-unique-id"="9d329ad59055e972____ZL8dg_bool3" "sycl-device-image-scope"="true" "sycl-host-access"="1" "sycl-implement-in-csr" "sycl-init-mode"="0" "sycl-device-global-size"="1" }
+attributes #3 = { "sycl-unique-id"="dda2bad52c45c432____ZL8dg_bool4" "sycl-device-image-scope" "sycl-host-access"="3" "sycl-device-global-size"="1" }
+attributes #4 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="1" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
 attributes #5 = { convergent nounwind }
 ; no sycl-device-global-size attribute, this is not a device global variable but it contains compile-time properties,
 ; a metadata node will be generated.
diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used.ll
index b4d3bf3f557e1..fa4a5200d3ba4 100644
--- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used.ll
+++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used.ll
@@ -33,11 +33,11 @@ attributes #6 = { convergent nounwind }
 !0 = !{!1, !2, !3}
 !1 = !{i32 6149, i32 1}
 !2 = !{i32 6148, i32 0}
-!3 = !{i32 6147, i32 1, !"6da74a122db9f35d____ZL7dg_int1"}
+!3 = !{i32 6188, i32 2, !"6da74a122db9f35d____ZL7dg_int1"}
 !4 = !{!5, !6, !7}
 !5 = !{i32 6149, i32 0}
 !6 = !{i32 6148, i32 1}
-!7 = !{i32 6147, i32 2, !"7da74a1187b9f35d____ZL7dg_int2"}
+!7 = !{i32 6188, i32 3, !"7da74a1187b9f35d____ZL7dg_int2"}
 !8 = !{!"libcpmt"}
 !9 = !{i32 1, !"wchar_size", i32 2}
 !10 = !{i32 7, !"frame-pointer", i32 2}
diff --git a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used_opaque_ptr.ll b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used_opaque_ptr.ll
index 771f8ab1bada8..8a6aab4751920 100644
--- a/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used_opaque_ptr.ll
+++ b/llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used_opaque_ptr.ll
@@ -19,8 +19,8 @@ target triple = "spir64-unknown-unknown"
 
 @_ZL16NotADeviceGlobal = internal addrspace(1) constant i8 zeroinitializer
 
-attributes #0 = { "sycl-device-global-size"="4" "sycl-device-image-scope"="false" "sycl-host-access"="1" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-unique-id"="6da74a122db9f35d____ZL7dg_int1" }
-attributes #1 = { "sycl-device-global-size"="4" "sycl-implement-in-csr"="false" "sycl-init-mode"="1" "sycl-unique-id"="7da74a1187b9f35d____ZL7dg_int2" }
+attributes #0 = { "sycl-device-global-size"="4" "sycl-device-image-scope"="false" "sycl-host-access"="2" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-unique-id"="6da74a122db9f35d____ZL7dg_int1" }
+attributes #1 = { "sycl-device-global-size"="4" "sycl-implement-in-csr"="false" "sycl-init-mode"="2" "sycl-unique-id"="7da74a1187b9f35d____ZL7dg_int2" }
 attributes #5 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
 attributes #6 = { convergent nounwind }
 
@@ -33,11 +33,11 @@ attributes #6 = { convergent nounwind }
 !0 = !{!1, !2, !3}
 !1 = !{i32 6149, i32 1}
 !2 = !{i32 6148, i32 0}
-!3 = !{i32 6147, i32 1, !"6da74a122db9f35d____ZL7dg_int1"}
+!3 = !{i32 6188, i32 2, !"6da74a122db9f35d____ZL7dg_int1"}
 !4 = !{!5, !6, !7}
 !5 = !{i32 6149, i32 0}
 !6 = !{i32 6148, i32 1}
-!7 = !{i32 6147, i32 2, !"7da74a1187b9f35d____ZL7dg_int2"}
+!7 = !{i32 6188, i32 3, !"7da74a1187b9f35d____ZL7dg_int2"}
 !8 = !{!"libcpmt"}
 !9 = !{i32 1, !"wchar_size", i32 2}
 !10 = !{i32 7, !"frame-pointer", i32 2}
diff --git a/llvm/test/tools/sycl-post-link/device-sanitizer/asan.ll b/llvm/test/tools/sycl-post-link/device-sanitizer/asan.ll
index f9ac67c26e93e..40df5ba8a3c8d 100644
--- a/llvm/test/tools/sycl-post-link/device-sanitizer/asan.ll
+++ b/llvm/test/tools/sycl-post-link/device-sanitizer/asan.ll
@@ -52,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"="0" "sycl-unique-id"="_Z20__AsanKernelMetadata" }
+attributes #2 = { "sycl-device-global-size"="16" "sycl-device-image-scope" "sycl-host-access"="1" "sycl-unique-id"="_Z20__AsanKernelMetadata" }
 
 !llvm.module.flags = !{!0, !1, !2}
 !opencl.spir.version = !{!3}
diff --git a/llvm/test/tools/sycl-post-link/device-sanitizer/msan.ll b/llvm/test/tools/sycl-post-link/device-sanitizer/msan.ll
index 763e49219a5ba..435adf88c4e4d 100644
--- a/llvm/test/tools/sycl-post-link/device-sanitizer/msan.ll
+++ b/llvm/test/tools/sycl-post-link/device-sanitizer/msan.ll
@@ -20,7 +20,7 @@ $_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8MyKernel = comdat any
 @__MsanKernelMetadata = appending dso_local local_unnamed_addr addrspace(1) global [1 x { i64, i64 }] [{ i64, i64 } { i64 ptrtoint (ptr addrspace(1) @__msan_kernel to i64), i64 54 }], !spirv.Decorations !9 #0
 ; CHECK-IR: @__MsanKernelMetadata = dso_local local_unnamed_addr addrspace(1) global %0 { {{.*}} }, !spirv.Decorations [[MD1:![0-9]+]] #{{.*}}
 ; CHECK-IR: [[MD1]] = !{[[MD2:![0-9]+]]}
-; CHECK-IR: [[MD2]] = !{i32 6147, i32 0, !"_Z20__SanitizerKernelMetadata"}
+; CHECK-IR: [[MD2]] = !{i32 6188, i32 0, !"_Z20__SanitizerKernelMetadata"}
 @__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
 @__msan_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"
 
@@ -69,7 +69,7 @@ declare i64 @__msan_get_shadow(i64, i32)
 declare void @__msan_maybe_warning_1(i8, i32, ptr addrspace(2), i32, ptr addrspace(2))
 declare void @__msan_maybe_warning_8(i8, i32, ptr addrspace(2), i32, ptr addrspace(2))
 
-attributes #0 = { "sycl-device-global-size"="16" "sycl-device-image-scope" "sycl-host-access"="0" "sycl-unique-id"="_Z20__MsanKernelMetadata" }
+attributes #0 = { "sycl-device-global-size"="16" "sycl-device-image-scope" "sycl-host-access"="1" "sycl-unique-id"="_Z20__MsanKernelMetadata" }
 attributes #1 = { mustprogress norecurse nounwind sanitize_memory uwtable "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="check_call.cpp" "sycl-single-task" "uniform-work-group-size"="true" }
 attributes #2 = { mustprogress noinline norecurse nounwind sanitize_memory uwtable "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
 
diff --git a/sycl-jit/jit-compiler/lib/translation/SPIRVLLVMTranslation.cpp b/sycl-jit/jit-compiler/lib/translation/SPIRVLLVMTranslation.cpp
index 86ecfc061afbd..8d84f468c3ad2 100644
--- a/sycl-jit/jit-compiler/lib/translation/SPIRVLLVMTranslation.cpp
+++ b/sycl-jit/jit-compiler/lib/translation/SPIRVLLVMTranslation.cpp
@@ -68,7 +68,8 @@ SPIRV::TranslatorOpts &SPIRVLLVMTranslator::translatorOpts() {
       SPIRV::ExtensionID::SPV_KHR_non_semantic_info,
       SPIRV::ExtensionID::SPV_KHR_cooperative_matrix,
       SPIRV::ExtensionID::SPV_EXT_shader_atomic_float16_add,
-      SPIRV::ExtensionID::SPV_INTEL_fp_max_error};
+      SPIRV::ExtensionID::SPV_INTEL_fp_max_error,
+      SPIRV::ExtensionID::SPV_INTEL_global_variable_host_access};
 
   static auto Opts = [&]() -> SPIRV::TranslatorOpts {
     // Options for translation between SPIR-V and LLVM IR.
diff --git a/sycl/include/sycl/ext/oneapi/device_global/properties.hpp b/sycl/include/sycl/ext/oneapi/device_global/properties.hpp
index abd5cd202a980..7ebf76e57b645 100644
--- a/sycl/include/sycl/ext/oneapi/device_global/properties.hpp
+++ b/sycl/include/sycl/ext/oneapi/device_global/properties.hpp
@@ -26,7 +26,14 @@ struct device_image_scope_key
   using value_t = property_value<device_image_scope_key>;
 };
 
-enum class host_access_enum : std::uint16_t { read, write, read_write, none };
+// Values of host_access_enum correspond to the access qualifiers in
+// SPV_INTEL_global_variable_host_access.
+enum class host_access_enum : std::uint16_t {
+  none = 0,
+  read = 1,
+  write = 2,
+  read_write = 3
+};
 
 struct host_access_key
     : detail::compile_time_property_key<detail::PropKind::HostAccess> {