Skip to content

[SYCL] Use global_var_host_access ext for device_global on SPIR-V #21351

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
Maetveis wants to merge 2 commits into
base: sycl
Choose a base branch
from
Open

[SYCL] Use global_var_host_access ext for device_global on SPIR-V #21351

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
44c5fa2
[SYCL] Use global_var_host_access ext for device_global on SPIR-V
Maetveis Feb 24, 2026
49ea3be
Apply code-review suggestions
Maetveis Feb 25, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions clang/lib/Driver/ToolChains/Clang.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10710,6 +10710,7 @@ static void getTripleBasedSPIRVTransOpts(Compilation &C,
",+SPV_INTEL_long_composites"
",+SPV_INTEL_arithmetic_fence"
",+SPV_INTEL_global_variable_decorations"
",+SPV_INTEL_global_variable_host_access"
",+SPV_INTEL_cache_controls"
",+SPV_INTEL_fpga_buffer_location"
",+SPV_INTEL_fpga_argument_interfaces"
Expand Down
1 change: 1 addition & 0 deletions clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -890,6 +890,7 @@ getTripleBasedSPIRVTransOpts(const ArgList &Args,
",+SPV_INTEL_long_composites"
",+SPV_INTEL_arithmetic_fence"
",+SPV_INTEL_global_variable_decorations"
",+SPV_INTEL_global_variable_host_access"
",+SPV_INTEL_cache_controls"
",+SPV_INTEL_fpga_buffer_location"
",+SPV_INTEL_fpga_argument_interfaces"
Expand Down
33 changes: 28 additions & 5 deletions llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -35,10 +35,10 @@ 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
// in the SPV_INTEL_global_variable_host_access extension:
// https://github.com/KhronosGroup/SPIRV-Registry/blob/main/extensions/INTEL/SPV_INTEL_global_variable_host_access.asciidoc#decoration
constexpr uint32_t SpirvHostAccessDecor = 6188;
constexpr uint32_t SpirvHostAccessDefaultValue = 3; // Read/Write

constexpr uint32_t SpirvInitiationIntervalDecor = 5917;
constexpr uint32_t SpirvPipelineEnableDecor = 5919;
Expand Down Expand Up @@ -559,6 +559,28 @@ void getUserListIgnoringCast(
}
}

// Map the "host access mode" enumeration values as defined by
// sycl_ext_oneapi_device_global and encoded in the "sycl-host-access" attribute
// to the SPIR-V decoration values as defined by
// SPV_INTEL_global_variable_host_access
uint32_t mapSYCLHostAccessToSPIRV(uint32_t HostAccess) {
// The mapping is as follows:
// 0 (read) -> 1 (ReadINTEL)
// 1 (write) -> 2 (WriteINTEL)
// 2 (read_write) -> 3 (ReadWriteINTEL)
// 3 (none) -> 0 (NoneINTEL)
if (HostAccess < 3)
return HostAccess + 1;

if (HostAccess == 3)
return 0;

// For values outside the defined range, keep input value to match previous
// behavior when the deprecated SPV_INTEL_global_variable_decorations
// extension was used.
return HostAccess;
}

} // anonymous namespace

PreservedAnalyses CompileTimePropertiesPass::run(Module &M,
Expand Down Expand Up @@ -589,7 +611,8 @@ PreservedAnalyses CompileTimePropertiesPass::run(Module &M,
if (isDeviceGlobalVariable(GV)) {
auto HostAccessDecorValue =
GV.hasAttribute(SyclHostAccessAttr)
? getAttributeAsInteger<uint32_t>(GV, SyclHostAccessAttr)
? mapSYCLHostAccessToSPIRV(
getAttributeAsInteger<uint32_t>(GV, SyclHostAccessAttr))
: SpirvHostAccessDefaultValue;
auto VarName = getGlobalVariableUniqueId(GV);
MDOps.push_back(buildSpirvDecorMetadata(Ctx, SpirvHostAccessDecor,
Expand Down
9 changes: 6 additions & 3 deletions llvm/lib/SYCLLowerIR/SanitizerPostOptimizer.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,7 +24,10 @@ using namespace llvm;
namespace llvm {

constexpr StringRef SPIRV_DECOR_MD_KIND = "spirv.Decorations";
constexpr uint32_t SPIRV_HOST_ACCESS_DECOR = 6147;

// Values as defined by SPV_INTEL_global_variable_host_access extension:
constexpr uint32_t SPIRV_HOST_ACCESS_DECOR = 6188;
constexpr uint32_t SPIRV_HOST_ACCESS_READ = 1; // ReadINTEL

struct EliminateDeadCheck : public InstVisitor<EliminateDeadCheck> {
void visitCallInst(CallInst &CI) {
Expand Down Expand Up @@ -102,8 +105,8 @@ static bool FixSanitizerKernelMetadata(Module &M) {
auto *Ty = Type::getInt32Ty(Ctx);
MD.push_back(ConstantAsMetadata::get(
Constant::getIntegerValue(Ty, APInt(32, SPIRV_HOST_ACCESS_DECOR))));
MD.push_back(
ConstantAsMetadata::get(Constant::getIntegerValue(Ty, APInt(32, 0))));
MD.push_back(ConstantAsMetadata::get(
Constant::getIntegerValue(Ty, APInt(32, SPIRV_HOST_ACCESS_READ))));
MD.push_back(MDString::get(Ctx, "_Z20__SanitizerKernelMetadata"));

MDOps.push_back(MDNode::get(Ctx, MD));
Expand Down
8 changes: 4 additions & 4 deletions llvm/test/SYCLLowerIR/CompileTimePropertiesPass/device-globals/test_global_variable.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
8 changes: 4 additions & 4 deletions llvm/test/tools/sycl-post-link/device-globals/test_global_variable.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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}
Expand Down
10 changes: 5 additions & 5 deletions llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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}
Expand Down
10 changes: 5 additions & 5 deletions llvm/test/tools/sycl-post-link/device-globals/test_global_variable_drop_used_opaque_ptr.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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}
Expand Down
8 changes: 4 additions & 4 deletions ...tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_img_scope.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -120,14 +120,14 @@ 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}
!7 = !{!"clang version 14.0.0"}
!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"}
4 changes: 2 additions & 2 deletions llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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}
Expand Down
4 changes: 2 additions & 2 deletions llvm/test/tools/sycl-post-link/device-globals/test_global_variable_trim_used_opaque_ptr.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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}
Expand Down
2 changes: 1 addition & 1 deletion llvm/test/tools/sycl-post-link/device-globals/test_no_metadata_generated.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -43,6 +43,6 @@ attributes #2 = { convergent nounwind }
!4 = !{i32 4, i32 100000}
!5 = !{!"clang version 14.0.0"}

; CHECK-NOT: ![[#MN1:]] = !{i32 6147
; CHECK-NOT: ![[#MN1:]] = !{i32 6188
; CHECK-NOT: ![[#MN2:]] = !{i32 6148
; CHECK-NOT: ![[#MN3:]] = !{i32 6149
2 changes: 1 addition & 1 deletion llvm/test/tools/sycl-post-link/device-sanitizer/msan.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 1, !"_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"

Expand Down
1 change: 1 addition & 0 deletions sycl-jit/jit-compiler/lib/translation/SPIRVLLVMTranslation.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -50,6 +50,7 @@ SPIRV::TranslatorOpts &SPIRVLLVMTranslator::translatorOpts() {
SPIRV::ExtensionID::SPV_INTEL_long_composites,
SPIRV::ExtensionID::SPV_INTEL_arithmetic_fence,
SPIRV::ExtensionID::SPV_INTEL_global_variable_decorations,
SPIRV::ExtensionID::SPV_INTEL_global_variable_host_access,
SPIRV::ExtensionID::SPV_INTEL_cache_controls,
SPIRV::ExtensionID::SPV_INTEL_fpga_buffer_location,
SPIRV::ExtensionID::SPV_INTEL_fpga_argument_interfaces,
Expand Down
Loading