[SYCL] Separated CompileTimePropertiesPass from sycl-post-link tool

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

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

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

llvm/tools/sycl-post-link/CompileTimePropertiesPass.cpp → 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"

llvm/tools/sycl-post-link/DeviceGlobals.cpp → 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"

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

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

llvm/test/tools/sycl-post-link/kernel-attributes/register-map-interface.ll → 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

llvm/test/tools/sycl-post-link/kernel-attributes/streaming-interface.ll → 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

llvm/test/tools/sycl-post-link/kernel-properties.ll → 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

llvm/test/tools/sycl-post-link/sycl-properties-ptr-annotations.ll → 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"