diff --git a/clang/lib/CodeGen/BackendUtil.cpp b/clang/lib/CodeGen/BackendUtil.cpp index 40093f9ff503d..173a3c4e3d8f7 100644 --- a/clang/lib/CodeGen/BackendUtil.cpp +++ b/clang/lib/CodeGen/BackendUtil.cpp @@ -46,6 +46,7 @@ #include "llvm/Passes/PassBuilder.h" #include "llvm/Passes/PassPlugin.h" #include "llvm/Passes/StandardInstrumentations.h" +#include "llvm/SYCLLowerIR/CompileTimePropertiesPass.h" #include "llvm/SYCLLowerIR/ESIMD/ESIMDVerifier.h" #include "llvm/SYCLLowerIR/LowerWGLocalMemory.h" #include "llvm/SYCLLowerIR/MutatePrintfAddrspace.h" @@ -1053,6 +1054,9 @@ void EmitAssemblyHelper::RunOptimizationPipeline( // Allocate static local memory in SYCL kernel scope for each allocation // call. MPM.addPass(SYCLLowerWGLocalMemoryPass()); + + // Process properties and annotations + MPM.addPass(CompileTimePropertiesPass()); } } diff --git a/clang/test/CodeGenSYCL/kernel-early-optimization-pipeline.cpp b/clang/test/CodeGenSYCL/kernel-early-optimization-pipeline.cpp index 0709f0dd43d11..17527b58e5a8e 100644 --- a/clang/test/CodeGenSYCL/kernel-early-optimization-pipeline.cpp +++ b/clang/test/CodeGenSYCL/kernel-early-optimization-pipeline.cpp @@ -9,3 +9,8 @@ // RUN: %clang_cc1 -O2 -fsycl-is-device -triple spir64-unknown-unknown %s -mdebug-pass Structure -emit-llvm -fno-sycl-early-optimizations -o /dev/null 2>&1 | FileCheck %s --check-prefix=CHECK-NEWPM-NOEARLYOPT // CHECK-NEWPM-NOEARLYOPT-NOT: ConstantMergePass // CHECK-NEWPM-NOEARLYOPT: SYCLMutatePrintfAddrspacePass + +// Checks that the compile time properties pass is added into the compilation pipeline +// +// RUN: %clang_cc1 -O2 -fsycl-is-device -triple spir64-unknown-unknown %s -mdebug-pass Structure -emit-llvm -o /dev/null 2>&1 | FileCheck %s --check-prefix=CHECK-COMPTIMEPROPS +// CHECK-COMPTIMEPROPS: Running pass: CompileTimePropertiesPass on [module] diff --git a/llvm/tools/sycl-post-link/CompileTimePropertiesPass.h b/llvm/include/llvm/SYCLLowerIR/CompileTimePropertiesPass.h similarity index 100% rename from llvm/tools/sycl-post-link/CompileTimePropertiesPass.h rename to llvm/include/llvm/SYCLLowerIR/CompileTimePropertiesPass.h diff --git a/llvm/tools/sycl-post-link/DeviceGlobals.h b/llvm/include/llvm/SYCLLowerIR/DeviceGlobals.h similarity index 100% rename from llvm/tools/sycl-post-link/DeviceGlobals.h rename to llvm/include/llvm/SYCLLowerIR/DeviceGlobals.h diff --git a/llvm/lib/Passes/PassBuilder.cpp b/llvm/lib/Passes/PassBuilder.cpp index 4f2e22393f921..2426981321fc7 100644 --- a/llvm/lib/Passes/PassBuilder.cpp +++ b/llvm/lib/Passes/PassBuilder.cpp @@ -80,6 +80,7 @@ #include "llvm/IR/SafepointIRVerifier.h" #include "llvm/IR/Verifier.h" #include "llvm/IRPrinter/IRPrintingPasses.h" +#include "llvm/SYCLLowerIR/CompileTimePropertiesPass.h" #include "llvm/SYCLLowerIR/ESIMD/ESIMDVerifier.h" #include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h" #include "llvm/SYCLLowerIR/LowerInvokeSimd.h" diff --git a/llvm/lib/Passes/PassRegistry.def b/llvm/lib/Passes/PassRegistry.def index 3210c5ab1bf13..ca96fd6550ce9 100644 --- a/llvm/lib/Passes/PassRegistry.def +++ b/llvm/lib/Passes/PassRegistry.def @@ -138,6 +138,7 @@ MODULE_PASS("deadargelim-sycl", DeadArgumentEliminationSYCLPass()) MODULE_PASS("sycllowerwglocalmemory", SYCLLowerWGLocalMemoryPass()) MODULE_PASS("lower-esimd-kernel-attrs", SYCLFixupESIMDKernelWrapperMDPass()) MODULE_PASS("sycl-propagate-aspects-usage", SYCLPropagateAspectsUsagePass()) +MODULE_PASS("compile-time-properties", CompileTimePropertiesPass()) #undef MODULE_PASS #ifndef MODULE_PASS_WITH_PARAMS diff --git a/llvm/lib/SYCLLowerIR/CMakeLists.txt b/llvm/lib/SYCLLowerIR/CMakeLists.txt index 79c0eff62f8bd..b694a82aa0641 100644 --- a/llvm/lib/SYCLLowerIR/CMakeLists.txt +++ b/llvm/lib/SYCLLowerIR/CMakeLists.txt @@ -52,6 +52,8 @@ add_llvm_component_library(LLVMSYCLLowerIR ESIMD/ESIMDVerifier.cpp ESIMD/LowerESIMD.cpp ESIMD/LowerESIMDKernelAttrs.cpp + CompileTimePropertiesPass.cpp + DeviceGlobals.cpp ESIMD/LowerESIMDVecArg.cpp ESIMD/LowerESIMDVLoadVStore.cpp ESIMD/LowerESIMDSlmReservation.cpp diff --git a/llvm/tools/sycl-post-link/CompileTimeProperties.def b/llvm/lib/SYCLLowerIR/CompileTimeProperties.def similarity index 100% rename from llvm/tools/sycl-post-link/CompileTimeProperties.def rename to llvm/lib/SYCLLowerIR/CompileTimeProperties.def diff --git a/llvm/tools/sycl-post-link/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp similarity index 99% rename from llvm/tools/sycl-post-link/CompileTimePropertiesPass.cpp rename to llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp index 27b778c8b4090..1d2fa554ed8ee 100644 --- a/llvm/tools/sycl-post-link/CompileTimePropertiesPass.cpp +++ b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp @@ -8,8 +8,8 @@ // See comments in the header. //===----------------------------------------------------------------------===// -#include "CompileTimePropertiesPass.h" -#include "DeviceGlobals.h" +#include "llvm/SYCLLowerIR/CompileTimePropertiesPass.h" +#include "llvm/SYCLLowerIR/DeviceGlobals.h" #include "llvm/ADT/APInt.h" #include "llvm/ADT/StringMap.h" diff --git a/llvm/tools/sycl-post-link/DeviceGlobals.cpp b/llvm/lib/SYCLLowerIR/DeviceGlobals.cpp similarity index 97% rename from llvm/tools/sycl-post-link/DeviceGlobals.cpp rename to llvm/lib/SYCLLowerIR/DeviceGlobals.cpp index 165d0f8ad8503..b28f0d63f8cbd 100644 --- a/llvm/tools/sycl-post-link/DeviceGlobals.cpp +++ b/llvm/lib/SYCLLowerIR/DeviceGlobals.cpp @@ -8,8 +8,8 @@ // See comments in the header. //===----------------------------------------------------------------------===// -#include "DeviceGlobals.h" -#include "CompileTimePropertiesPass.h" +#include "llvm/SYCLLowerIR/DeviceGlobals.h" +#include "llvm/SYCLLowerIR/CompileTimePropertiesPass.h" #include "llvm/ADT/STLExtras.h" #include "llvm/ADT/StringRef.h" diff --git a/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/device-globals/test_global_variable.ll b/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/device-globals/test_global_variable.ll new file mode 100644 index 0000000000000..261fa188bc4b9 --- /dev/null +++ b/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/device-globals/test_global_variable.ll @@ -0,0 +1,88 @@ +; RUN: opt -passes=compile-time-properties %s -S | FileCheck %s --check-prefix CHECK-IR + +; This test is intended to check that DeviceGlobalPass adds all the required +; metadata nodes to every device global variable as well as the required +; properties in the 'SYCL/device globals' property set and handles the +; 'sycl-device-image-scope' attribute written in any allowed form. + +source_filename = "test_global_variable.cpp" +target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64" +target triple = "spir64-unknown-unknown" + +%"class.cl::sycl::ext::oneapi::device_global.0" = type { i32 addrspace(4)* } +%"class.cl::sycl::ext::oneapi::device_global.1" = type { i8 } +%class.anon.0 = type { i8 } + +@_ZL7dg_int1 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8 #0 +; CHECK-IR: @_ZL7dg_int1 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8, !spirv.Decorations ![[#MN0:]] +@_ZL7dg_int2 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8 #1 +; CHECK-IR: @_ZL7dg_int2 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8, !spirv.Decorations ![[#MN4:]] +@_ZL8dg_bool3 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.1" zeroinitializer, align 1 #2 +; CHECK-IR: @_ZL8dg_bool3 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.1" zeroinitializer, align 1, !spirv.Decorations ![[#MN8:]] +@_ZL8dg_bool4 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.1" zeroinitializer, align 1 #3 +; CHECK-IR: @_ZL8dg_bool4 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.1" zeroinitializer, align 1, !spirv.Decorations ![[#MN10:]] +@_ZL7no_dg_int1 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8 #6 +; CHECK-IR: @_ZL7no_dg_int1 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8, !spirv.Decorations ![[#MN12:]] + +define internal spir_func void @_ZZ4mainENKUlvE_clEv(%class.anon.0 addrspace(4)* align 1 dereferenceable_or_null(1) %this) #4 align 2 { +entry: + %this.addr = alloca %class.anon.0 addrspace(4)*, align 8 + %this.addr.ascast = addrspacecast %class.anon.0 addrspace(4)** %this.addr to %class.anon.0 addrspace(4)* addrspace(4)* + store %class.anon.0 addrspace(4)* %this, %class.anon.0 addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %this1 = load %class.anon.0 addrspace(4)*, %class.anon.0 addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 + %call1 = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_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.0" addrspace(4)* align 8 dereferenceable_or_null(8) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(1)* @_ZL7dg_int1 to %"class.cl::sycl::ext::oneapi::device_global.0" addrspace(4)*)) #5 + %call2 = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_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.0" addrspace(4)* align 8 dereferenceable_or_null(8) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(1)* @_ZL7dg_int2 to %"class.cl::sycl::ext::oneapi::device_global.0" addrspace(4)*)) #5 + %call3 = call 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) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.1" addrspace(1)* @_ZL8dg_bool3 to %"class.cl::sycl::ext::oneapi::device_global.1" addrspace(4)*)) #5 + %call4 = call 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) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.1" addrspace(1)* @_ZL8dg_bool4 to %"class.cl::sycl::ext::oneapi::device_global.1" addrspace(4)*)) #5 + ret void +} + +; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +declare spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_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.0" addrspace(4)* align 8 dereferenceable_or_null(8) %this) #4 align 2 + +; 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 #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. +attributes #6 = { "sycl-unique-id"="6da74a122db9f35d____ZL7no_dg_int1" "sycl-device-image-scope"="false" "sycl-host-access"="1" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" } + +!llvm.dependent-libraries = !{!0} +!llvm.module.flags = !{!1, !2} +!opencl.spir.version = !{!3} +!spirv.Source = !{!4} +!llvm.ident = !{!5} + +!0 = !{!"libcpmt"} +!1 = !{i32 1, !"wchar_size", i32 2} +!2 = !{i32 7, !"frame-pointer", i32 2} +!3 = !{i32 1, i32 2} +!4 = !{i32 4, i32 100000} +!5 = !{!"clang version 14.0.0"} + +; Ensure that the generated metadata nodes are correct +; 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: ![[#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: ![[#MN8]] = !{![[#MN1]], ![[#MN2]], ![[#MN9:]]} +; CHECK-IR-DAG: ![[#MN9]] = !{i32 6147, i32 0, !"9d329ad59055e972____ZL8dg_bool3"} + +; CHECK-IR-DAG: ![[#MN10]] = !{![[#MN11:]]} +; 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 +; CHECK-IR-DAG: ![[#MN12]] = !{![[#MN1]], ![[#MN2]]} diff --git a/llvm/test/tools/sycl-post-link/kernel-attributes/kernel-pipelined.ll b/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/kernel-attributes/kernel-pipelined.ll similarity index 94% rename from llvm/test/tools/sycl-post-link/kernel-attributes/kernel-pipelined.ll rename to llvm/test/SYCLLowerIR/CompileTimePropertiesPass/kernel-attributes/kernel-pipelined.ll index 131bbfc270b37..97dc178aa2d9a 100644 --- a/llvm/test/tools/sycl-post-link/kernel-attributes/kernel-pipelined.ll +++ b/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/kernel-attributes/kernel-pipelined.ll @@ -1,5 +1,5 @@ ; Check conversion of sycl-pipelined attribute -; RUN: sycl-post-link --device-globals --ir-output-only %s -S -o - | FileCheck %s --check-prefix CHECK-IR +; RUN: opt -passes="compile-time-properties" %s -S -o - | FileCheck %s --check-prefix CHECK-IR ; CHECK-IR-DAG: @pipelineNegative() #0 {{.*}}!spirv.Decorations [[DEFAULT_PIPELINE:![0-9]+]] { ; Function Attrs: convergent norecurse diff --git a/llvm/test/tools/sycl-post-link/kernel-attributes/register-map-interface.ll b/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/kernel-attributes/register-map-interface.ll similarity index 92% rename from llvm/test/tools/sycl-post-link/kernel-attributes/register-map-interface.ll rename to llvm/test/SYCLLowerIR/CompileTimePropertiesPass/kernel-attributes/register-map-interface.ll index c9bda1e5e5347..7512baf1b23e9 100644 --- a/llvm/test/tools/sycl-post-link/kernel-attributes/register-map-interface.ll +++ b/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/kernel-attributes/register-map-interface.ll @@ -1,5 +1,5 @@ ; Check conversion of sycl-register-map-interface attribute -; RUN: sycl-post-link --device-globals --ir-output-only %s -S -o - | FileCheck %s --check-prefix CHECK-IR +; RUN: opt -passes="compile-time-properties" %s -S -o - | FileCheck %s --check-prefix CHECK-IR ; CHECK-IR-DAG: @pStreaming() #0 {{.*}}!ip_interface [[REGISTER_MAP:![0-9]+]] { ; Function Attrs: convergent norecurse diff --git a/llvm/test/tools/sycl-post-link/kernel-attributes/streaming-interface.ll b/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/kernel-attributes/streaming-interface.ll similarity index 92% rename from llvm/test/tools/sycl-post-link/kernel-attributes/streaming-interface.ll rename to llvm/test/SYCLLowerIR/CompileTimePropertiesPass/kernel-attributes/streaming-interface.ll index 3a1c5b828996d..052d6618ec754 100644 --- a/llvm/test/tools/sycl-post-link/kernel-attributes/streaming-interface.ll +++ b/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/kernel-attributes/streaming-interface.ll @@ -1,5 +1,5 @@ ; Check conversion of sycl-streaming-interface attribute -; RUN: sycl-post-link --device-globals --ir-output-only %s -S -o - | FileCheck %s --check-prefix CHECK-IR +; RUN: opt -passes="compile-time-properties" %s -S -o - | FileCheck %s --check-prefix CHECK-IR ; CHECK-IR-DAG: @pStreaming() #0 {{.*}}!ip_interface [[STREAMING:![0-9]+]] { ; Function Attrs: convergent norecurse diff --git a/llvm/test/tools/sycl-post-link/kernel-properties.ll b/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/kernel-properties.ll similarity index 96% rename from llvm/test/tools/sycl-post-link/kernel-properties.ll rename to llvm/test/SYCLLowerIR/CompileTimePropertiesPass/kernel-properties.ll index 02f3357607845..39e40a21c4e5f 100644 --- a/llvm/test/tools/sycl-post-link/kernel-properties.ll +++ b/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/kernel-properties.ll @@ -1,4 +1,4 @@ -; RUN: sycl-post-link --ir-output-only --device-globals %s -S -o - | FileCheck %s --check-prefix CHECK-IR +; RUN: opt -passes=compile-time-properties %s -S | FileCheck %s --check-prefix CHECK-IR ; CHECK-IR-DAG: @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE9TheKernel0"() #0 {{.*}}!intel_reqd_sub_group_size ![[SGSizeMD0:[0-9]+]] {{.*}}!reqd_work_group_size ![[WGSizeMD0:[0-9]+]]{{.*}}!work_group_size_hint ![[WGSizeHintMD0:[0-9]+]] ; Function Attrs: convergent norecurse diff --git a/llvm/test/tools/sycl-post-link/sycl-properties-ptr-annotations.ll b/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/sycl-properties-ptr-annotations.ll similarity index 94% rename from llvm/test/tools/sycl-post-link/sycl-properties-ptr-annotations.ll rename to llvm/test/SYCLLowerIR/CompileTimePropertiesPass/sycl-properties-ptr-annotations.ll index 330c5fc01971d..4418fc334e34f 100644 --- a/llvm/test/tools/sycl-post-link/sycl-properties-ptr-annotations.ll +++ b/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/sycl-properties-ptr-annotations.ll @@ -1,15 +1,4 @@ -; RUN: sycl-post-link --device-globals --ir-output-only -S %s -o %t.ll -; RUN: FileCheck %s -input-file=%t.ll -; -; TODO: Remove --device-globals once other features start using compile-time -; properties. -; -; Tests the translation of "sycl-properties" pointer annotations to pointer -; annotations the SPIR-V translator will produce decorations from. -; NOTE: These use SYCL property meta-names that are currently only intended for -; use in attributes-to-metadata translations, but sycl-post-link does not -; currently make the distinction so we will use them for the purpose of -; testing the transformations. +; RUN: opt -passes=compile-time-properties %s -S | FileCheck %s target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64" target triple = "spir64-unknown-unknown" 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 df5461d9be2ef..d33088898c52e 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 @@ -1,13 +1,5 @@ ; RUN: sycl-post-link --device-globals -S %s -o %t.files.table ; RUN: FileCheck %s -input-file=%t.files_0.prop --check-prefix CHECK-PROP -; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-IR -; RUN: sycl-post-link --ir-output-only --device-globals %s -S -o - | FileCheck %s --check-prefix CHECK-IR - -; This test is intended to check that DeviceGlobalPass adds all the required -; metadata nodes to every device global variable as well as the required -; properties in the 'SYCL/device globals' property set and handles the -; 'sycl-device-image-scope' attribute written in any allowed form. - source_filename = "test_global_variable.cpp" target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64" target triple = "spir64-unknown-unknown" @@ -16,79 +8,65 @@ target triple = "spir64-unknown-unknown" %"class.cl::sycl::ext::oneapi::device_global.1" = type { i8 } %class.anon.0 = type { i8 } -@_ZL7dg_int1 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8 #0 -; CHECK-IR: @_ZL7dg_int1 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8, !spirv.Decorations ![[#MN0:]] -@_ZL7dg_int2 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8 #1 -; CHECK-IR: @_ZL7dg_int2 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8, !spirv.Decorations ![[#MN4:]] -@_ZL8dg_bool3 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.1" zeroinitializer, align 1 #2 -; CHECK-IR: @_ZL8dg_bool3 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.1" zeroinitializer, align 1, !spirv.Decorations ![[#MN8:]] -@_ZL8dg_bool4 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.1" zeroinitializer, align 1 #3 -; CHECK-IR: @_ZL8dg_bool4 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.1" zeroinitializer, align 1, !spirv.Decorations ![[#MN10:]] -@_ZL7no_dg_int1 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8 #6 -; CHECK-IR: @_ZL7no_dg_int1 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8, !spirv.Decorations ![[#MN12:]] +@_ZL7dg_int1 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8, !spirv.Decorations !0 #0 +@_ZL7dg_int2 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8, !spirv.Decorations !4 #1 +@_ZL8dg_bool3 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.1" zeroinitializer, align 1, !spirv.Decorations !8 #2 +@_ZL8dg_bool4 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.1" zeroinitializer, align 1, !spirv.Decorations !10 #3 +@_ZL7no_dg_int1 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8, !spirv.Decorations !12 #4 -define internal spir_func void @_ZZ4mainENKUlvE_clEv(%class.anon.0 addrspace(4)* align 1 dereferenceable_or_null(1) %this) #4 align 2 { +; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +define internal spir_func void @_ZZ4mainENKUlvE_clEv(%class.anon.0 addrspace(4)* align 1 dereferenceable_or_null(1) %this) #5 align 2 { entry: %this.addr = alloca %class.anon.0 addrspace(4)*, align 8 %this.addr.ascast = addrspacecast %class.anon.0 addrspace(4)** %this.addr to %class.anon.0 addrspace(4)* addrspace(4)* store %class.anon.0 addrspace(4)* %this, %class.anon.0 addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 %this1 = load %class.anon.0 addrspace(4)*, %class.anon.0 addrspace(4)* addrspace(4)* %this.addr.ascast, align 8 - %call1 = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_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.0" addrspace(4)* align 8 dereferenceable_or_null(8) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(1)* @_ZL7dg_int1 to %"class.cl::sycl::ext::oneapi::device_global.0" addrspace(4)*)) #5 - %call2 = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_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.0" addrspace(4)* align 8 dereferenceable_or_null(8) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(1)* @_ZL7dg_int2 to %"class.cl::sycl::ext::oneapi::device_global.0" addrspace(4)*)) #5 - %call3 = call 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) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.1" addrspace(1)* @_ZL8dg_bool3 to %"class.cl::sycl::ext::oneapi::device_global.1" addrspace(4)*)) #5 - %call4 = call 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) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.1" addrspace(1)* @_ZL8dg_bool4 to %"class.cl::sycl::ext::oneapi::device_global.1" addrspace(4)*)) #5 + %call1 = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_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.0" addrspace(4)* align 8 dereferenceable_or_null(8) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(1)* @_ZL7dg_int1 to %"class.cl::sycl::ext::oneapi::device_global.0" addrspace(4)*)) #6 + %call2 = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_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.0" addrspace(4)* align 8 dereferenceable_or_null(8) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(1)* @_ZL7dg_int2 to %"class.cl::sycl::ext::oneapi::device_global.0" addrspace(4)*)) #6 + %call3 = call 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) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.1" addrspace(1)* @_ZL8dg_bool3 to %"class.cl::sycl::ext::oneapi::device_global.1" addrspace(4)*)) #6 + %call4 = call 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) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.1" addrspace(1)* @_ZL8dg_bool4 to %"class.cl::sycl::ext::oneapi::device_global.1" addrspace(4)*)) #6 ret void } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -declare spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_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.0" addrspace(4)* align 8 dereferenceable_or_null(8) %this) #4 align 2 +declare spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_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.0" addrspace(4)* align 8 dereferenceable_or_null(8)) #5 align 2 ; 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 #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. -attributes #6 = { "sycl-unique-id"="6da74a122db9f35d____ZL7no_dg_int1" "sycl-device-image-scope"="false" "sycl-host-access"="1" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" } - -!llvm.dependent-libraries = !{!0} -!llvm.module.flags = !{!1, !2} -!opencl.spir.version = !{!3} -!spirv.Source = !{!4} -!llvm.ident = !{!5} - -!0 = !{!"libcpmt"} -!1 = !{i32 1, !"wchar_size", i32 2} -!2 = !{i32 7, !"frame-pointer", i32 2} -!3 = !{i32 1, i32 2} -!4 = !{i32 4, i32 100000} -!5 = !{!"clang version 14.0.0"} - -; Ensure that the generated metadata nodes are correct -; 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: ![[#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: ![[#MN8]] = !{![[#MN1]], ![[#MN2]], ![[#MN9:]]} -; CHECK-IR-DAG: ![[#MN9]] = !{i32 6147, i32 0, !"9d329ad59055e972____ZL8dg_bool3"} - -; CHECK-IR-DAG: ![[#MN10]] = !{![[#MN11:]]} -; 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 -; CHECK-IR-DAG: ![[#MN12]] = !{![[#MN1]], ![[#MN2]]} +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 #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 } + +!llvm.dependent-libraries = !{!13} +!llvm.module.flags = !{!14, !15} +!opencl.spir.version = !{!16} +!spirv.Source = !{!17} +!llvm.ident = !{!18} + +!0 = !{!1, !2, !3} +!1 = !{i32 6149, i32 1} +!2 = !{i32 6148, i32 0} +!3 = !{i32 6147, i32 1, !"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"} +!8 = !{!1, !2, !9} +!9 = !{i32 6147, i32 0, !"9d329ad59055e972____ZL8dg_bool3"} +!10 = !{!11} +!11 = !{i32 6147, i32 2, !"dda2bad52c45c432____ZL8dg_bool4"} +!12 = !{!1, !2} +!13 = !{!"libcpmt"} +!14 = !{i32 1, !"wchar_size", i32 2} +!15 = !{i32 7, !"frame-pointer", i32 2} +!16 = !{i32 1, i32 2} +!17 = !{i32 4, i32 100000} +!18 = !{!"clang version 14.0.0"} ; Ensure that the default values are correct. ; ABAAAAAAAAABAAAAAxxxxx is decoded to 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 a8627d63fe19a..5a19bd505f9fc 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 @@ -3,34 +3,35 @@ ; RUN: FileCheck %s -input-file=%t.files_1.ll --check-prefix CHECK-MOD1 ; RUN: FileCheck %s -input-file=%t.files_2.ll --check-prefix CHECK-MOD2 -; This test is intended to check that sycl-post-link generates no errors even -; when a single device global variable but without the 'device_image_scope' -; property is used from multiple device images. - +; ModuleID = 'llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_img_scope.ll' +source_filename = "llvm/test/tools/sycl-post-link/device-globals/test_global_variable_many_modules_no_dev_img_scope.ll" target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64" target triple = "spir64-unknown-unknown" %"class.cl::sycl::ext::oneapi::device_global" = type { i32 } %"class.cl::sycl::detail::accessor_common" = type { i8 } +$_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_ = comdat any + $_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_ = comdat any + $_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_ = comdat any -$_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_ = comdat any $dg_int2 = comdat any -@dg_int2 = linkonce_odr dso_local addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, comdat, align 4 #0 + +@dg_int2 = linkonce_odr dso_local addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, comdat, align 4, !spirv.Decorations !0 #0 ; CHECK-MOD0: @dg_int2 = linkonce_odr dso_local addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, comdat, align 4 ; CHECK-MOD1: @dg_int2 = linkonce_odr dso_local addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global" zeroinitializer, comdat, align 4 ; CHECK-MOD2-NOT: @dg_int2 -; Third kernel that uses no device-global variables -define weak_odr dso_local spir_kernel void @_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_() #4 comdat !kernel_arg_buffer_location !6 { +; Function Attrs: convergent mustprogress noinline norecurse optnone +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel3RN2cl4sycl5queueEEUlvE_() #1 comdat !kernel_arg_buffer_location !10 { entry: ret void } ; Function Attrs: convergent mustprogress noinline norecurse optnone -define weak_odr dso_local spir_kernel void @_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_() #2 comdat !kernel_arg_buffer_location !6 { +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel1RN2cl4sycl5queueEEUlvE_() #2 comdat !kernel_arg_buffer_location !10 { entry: %0 = alloca %"class.cl::sycl::detail::accessor_common", align 1 %1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)* @@ -39,7 +40,7 @@ entry: } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define internal spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #1 align 2 { +define internal spir_func void @_ZZ7kernel1RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #3 align 2 { entry: %this.addr = alloca %"class.cl::sycl::detail::accessor_common" addrspace(4)*, align 8 %this.addr.ascast = addrspacecast %"class.cl::sycl::detail::accessor_common" addrspace(4)** %this.addr to %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* @@ -50,14 +51,14 @@ entry: } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define dso_local spir_func void @_Z14kernel1_level1v() #1 { +define dso_local spir_func void @_Z14kernel1_level1v() #3 { entry: %call = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global" addrspace(1)* @dg_int2 to %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*)) #6 ret void } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #1 align 2 { +define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #3 align 2 { entry: %retval = alloca i32 addrspace(4)*, align 8 %this.addr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8 @@ -70,7 +71,7 @@ entry: } ; Function Attrs: convergent mustprogress noinline norecurse optnone -define weak_odr dso_local spir_kernel void @_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_() #3 comdat !kernel_arg_buffer_location !6 { +define weak_odr dso_local spir_kernel void @_ZTSZ7kernel2RN2cl4sycl5queueEEUlvE_() #4 comdat !kernel_arg_buffer_location !10 { entry: %0 = alloca %"class.cl::sycl::detail::accessor_common", align 1 %1 = addrspacecast %"class.cl::sycl::detail::accessor_common"* %0 to %"class.cl::sycl::detail::accessor_common" addrspace(4)* @@ -79,7 +80,7 @@ entry: } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define internal spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #1 align 2 { +define internal spir_func void @_ZZ7kernel2RN2cl4sycl5queueEENKUlvE_clEv(%"class.cl::sycl::detail::accessor_common" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #3 align 2 { entry: %this.addr = alloca %"class.cl::sycl::detail::accessor_common" addrspace(4)*, align 8 %this.addr.ascast = addrspacecast %"class.cl::sycl::detail::accessor_common" addrspace(4)** %this.addr to %"class.cl::sycl::detail::accessor_common" addrspace(4)* addrspace(4)* @@ -90,7 +91,7 @@ entry: } ; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #1 align 2 { +define internal spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIPKcXadsoS5_L_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_I11host_accessXadsoS5_L_ZL5Name2EEELS8_1EEENS4_IS6_XadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IS6_XadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv.2(%"class.cl::sycl::ext::oneapi::device_global" addrspace(4)* align 4 dereferenceable_or_null(4) %this) #3 align 2 { entry: %retval = alloca i32 addrspace(4)*, align 8 %this.addr = alloca %"class.cl::sycl::ext::oneapi::device_global" addrspace(4)*, align 8 @@ -102,28 +103,31 @@ entry: ret i32 addrspace(4)* %val } -; This device_global variable has no "device_image_scope" property -attributes #0 = { "sycl-unique-id"="dg_int2" "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 #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 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" } -attributes #4 = { 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 #3 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #4 = { 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" } attributes #5 = { convergent } attributes #6 = { convergent nounwind } -!llvm.dependent-libraries = !{!0} -!opencl.spir.version = !{!1} -!spirv.Source = !{!2} -!llvm.ident = !{!3} -!llvm.module.flags = !{!4, !5} - -!0 = !{!"libcpmt"} -!1 = !{i32 1, i32 2} -!2 = !{i32 4, i32 100000} -!3 = !{!"clang version 14.0.0"} -!4 = !{i32 1, !"wchar_size", i32 2} -!5 = !{i32 7, !"frame-pointer", i32 2} -!6 = !{} +!llvm.dependent-libraries = !{!4} +!opencl.spir.version = !{!5} +!spirv.Source = !{!6} +!llvm.ident = !{!7} +!llvm.module.flags = !{!8, !9} + +!0 = !{!1, !2, !3} +!1 = !{i32 6149, i32 1} +!2 = !{i32 6148, i32 0} +!3 = !{i32 6147, i32 1, !"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"} diff --git a/llvm/test/tools/sycl-post-link/emit_program_metadata.ll b/llvm/test/tools/sycl-post-link/emit_program_metadata.ll index 7b2ad12057af7..6fff2d8ea459e 100644 --- a/llvm/test/tools/sycl-post-link/emit_program_metadata.ll +++ b/llvm/test/tools/sycl-post-link/emit_program_metadata.ll @@ -6,10 +6,6 @@ target triple = "spir64-unknown-unknown" -attributes #0 = { "sycl-work-group-size"="4,2,1" } -attributes #1 = { "sycl-work-group-size"="4,2" } -attributes #2 = { "sycl-work-group-size"="4" } - !0 = !{i32 1, i32 2, i32 4} !1 = !{i32 2, i32 4} !2 = !{i32 4} @@ -19,27 +15,12 @@ define weak_odr spir_kernel void @SpirKernel1(float %arg1) !reqd_work_group_size ret void } -define weak_odr spir_kernel void @SpirKernel2(float %arg1) #0 { - call void @foo(float %arg1) - ret void -} - -define weak_odr spir_kernel void @SpirKernel3(float %arg1) !reqd_work_group_size !1 { - call void @foo(float %arg1) - ret void -} - -define weak_odr spir_kernel void @SpirKernel4(float %arg1) #1 { - call void @foo(float %arg1) - ret void -} - -define weak_odr spir_kernel void @SpirKernel5(float %arg1) !reqd_work_group_size !2 { +define weak_odr spir_kernel void @SpirKernel2(float %arg1) !reqd_work_group_size !1 { call void @foo(float %arg1) ret void } -define weak_odr spir_kernel void @SpirKernel6(float %arg1) #2 { +define weak_odr spir_kernel void @SpirKernel3(float %arg1) !reqd_work_group_size !2 { call void @foo(float %arg1) ret void } @@ -49,11 +30,8 @@ declare void @foo(float) ; CHECK-PROP: [SYCL/program metadata] ; // Base64 encoding in the prop file (including 8 bytes length): ; CHECK-PROP-NEXT: SpirKernel1@reqd_work_group_size=2|gBAAAAAAAAQAAAAACAAAAQAAAAA -; CHECK-PROP-NEXT: SpirKernel2@reqd_work_group_size=2|gBAAAAAAAAQAAAAACAAAAQAAAAA -; CHECK-PROP-NEXT: SpirKernel3@reqd_work_group_size=2|ABAAAAAAAAgAAAAAEAAAAA -; CHECK-PROP-NEXT: SpirKernel4@reqd_work_group_size=2|ABAAAAAAAAgAAAAAEAAAAA -; CHECK-PROP-NEXT: SpirKernel5@reqd_work_group_size=2|gAAAAAAAAAABAAAA -; CHECK-PROP-NEXT: SpirKernel6@reqd_work_group_size=2|gAAAAAAAAAABAAAA +; CHECK-PROP-NEXT: SpirKernel2@reqd_work_group_size=2|ABAAAAAAAAgAAAAAEAAAAA +; CHECK-PROP-NEXT: SpirKernel3@reqd_work_group_size=2|gAAAAAAAAAABAAAA ; CHECK-TABLE: [Code|Properties] ; CHECK-TABLE-NEXT: {{.*}}files_0.prop diff --git a/llvm/tools/sycl-post-link/CMakeLists.txt b/llvm/tools/sycl-post-link/CMakeLists.txt index 60801899d4f81..976b20534e903 100644 --- a/llvm/tools/sycl-post-link/CMakeLists.txt +++ b/llvm/tools/sycl-post-link/CMakeLists.txt @@ -23,8 +23,6 @@ include_directories( add_llvm_tool(sycl-post-link sycl-post-link.cpp - CompileTimePropertiesPass.cpp - DeviceGlobals.cpp ModuleSplitter.cpp SpecConstants.cpp SYCLDeviceLibReqMask.cpp diff --git a/llvm/tools/sycl-post-link/ModuleSplitter.cpp b/llvm/tools/sycl-post-link/ModuleSplitter.cpp index cae151816cd2b..a517d7639fffc 100644 --- a/llvm/tools/sycl-post-link/ModuleSplitter.cpp +++ b/llvm/tools/sycl-post-link/ModuleSplitter.cpp @@ -9,7 +9,6 @@ //===----------------------------------------------------------------------===// #include "ModuleSplitter.h" -#include "DeviceGlobals.h" #include "Support.h" #include "llvm/ADT/SetVector.h" @@ -19,6 +18,7 @@ #include "llvm/IR/Instructions.h" #include "llvm/IR/LegacyPassManager.h" #include "llvm/IR/Module.h" +#include "llvm/SYCLLowerIR/DeviceGlobals.h" #include "llvm/SYCLLowerIR/LowerInvokeSimd.h" #include "llvm/SYCLLowerIR/LowerKernelProps.h" #include "llvm/SYCLLowerIR/SYCLUtils.h" diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index f8c426e641c5d..f9110752bc331 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -13,8 +13,6 @@ // - specialization constant intrinsic transformation //===----------------------------------------------------------------------===// -#include "CompileTimePropertiesPass.h" -#include "DeviceGlobals.h" #include "ModuleSplitter.h" #include "SYCLDeviceLibReqMask.h" #include "SYCLDeviceRequirements.h" @@ -37,6 +35,7 @@ #include "llvm/IRReader/IRReader.h" #include "llvm/Linker/Linker.h" #include "llvm/Passes/PassBuilder.h" +#include "llvm/SYCLLowerIR/DeviceGlobals.h" #include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h" #include "llvm/SYCLLowerIR/LowerInvokeSimd.h" #include "llvm/SYCLLowerIR/LowerKernelProps.h" @@ -846,11 +845,6 @@ processInputModule(std::unique_ptr M) { DUMP_ENTRY_POINTS(MDesc2.entries(), MDesc2.Name.c_str(), 3); Modified |= processSpecConstants(MDesc2); - // TODO: detach compile-time properties from device globals. - if (DeviceGlobals.getNumOccurrences() > 0) { - Modified |= - runModulePass(MDesc2.getModule()); - } if (!MDesc2.isSYCL() && LowerEsimd) { assert(MDesc2.isESIMD() && "NYI"); // ESIMD lowering also detects large-GRF kernels, so it must happen