[SYCL] Optional kernel features: implement split based on reqd-sub-group-size (take 2)

llvm/test/tools/sycl-post-link/device-code-split/per-reqd-sub-group-size-split-1.ll

@@ -0,0 +1,133 @@
+; This test emulates two translation units with 3 kernels:
+; TU0_kernel0 - 1st translation unit, no reqd_sub_group_size attribute used
+; TU0_kernel1 - 1st translation unit, reqd_sub_group_size attribute is used
+; TU1_kernel2 - 2nd translation unit, no reqd_sub_group_size attribute used
+
+; The test is intended to check that sycl-post-link correctly separates kernels
+; that use reqd_sub_group_size attributes from kernels which doesn't use them
+; regardless of device code split mode
+
+; RUN: sycl-post-link -split=auto -symbols -S %s -o %t.table
+; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-M0-IR \
+; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1
+; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M1-IR \
+; RUN: --implicit-check-not kernel0 --implicit-check-not kernel2
+; RUN: FileCheck %s -input-file=%t_2.ll --check-prefixes CHECK-M2-IR \
+; RUN: --implicit-check-not kernel1 --implicit-check-not kernel2
+; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-M0-SYMS \
+; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1
+; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-M1-SYMS \
+; RUN: --implicit-check-not kernel0 --implicit-check-not kernel2
+; RUN: FileCheck %s -input-file=%t_2.sym --check-prefixes CHECK-M2-SYMS \
+; RUN: --implicit-check-not kernel1 --implicit-check-not kernel2
+
+; RUN: sycl-post-link -split=kernel -symbols -S %s -o %t.table
+; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-M0-IR \
+; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1
+; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M1-IR \
+; RUN: --implicit-check-not kernel0 --implicit-check-not kernel2
+; RUN: FileCheck %s -input-file=%t_2.ll --check-prefixes CHECK-M2-IR \
+; RUN: --implicit-check-not kernel1 --implicit-check-not kernel2
+; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-M0-SYMS \
+; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1
+; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-M1-SYMS \
+; RUN: --implicit-check-not kernel0 --implicit-check-not kernel2
+; RUN: FileCheck %s -input-file=%t_2.sym --check-prefixes CHECK-M2-SYMS \
+; RUN: --implicit-check-not kernel1 --implicit-check-not kernel2
+
+; RUN: sycl-post-link -split=source -symbols -S %s -o %t.table
+; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-M0-IR \
+; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1
+; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M1-IR \
+; RUN: --implicit-check-not kernel0 --implicit-check-not kernel2
+; RUN: FileCheck %s -input-file=%t_2.ll --check-prefixes CHECK-M2-IR \
+; RUN: --implicit-check-not kernel1 --implicit-check-not kernel2
+; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-M0-SYMS \
+; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1
+; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-M1-SYMS \
+; RUN: --implicit-check-not kernel0 --implicit-check-not kernel2
+; RUN: FileCheck %s -input-file=%t_2.sym --check-prefixes CHECK-M2-SYMS \
+; RUN: --implicit-check-not kernel1 --implicit-check-not kernel2
+
+; Regardless of device code split mode, each kernel should go into a separate
+; device image
+
+; CHECK-M2-IR: define {{.*}} @TU0_kernel0
+; CHECK-M2-SYMS: TU0_kernel0
+
+; CHECK-M1-IR: define {{.*}} @TU0_kernel1
+; CHECK-M1-SYMS: TU0_kernel1
+
+; CHECK-M0-IR: define {{.*}} @TU1_kernel2
+; CHECK-M0-SYMS: TU1_kernel2
+
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
+target triple = "spir64-unknown-linux"
+
+; FIXME: device globals should also be properly distributed across device images
+; if they are of optional type
+@_ZL2GV = internal addrspace(1) constant [1 x i32] [i32 42], align 4
+
+define dso_local spir_kernel void @TU0_kernel0() #0 {
+entry:
+  call spir_func void @foo()
+  ret void
+}
+
+define dso_local spir_func void @foo() {
+entry:
+  %a = alloca i32, align 4
+  %call = call spir_func i32 @bar(i32 1)
+  %add = add nsw i32 2, %call
+  store i32 %add, i32* %a, align 4
+  ret void
+}
+
+; Function Attrs: nounwind
+define linkonce_odr dso_local spir_func i32 @bar(i32 %arg) {
+entry:
+  %arg.addr = alloca i32, align 4
+  store i32 %arg, i32* %arg.addr, align 4
+  %0 = load i32, i32* %arg.addr, align 4
+  ret i32 %0
+}
+
+define dso_local spir_kernel void @TU0_kernel1() #0 !intel_reqd_sub_group_size !2 {
+entry:
+  call spir_func void @foo1()
+  ret void
+}
+
+; Function Attrs: nounwind
+define dso_local spir_func void @foo1() {
+entry:
+  %a = alloca i32, align 4
+  store i32 2, i32* %a, align 4
+  ret void
+}
+
+define dso_local spir_kernel void @TU1_kernel2() #1 {
+entry:
+  call spir_func void @foo2()
+  ret void
+}
+
+; Function Attrs: nounwind
+define dso_local spir_func void @foo2() {
+entry:
+  %a = alloca i32, align 4
+  %0 = load i32, i32 addrspace(4)* getelementptr inbounds ([1 x i32], [1 x i32] addrspace(4)* addrspacecast ([1 x i32] addrspace(1)* @_ZL2GV to [1 x i32] addrspace(4)*), i64 0, i64 0), align 4
+  %add = add nsw i32 4, %0
+  store i32 %add, i32* %a, align 4
+  ret void
+}
+
+attributes #0 = { "sycl-module-id"="TU1.cpp" }
+attributes #1 = { "sycl-module-id"="TU2.cpp" }
+
+!opencl.spir.version = !{!0, !0}
+!spirv.Source = !{!1, !1}
+
+!0 = !{i32 1, i32 2}
+!1 = !{i32 4, i32 100000}
+!2 = !{i32 32}

llvm/test/tools/sycl-post-link/device-code-split/per-reqd-sub-group-size-split-2.ll

@@ -0,0 +1,60 @@
+; The test is intended to check that sycl-post-link correctly groups kernels
+; by unique reqd_sub_group_size values used in them
+
+; RUN: sycl-post-link -split=auto -symbols -S %s -o %t.table
+; RUN: FileCheck %s -input-file=%t.table --check-prefix CHECK-TABLE
+;
+; RUN: FileCheck %s -input-file=%t_0.sym --check-prefix CHECK-M0-SYMS \
+; RUN:     --implicit-check-not kernel0 --implicit-check-not kernel3
+;
+; RUN: FileCheck %s -input-file=%t_1.sym --check-prefix CHECK-M1-SYMS \
+; RUN:     --implicit-check-not kernel1 --implicit-check-not kernel2 \
+; RUN:     --implicit-check-not kernel3
+
+;
+; RUN: FileCheck %s -input-file=%t_2.sym --check-prefix CHECK-M2-SYMS \
+; RUN:     --implicit-check-not kernel0 --implicit-check-not kernel1 \
+; RUN:     --implicit-check-not kernel2
+
+; CHECK-TABLE: Code
+; CHECK-TABLE-NEXT: _0.sym
+; CHECK-TABLE-NEXT: _1.sym
+; CHECK-TABLE-NEXT: _2.sym
+; CHECK-TABLE-EMPTY:
+
+; CHECK-M0-SYMS: kernel1
+; CHECK-M0-SYMS: kernel2
+
+; CHECK-M1-SYMS: kernel0
+
+; CHECK-M2-SYMS: kernel3
+
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
+target triple = "spir64-unknown-linux"
+
+define dso_local spir_kernel void @kernel0() #0 !intel_reqd_sub_group_size !1 {
+entry:
+  ret void
+}
+
+define dso_local spir_kernel void @kernel1() #0 !intel_reqd_sub_group_size !2 {
+entry:
+  ret void
+}
+
+define dso_local spir_kernel void @kernel2() #0 !intel_reqd_sub_group_size !3 {
+entry:
+  ret void
+}
+
+define dso_local spir_kernel void @kernel3() #0 !intel_reqd_sub_group_size !4 {
+entry:
+  ret void
+}
+
+attributes #0 = { "sycl-module-id"="TU1.cpp" }
+
+!1 = !{i32 32}
+!2 = !{i32 64}
+!3 = !{i32 64}
+!4 = !{i32 16}

llvm/test/tools/sycl-post-link/device-requirements/reqd-sub-group-size.ll

@@ -0,0 +1,126 @@
+; Original code:
+; #include <sycl/sycl.hpp>
+
+; int main() {
+;   sycl::queue q;
+;   q.submit([&](sycl::handler &h) {
+;       h.parallel_for<class KernelA>(
+;           sycl::range<1>(32),
+;           [=](sycl::item<1> it) [[sycl::reqd_sub_group_size(16)]] {});
+;     });
+;   q.submit([&](sycl::handler &h) {
+;       h.parallel_for<class KernelB>(
+;           sycl::range<1>(32),
+;           [=](sycl::item<1> it) [[sycl::reqd_sub_group_size(32)]] {});
+;     });
+;   q.submit([&](sycl::handler &h) {
+;       h.parallel_for<class KernelC>(
+;           sycl::range<1>(32),
+;           [=](sycl::item<1> it) [[sycl::reqd_sub_group_size(16)]] {});
+;     });
+;   return 0;
+; }
+
+; RUN: sycl-post-link -split=auto %s -o %t.table
+; RUN: FileCheck %s -input-file=%t_0.prop --check-prefix CHECK-PROP-AUTO-SPLIT-0
+; RUN: FileCheck %s -input-file=%t_1.prop --check-prefix CHECK-PROP-AUTO-SPLIT-1
+
+; CHECK-PROP-AUTO-SPLIT-0: [SYCL/device requirements]
+; CHECK-PROP-AUTO-SPLIT-0: reqd_sub_group_size=1|32
+
+; CHECK-PROP-AUTO-SPLIT-1: [SYCL/device requirements]
+; CHECK-PROP-AUTO-SPLIT-1: reqd_sub_group_size=1|16
+
+; ModuleID = 'foo.cpp'
+source_filename = "foo.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"
+
+$_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E7KernelA = comdat any
+
+$_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_E7KernelB = comdat any
+
+$_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE1_clES2_E7KernelC = comdat any
+
+; Function Attrs: norecurse nounwind
+define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E7KernelA() local_unnamed_addr #0 comdat !srcloc !48 !kernel_arg_buffer_location !49 !intel_reqd_sub_group_size !50 !sycl_fixed_targets !49 !sycl_kernel_omit_args !49 {
+entry:
+  ret void
+}
+
+; Function Attrs: norecurse nounwind
+define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_E7KernelB() local_unnamed_addr #0 comdat !srcloc !51 !kernel_arg_buffer_location !49 !intel_reqd_sub_group_size !52 !sycl_fixed_targets !49 !sycl_kernel_omit_args !49 {
+entry:
+  ret void
+}
+
+; Function Attrs: norecurse nounwind
+define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE1_clES2_E7KernelC() local_unnamed_addr #0 comdat !srcloc !53 !kernel_arg_buffer_location !49 !intel_reqd_sub_group_size !50 !sycl_fixed_targets !49 !sycl_kernel_omit_args !49 {
+entry:
+  ret void
+}
+
+declare dso_local spir_func i32 @_Z18__spirv_ocl_printfPU3AS2Kcz(i8 addrspace(2)*, ...)
+
+attributes #0 = { norecurse nounwind "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="foo.cpp" "sycl-optlevel"="2" "uniform-work-group-size"="true" }
+
+!llvm.module.flags = !{!0, !1}
+!opencl.spir.version = !{!2}
+!spirv.Source = !{!3}
+!sycl_aspects = !{!4, !5, !6, !7, !8, !9, !10, !11, !12, !13, !14, !15, !16, !17, !18, !19, !20, !21, !22, !23, !24, !25, !26, !27, !28, !29, !30, !31, !32, !33, !34, !35, !36, !37, !38, !39, !40, !41, !42, !43, !44, !45, !46}
+!llvm.ident = !{!47}
+
+!0 = !{i32 1, !"wchar_size", i32 4}
+!1 = !{i32 7, !"frame-pointer", i32 2}
+!2 = !{i32 1, i32 2}
+!3 = !{i32 4, i32 100000}
+!4 = !{!"cpu", i32 1}
+!5 = !{!"gpu", i32 2}
+!6 = !{!"accelerator", i32 3}
+!7 = !{!"custom", i32 4}
+!8 = !{!"fp16", i32 5}
+!9 = !{!"fp64", i32 6}
+!10 = !{!"image", i32 9}
+!11 = !{!"online_compiler", i32 10}
+!12 = !{!"online_linker", i32 11}
+!13 = !{!"queue_profiling", i32 12}
+!14 = !{!"usm_device_allocations", i32 13}
+!15 = !{!"usm_host_allocations", i32 14}
+!16 = !{!"usm_shared_allocations", i32 15}
+!17 = !{!"usm_system_allocations", i32 17}
+!18 = !{!"ext_intel_pci_address", i32 18}
+!19 = !{!"ext_intel_gpu_eu_count", i32 19}
+!20 = !{!"ext_intel_gpu_eu_simd_width", i32 20}
+!21 = !{!"ext_intel_gpu_slices", i32 21}
+!22 = !{!"ext_intel_gpu_subslices_per_slice", i32 22}
+!23 = !{!"ext_intel_gpu_eu_count_per_subslice", i32 23}
+!24 = !{!"ext_intel_max_mem_bandwidth", i32 24}
+!25 = !{!"ext_intel_mem_channel", i32 25}
+!26 = !{!"usm_atomic_host_allocations", i32 26}
+!27 = !{!"usm_atomic_shared_allocations", i32 27}
+!28 = !{!"atomic64", i32 28}
+!29 = !{!"ext_intel_device_info_uuid", i32 29}
+!30 = !{!"ext_oneapi_srgb", i32 30}
+!31 = !{!"ext_oneapi_native_assert", i32 31}
+!32 = !{!"host_debuggable", i32 32}
+!33 = !{!"ext_intel_gpu_hw_threads_per_eu", i32 33}
+!34 = !{!"ext_oneapi_cuda_async_barrier", i32 34}
+!35 = !{!"ext_oneapi_bfloat16_math_functions", i32 35}
+!36 = !{!"ext_intel_free_memory", i32 36}
+!37 = !{!"ext_intel_device_id", i32 37}
+!38 = !{!"ext_intel_memory_clock_rate", i32 38}
+!39 = !{!"ext_intel_memory_bus_width", i32 39}
+!40 = !{!"emulated", i32 40}
+!41 = !{!"ext_intel_legacy_image", i32 41}
+!42 = !{!"int64_base_atomics", i32 7}
+!43 = !{!"int64_extended_atomics", i32 8}
+!44 = !{!"usm_system_allocator", i32 17}
+!45 = !{!"usm_restricted_shared_allocations", i32 16}
+!46 = !{!"host", i32 0}
+!47 = !{!"clang version 17.0.0 (https://github.com/jzc/llvm eed5b5576bef314433e8ae7313620dae399c9d22)"}
+!48 = !{i32 170}
+!49 = !{}
+!50 = !{i32 16}
+!51 = !{i32 351}
+!52 = !{i32 32}
+!53 = !{i32 532}

llvm/test/tools/sycl-post-link/registerallocmode.ll

@@ -2,41 +2,49 @@
 
 ; RUN: sycl-post-link -split=source -symbols -split-esimd -lower-esimd -S < %s -o %t.table
 ; RUN: FileCheck %s -input-file=%t.table
-; RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes CHECK-ESIMD-LargeGRF-IR
-; RUN: FileCheck %s -input-file=%t_esimd_0.prop --check-prefixes CHECK-ESIMD-LargeGRF-PROP
-; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-SYCL-LargeGRF-IR
-; RUN: FileCheck %s -input-file=%t_0.prop --check-prefixes CHECK-SYCL-LargeGRF-PROP
-; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-SYCL-LargeGRF-SYM
-; RUN: FileCheck %s -input-file=%t_1.prop --check-prefixes CHECK-SYCL-PROP
-; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-SYCL-SYM
-; RUN: FileCheck %s -input-file=%t_esimd_1.prop --check-prefixes CHECK-ESIMD-PROP
-; RUN: FileCheck %s -input-file=%t_esimd_0.sym --check-prefixes CHECK-ESIMD-LargeGRF-SYM
 
 ; CHECK: [Code|Properties|Symbols]
-; CHECK: {{.*}}_esimd_0.ll|{{.*}}_esimd_0.prop|{{.*}}_esimd_0.sym
-; CHECK: {{.*}}_esimd_1.ll|{{.*}}_esimd_1.prop|{{.*}}_esimd_1.sym
-; CHECK: {{.*}}_1.ll|{{.*}}_1.prop|{{.*}}_1.sym
+; CHECK-NEXT: {{.*}}_esimd_0.ll|{{.*}}_esimd_0.prop|{{.*}}_esimd_0.sym
+; CHECK-NEXT: {{.*}}_1.ll|{{.*}}_1.prop|{{.*}}_1.sym
+; CHECK-NEXT: {{.*}}_esimd_2.ll|{{.*}}_esimd_2.prop|{{.*}}_esimd_2.sym
+; CHECK-NEXT: {{.*}}_3.ll|{{.*}}_3.prop|{{.*}}_3.sym
 
-; CHECK-ESIMD-LargeGRF-PROP: isEsimdImage=1|1
-; CHECK-ESIMD-LargeGRF-PROP: sycl-register-alloc-mode=1|2
+; RUN: FileCheck %s -input-file=%t_esimd_0.ll   --check-prefixes CHECK-ESIMD-LargeGRF-IR
+; RUN: FileCheck %s -input-file=%t_esimd_0.prop --check-prefixes CHECK-ESIMD-LargeGRF-PROP
+; RUN: FileCheck %s -input-file=%t_esimd_0.sym  --check-prefixes CHECK-ESIMD-LargeGRF-SYM
 
-; CHECK-SYCL-LargeGRF-PROP: sycl-register-alloc-mode=1|2
+; CHECK-ESIMD-LargeGRF-SYM: __ESIMD_large_grf_kernel
+; CHECK-ESIMD-LargeGRF-SYM-EMPTY:
 
-; CHECK-SYCL-PROP-NOT: sycl-register-alloc-mode
+; CHECK-ESIMD-LargeGRF-PROP: isEsimdImage=1|1
+; CHECK-ESIMD-LargeGRF-PROP: sycl-register-alloc-mode=1|2
 
-; CHECK-SYCL-SYM: __SYCL_kernel
-; CHECK-SYCL-SYM-EMPTY:
+; RUN: FileCheck %s -input-file=%t_1.ll   --check-prefixes CHECK-SYCL-LargeGRF-IR
+; RUN: FileCheck %s -input-file=%t_1.prop --check-prefixes CHECK-SYCL-LargeGRF-PROP
+; RUN: FileCheck %s -input-file=%t_1.sym  --check-prefixes CHECK-SYCL-LargeGRF-SYM
 
 ; CHECK-SYCL-LargeGRF-SYM: __SYCL_kernel_large_grf
 ; CHECK-SYCL-LargeGRF-SYM-EMPTY:
 
+; CHECK-SYCL-LargeGRF-PROP: sycl-register-alloc-mode=1|2
+
+; RUN: FileCheck %s -input-file=%t_esimd_2.prop --check-prefixes CHECK-ESIMD-PROP 
+; RUN: FileCheck %s -input-file=%t_esimd_2.sym  --check-prefixes CHECK-ESIMD-SYM 
+
 ; CHECK-ESIMD-SYM: __ESIMD_kernel
 ; CHECK-ESIMD-SYM-EMPTY:
 
 ; CHECK-ESIMD-PROP-NOT: sycl-register-alloc-mode
 
-; CHECK-ESIMD-LargeGRF-SYM: __ESIMD_large_grf_kernel
-; CHECK-ESIMD-LargeGRF-SYM-EMPTY:
+; RUN: FileCheck %s -input-file=%t_3.prop --check-prefixes CHECK-SYCL-PROP
+; RUN: FileCheck %s -input-file=%t_3.sym  --check-prefixes CHECK-SYCL-SYM
+
+; CHECK-SYCL-SYM: __SYCL_kernel
+; CHECK-SYCL-SYM-EMPTY:
+
+; CHECK-SYCL-PROP-NOT: sycl-register-alloc-mode
+
+
 
 ; ModuleID = 'large_grf.bc'
 source_filename = "llvm-link"