[SYCL] Separate CompileTimePropertiesPass from sycl-post-link tool (#7527)

Extracted the `CompileTimeProperties` pass from the `sycl-post-link`
tool up into the `llvm` directory with the rest of the passes and added
it as a step in the code generation pipeline. Moving this into a more
generalised location will allow for new tools to be developed that
require its logic.

Due to a dependency on `DeviceGlobals`, that also had to be moved and
includes adjusted in `sycl-post-link` tool. This pass can be invoked via
opt as well for testing.

Several `sycl-post-link` tests had to be transferred to
`llvm/test/SYCLLowerIR` with some being modified as well to account for
the change. Some tests had to be split across the two directories so as
to maintain coverage. Particular tests that focused on the result of
running the pass over some iput bytecode and testing the result have
been turned into their own tests under the
`llvm/test/SYCLLowerIR/CompileTimePropertiesPass` test directory.

Author: Andrew Lamzed-Short

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());
     }
   }
 

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]

llvm/tools/sycl-post-link/CompileTimePropertiesPass.h renamed to llvm/include/llvm/SYCLLowerIR/CompileTimePropertiesPass.h

@@ -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"

llvm/tools/sycl-post-link/DeviceGlobals.h renamed to llvm/include/llvm/SYCLLowerIR/DeviceGlobals.h

@@ -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

llvm/lib/Passes/PassBuilder.cpp

@@ -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

llvm/lib/Passes/PassRegistry.def

@@ -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"

llvm/lib/SYCLLowerIR/CMakeLists.txt

@@ -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"

llvm/tools/sycl-post-link/CompileTimeProperties.def renamed to llvm/lib/SYCLLowerIR/CompileTimeProperties.def

@@ -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]]}

llvm/tools/sycl-post-link/CompileTimePropertiesPass.cpp renamed to llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp

@@ -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

llvm/tools/sycl-post-link/DeviceGlobals.cpp renamed to llvm/lib/SYCLLowerIR/DeviceGlobals.cpp

@@ -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

llvm/test/SYCLLowerIR/CompileTimePropertiesPass/device-globals/test_global_variable.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

llvm/test/tools/sycl-post-link/kernel-attributes/kernel-pipelined.ll renamed to llvm/test/SYCLLowerIR/CompileTimePropertiesPass/kernel-attributes/kernel-pipelined.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

llvm/test/tools/sycl-post-link/kernel-attributes/register-map-interface.ll renamed to llvm/test/SYCLLowerIR/CompileTimePropertiesPass/kernel-attributes/register-map-interface.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"