diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
index 3f5a9ab07a3da..ec27d83b77c43 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -10769,8 +10769,7 @@ 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_global_variable_host_access";
+                            ",+SPV_INTEL_memory_access_aliasing";
   std::string KHRExtArg = ",+SPV_KHR_16bit_storage"
                           ",+SPV_KHR_cooperative_matrix"
                           ",+SPV_KHR_expect_assume"
@@ -10865,8 +10864,7 @@ 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_global_variable_host_access";
+      ",+SPV_INTEL_memory_access_aliasing";
 
   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 d0b1270d46e62..148f24548b3cd 100644
--- a/clang/test/Driver/sycl-spirv-backend.cpp
+++ b/clang/test/Driver/sycl-spirv-backend.cpp
@@ -23,7 +23,6 @@
 // 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 4fbb64f831f84..0b2bf88525934 100644
--- a/clang/test/Driver/sycl-spirv-ext-old-model.cpp
+++ b/clang/test/Driver/sycl-spirv-ext-old-model.cpp
@@ -50,4 +50,3 @@
 // CHECK-DEFAULT-SAME:,+SPV_EXT_shader_atomic_float16_add
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_fp_max_error
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_memory_access_aliasing
-// CHECK-DEFAULT-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 60253c7efaef0..2e38c9921dbcc 100644
--- a/clang/test/Driver/sycl-spirv-ext.cpp
+++ b/clang/test/Driver/sycl-spirv-ext.cpp
@@ -65,4 +65,3 @@
 // CHECK-DEFAULT-SAME:,+SPV_EXT_shader_atomic_float16_add
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_fp_max_error
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_memory_access_aliasing
-// CHECK-DEFAULT-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 935e11e9fa6c2..2e19dd9ed1dc4 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_global_variable_host_access"
+// CHECK-WITH-SAME: "-spirv-ext=-all,{{.*}},+SPV_INTEL_memory_access_aliasing"
 
 // 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 6e9352b1cd4e3..7a5f6215e31ae 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_global_variable_host_access"
+// SPIRV_DEVICE_OBJ-SAME: "-spirv-ext=-all,{{.*}},+SPV_INTEL_memory_access_aliasing"
 // 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 a8afd57c55c10..975099db36413 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_global_variable_host_access"
+// SPIRV_DEVICE_OBJ-SAME: "-spirv-ext=-all,{{.*}},+SPV_INTEL_memory_access_aliasing"
 // 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 cae212a91fecd..1cf035f7ae7d6 100644
--- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
+++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
@@ -1020,8 +1020,7 @@ 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_global_variable_host_access";
+      ",+SPV_INTEL_memory_access_aliasing";
   TranslatorArgs.push_back(Args.MakeArgString(ExtArg));
 }
 
diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
index 9aa7cceeaa46d..48c64e2d41976 100644
--- a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
+++ b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
@@ -34,9 +34,11 @@ constexpr StringRef SpirvDecorMdKind = "spirv.Decorations";
 constexpr StringRef SpirvDecorCacheControlMdKind =
     "spirv.DecorationCacheControlINTEL";
 constexpr StringRef SpirvParamDecorMdKind = "spirv.ParameterDecorations";
-
-constexpr uint32_t SpirvHostAccessDecor = 6188;
-constexpr uint32_t SpirvHostAccessDefaultValue = 3; // Read/Write
+// 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 SpirvInitiationIntervalDecor = 5917;
 constexpr uint32_t SpirvPipelineEnableDecor = 5919;
diff --git a/llvm/lib/SYCLLowerIR/SanitizerPostOptimizer.cpp b/llvm/lib/SYCLLowerIR/SanitizerPostOptimizer.cpp
index aeaa6181fff4d..75b74f1d7b7f5 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 = 6188;
+constexpr uint32_t SPIRV_HOST_ACCESS_DECOR = 6147;
 
 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 2d95278c02d0a..c1d1b8bfe9d3e 100644
--- a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp
+++ b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp
@@ -1484,7 +1484,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", "1"); // read only
+  AsanSpirKernelMetadata->addAttribute("sycl-host-access", "0"); // 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 5afd820bcd18d..f6d56304381e4 100644
--- a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp
+++ b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp
@@ -1307,7 +1307,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",
-                                       "1"); // read only
+                                       "0"); // 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 63e0bbd15926a..6eaf2519d44b8 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", "1"); // read only
+  TsanSpirKernelMetadata->addAttribute("sycl-host-access", "0"); // 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 02a052c04e7fb..7f844a10ed49f 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 = 1 read-only
-; CHECK: attributes [[ATTR0]] = { "sycl-device-global-size"="48" "sycl-device-image-scope" "sycl-host-access"="1" "sycl-unique-id"="__AsanKernelMetadata833c47834a0b74946e370c23c39607cc" }
+;; 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" }
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 b50f50c24afa4..bf3496d6667ba 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"="1" "sycl-unique-id"="__MsanKernelMetadata3ff767e9a7a43f1f3968062dbb4ee3b4"
+; CHECK-SAME: "sycl-device-global-size"="24" "sycl-device-image-scope" "sycl-host-access"="0" "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 9dcfd9ce3ad34..d20b133efff2a 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"="1" "sycl-unique-id"="__MsanKernelMetadata3ff767e9a7a43f1f3968062dbb4ee3b4"
+; CHECK-SAME: "sycl-device-global-size"="24" "sycl-device-image-scope" "sycl-host-access"="0" "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 988fe1f9e5c57..aaa5681ab83f1 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"="2" "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"="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"="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 #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 #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 6188, i32 2, !"6da74a122db9f35d____ZL7dg_int1"}
+; CHECK-IR-DAG: ![[#MN3]] = !{i32 6147, i32 1, !"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 6188, i32 3, !"7da74a1187b9f35d____ZL7dg_int2"}
+; CHECK-IR-DAG: ![[#MN7]] = !{i32 6147, i32 2, !"7da74a1187b9f35d____ZL7dg_int2"}
 
 ; CHECK-IR-DAG: ![[#MN8]] = !{![[#MN1]], ![[#MN2]], ![[#MN9:]]}
-; CHECK-IR-DAG: ![[#MN9]] = !{i32 6188, i32 1, !"9d329ad59055e972____ZL8dg_bool3"}
+; CHECK-IR-DAG: ![[#MN9]] = !{i32 6147, i32 0, !"9d329ad59055e972____ZL8dg_bool3"}
 
 ; CHECK-IR-DAG: ![[#MN10]] = !{![[#MN11:]]}
-; CHECK-IR-DAG: ![[#MN11]] = !{i32 6188, i32 3, !"dda2bad52c45c432____ZL8dg_bool4"}
+; CHECK-IR-DAG: ![[#MN11]] = !{i32 6147, i32 2, !"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 05a1f277af94a..cd07c948fb91d 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"="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 #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 #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 6188, i32 2, !"6da74a122db9f35d____ZL7dg_int1"}
+!3 = !{i32 6147, i32 1, !"6da74a122db9f35d____ZL7dg_int1"}
 !4 = !{!5, !6, !7}
 !5 = !{i32 6149, i32 0}
 !6 = !{i32 6148, i32 1}
-!7 = !{i32 6188, i32 3, !"7da74a1187b9f35d____ZL7dg_int2"}
+!7 = !{i32 6147, i32 2, !"7da74a1187b9f35d____ZL7dg_int2"}
 !8 = !{!1, !2, !9}
-!9 = !{i32 6188, i32 1, !"9d329ad59055e972____ZL8dg_bool3"}
+!9 = !{i32 6147, i32 0, !"9d329ad59055e972____ZL8dg_bool3"}
 !10 = !{!11}
-!11 = !{i32 6188, i32 3, !"dda2bad52c45c432____ZL8dg_bool4"}
+!11 = !{i32 6147, i32 2, !"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 60a89af39b726..2e2c6127083e1 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"="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 #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 #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 6188, i32 2, !"6da74a122db9f35d____ZL7dg_int1"}
+!3 = !{i32 6147, i32 1, !"6da74a122db9f35d____ZL7dg_int1"}
 !4 = !{!5, !6, !7}
 !5 = !{i32 6149, i32 0}
 !6 = !{i32 6148, i32 1}
-!7 = !{i32 6188, i32 3, !"7da74a1187b9f35d____ZL7dg_int2"}
+!7 = !{i32 6147, i32 2, !"7da74a1187b9f35d____ZL7dg_int2"}
 !8 = !{!1, !2, !9}
-!9 = !{i32 6188, i32 1, !"9d329ad59055e972____ZL8dg_bool3"}
+!9 = !{i32 6147, i32 0, !"9d329ad59055e972____ZL8dg_bool3"}
 !10 = !{!11}
-!11 = !{i32 6188, i32 3, !"dda2bad52c45c432____ZL8dg_bool4"}
-!12 = !{i32 6188, i32 2, !"6da74a122db9f35d____ZL7no_dg_int1"}
+!11 = !{i32 6147, i32 2, !"dda2bad52c45c432____ZL8dg_bool4"}
+!12 = !{i32 6147, i32 1, !"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 7728e394554f0..80905a39b9518 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"="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 #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 #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 6188, i32 2, !"6da74a122db9f35d____ZL7dg_int1"}
+!3 = !{i32 6147, i32 1, !"6da74a122db9f35d____ZL7dg_int1"}
 !4 = !{!5, !6, !7}
 !5 = !{i32 6149, i32 0}
 !6 = !{i32 6148, i32 1}
-!7 = !{i32 6188, i32 3, !"7da74a1187b9f35d____ZL7dg_int2"}
+!7 = !{i32 6147, i32 2, !"7da74a1187b9f35d____ZL7dg_int2"}
 !8 = !{!1, !2, !9}
-!9 = !{i32 6188, i32 1, !"9d329ad59055e972____ZL8dg_bool3"}
+!9 = !{i32 6147, i32 0, !"9d329ad59055e972____ZL8dg_bool3"}
 !10 = !{!11}
-!11 = !{i32 6188, i32 3, !"dda2bad52c45c432____ZL8dg_bool4"}
-!12 = !{i32 6188, i32 2, !"6da74a122db9f35d____ZL7no_dg_int1"}
+!11 = !{i32 6147, i32 2, !"dda2bad52c45c432____ZL8dg_bool4"}
+!12 = !{i32 6147, i32 1, !"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 86345299be632..9934318e77793 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"="2" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-unique-id"="dg_int2" }
+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 #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 6188, i32 2, !"dg_int2"}
+!3 = !{i32 6147, i32 1, !"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 6188, i32 2, !"dg_int2"}
-; CHECK-MOD1: !{i32 6188, i32 2, !"dg_int2"}
-; CHECK-MOD2-NOT: !{i32 6188, i32 2, !"dg_int2"}
+; 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"}
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 4ff458cf9626b..715f015018b34 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"="2" "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"="1" "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 9ee57cf2e2337..f3d8c6e0bff07 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"="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 #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 #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 fa4a5200d3ba4..b4d3bf3f557e1 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 6188, i32 2, !"6da74a122db9f35d____ZL7dg_int1"}
+!3 = !{i32 6147, i32 1, !"6da74a122db9f35d____ZL7dg_int1"}
 !4 = !{!5, !6, !7}
 !5 = !{i32 6149, i32 0}
 !6 = !{i32 6148, i32 1}
-!7 = !{i32 6188, i32 3, !"7da74a1187b9f35d____ZL7dg_int2"}
+!7 = !{i32 6147, i32 2, !"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 8a6aab4751920..771f8ab1bada8 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"="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 #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 #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 6188, i32 2, !"6da74a122db9f35d____ZL7dg_int1"}
+!3 = !{i32 6147, i32 1, !"6da74a122db9f35d____ZL7dg_int1"}
 !4 = !{!5, !6, !7}
 !5 = !{i32 6149, i32 0}
 !6 = !{i32 6148, i32 1}
-!7 = !{i32 6188, i32 3, !"7da74a1187b9f35d____ZL7dg_int2"}
+!7 = !{i32 6147, i32 2, !"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 40df5ba8a3c8d..f9ac67c26e93e 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"="1" "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}
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 435adf88c4e4d..763e49219a5ba 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 6188, i32 0, !"_Z20__SanitizerKernelMetadata"}
+; CHECK-IR: [[MD2]] = !{i32 6147, 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"="1" "sycl-unique-id"="_Z20__MsanKernelMetadata" }
+attributes #0 = { "sycl-device-global-size"="16" "sycl-device-image-scope" "sycl-host-access"="0" "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 74c532e9bae59..0b80bdd61dfea 100644
--- a/sycl-jit/jit-compiler/lib/translation/SPIRVLLVMTranslation.cpp
+++ b/sycl-jit/jit-compiler/lib/translation/SPIRVLLVMTranslation.cpp
@@ -67,8 +67,7 @@ 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_global_variable_host_access};
+      SPIRV::ExtensionID::SPV_INTEL_fp_max_error};
 
   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 7ebf76e57b645..abd5cd202a980 100644
--- a/sycl/include/sycl/ext/oneapi/device_global/properties.hpp
+++ b/sycl/include/sycl/ext/oneapi/device_global/properties.hpp
@@ -26,14 +26,7 @@ struct device_image_scope_key
   using value_t = property_value<device_image_scope_key>;
 };
 
-// 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
-};
+enum class host_access_enum : std::uint16_t { read, write, read_write, none };
 
 struct host_access_key
     : detail::compile_time_property_key<detail::PropKind::HostAccess> {