Skip to content

Commit 22ca174

Browse files
author
iclsrc
committed
Merge from 'sycl' to 'sycl-web'
2 parents ab83849 + 9dfaf27 commit 22ca174

File tree

16 files changed

+1572
-531
lines changed

16 files changed

+1572
-531
lines changed

‎llvm/test/tools/sycl-post-link/device-requirements/joint-matrix.ll

Lines changed: 82 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,82 @@
1+
; Original code:
2+
;
3+
; // Compiled with clang++ -fsycl -fsycl-device-only -fno-sycl-instrument-device-code -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 -O2 -S -emit-llvm -o - %s
4+
; int main() {
5+
; sycl::queue q;
6+
; q.submit([&](sycl::handler &cgh) {
7+
; cgh.parallel_for(nd_range<2>({1, 16}, {1, 16}), [=](nd_item<2> it) {
8+
; joint_matrix<sycl::sub_group, float, use::a, 8, 16, layout::row_major> a;
9+
; joint_matrix<sycl::sub_group, float, use::b, 16, 16, layout::row_major> b;
10+
; joint_matrix<sycl::sub_group, float, use::accumulator, 8, 16> c;
11+
; sub_group sg = it.get_sub_group();
12+
; joint_matrix_mad(sg, c, a, b, c);
13+
; });
14+
; });
15+
; q.submit([&](sycl::handler &cgh) {
16+
; cgh.parallel_for(nd_range<2>({1, 16}, {1, 16}), [=](nd_item<2> it) {
17+
; joint_matrix<sycl::sub_group, double, use::a, 16, 16, layout::row_major>
18+
; a;
19+
; });
20+
; });
21+
; return 0;
22+
; }
23+
24+
; RUN: sycl-post-link -split=kernel %s -o %t.table
25+
; RUN: FileCheck %s -input-file=%t_0.prop --check-prefix CHECK-PROP-KERNEL-SPLIT-0
26+
; RUN: FileCheck %s -input-file=%t_1.prop --check-prefix CHECK-PROP-KERNEL-SPLIT-1
27+
28+
; CHECK-PROP-KERNEL-SPLIT-0: [SYCL/device requirements]
29+
; CHECK-PROP-KERNEL-SPLIT-0: joint_matrix=2|gMAAAAAAAAQbhRncph3X0lHclpjOmB3MywSdzVmO6EGL4wSM2sTbhRncph3X0lHclpjOmB3MywSdzVmO6E2YjVXb1xWY09mcsgDLxYzOtFGdylGefRXewVmO6YGczIDL1NXZ6ojYsEjNsEjN
30+
; CHECK-PROP-KERNEL-SPLIT-0-NEXT: joint_matrix_mad=2|4JAAAAAAAAQbhRncph3X0lHclpjOmB3MywSbhRncph3X0lHclpjOmB3MywSbhRncph3X0lHclpjOmB3MywSbhRncph3X0lHclpjOmB3MywCOsEjNsEjN
31+
32+
; CHECK-PROP-KERNEL-SPLIT-1: [SYCL/device requirements]
33+
; CHECK-PROP-KERNEL-SPLIT-1: joint_matrix=2|wDAAAAAAAAQbhRncph3X0lHclpjOmBnN0wSdzVmO6EGLxYDLxYD
34+
35+
; ModuleID = '/tmp/test.bc'
36+
source_filename = "llvm-link"
37+
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"
38+
target triple = "spir64-unknown-unknown"
39+
40+
$_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlNS0_7nd_itemILi2EEEE_ = comdat any
41+
42+
$_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_EUlNS0_7nd_itemILi2EEEE_ = comdat any
43+
44+
; Function Attrs: convergent norecurse nounwind
45+
define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlNS0_7nd_itemILi2EEEE_() local_unnamed_addr #0 comdat !srcloc !5 !kernel_arg_buffer_location !6 !sycl_fixed_targets !6 !sycl_joint_matrix !7 !sycl_joint_matrix_mad !8 !sycl_kernel_omit_args !6 {
46+
entry:
47+
%call.i.i = tail call spir_func noundef target("spirv.JointMatrixINTEL", float, 8, 16, 3, 3, 2) @_Z27__spirv_JointMatrixMadINTELIffLm8ELm16ELm16ELN5__spv9MatrixUseE0ELS1_1ELS1_2ELNS0_12MatrixLayoutE0ELS2_0ELS2_3ELNS0_5Scope4FlagE3EEPNS0_24__spirv_JointMatrixINTELIT0_XT1_EXT3_EXT9_EXT10_EXT6_EEEPNS5_IT_XT1_EXT2_EXT7_EXT10_EXT4_EEEPNS5_IS9_XT2_EXT3_EXT8_EXT10_EXT5_EEES8_S4_(target("spirv.JointMatrixINTEL", float, 8, 16, 0, 3, 0) noundef undef, target("spirv.JointMatrixINTEL", float, 16, 16, 0, 3, 1) noundef undef, target("spirv.JointMatrixINTEL", float, 8, 16, 3, 3, 2) noundef undef, i32 noundef 3) #3
48+
ret void
49+
}
50+
51+
; Function Attrs: convergent nounwind
52+
declare dso_local spir_func noundef target("spirv.JointMatrixINTEL", float, 8, 16, 3, 3, 2) @_Z27__spirv_JointMatrixMadINTELIffLm8ELm16ELm16ELN5__spv9MatrixUseE0ELS1_1ELS1_2ELNS0_12MatrixLayoutE0ELS2_0ELS2_3ELNS0_5Scope4FlagE3EEPNS0_24__spirv_JointMatrixINTELIT0_XT1_EXT3_EXT9_EXT10_EXT6_EEEPNS5_IT_XT1_EXT2_EXT7_EXT10_EXT4_EEEPNS5_IS9_XT2_EXT3_EXT8_EXT10_EXT5_EEES8_S4_(target("spirv.JointMatrixINTEL", float, 8, 16, 0, 3, 0) noundef, target("spirv.JointMatrixINTEL", float, 16, 16, 0, 3, 1) noundef, target("spirv.JointMatrixINTEL", float, 8, 16, 3, 3, 2) noundef, i32 noundef) local_unnamed_addr #1
53+
54+
; Function Attrs: norecurse nounwind
55+
define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_EUlNS0_7nd_itemILi2EEEE_() local_unnamed_addr #2 comdat !srcloc !9 !kernel_arg_buffer_location !6 !sycl_fixed_targets !6 !sycl_joint_matrix !10 !sycl_kernel_omit_args !6 {
56+
entry:
57+
ret void
58+
}
59+
60+
declare dso_local spir_func i32 @_Z18__spirv_ocl_printfPU3AS2Kcz(ptr addrspace(2), ...)
61+
62+
attributes #0 = { convergent norecurse nounwind "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="sss.cpp" "sycl-optlevel"="2" "uniform-work-group-size"="true" }
63+
attributes #1 = { convergent nounwind "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
64+
attributes #2 = { norecurse nounwind "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="sss.cpp" "sycl-optlevel"="2" "uniform-work-group-size"="true" }
65+
attributes #3 = { convergent nounwind }
66+
67+
!llvm.module.flags = !{!0, !1}
68+
!opencl.spir.version = !{!2}
69+
!spirv.Source = !{!3}
70+
!llvm.ident = !{!4}
71+
72+
!0 = !{i32 1, !"wchar_size", i32 4}
73+
!1 = !{i32 7, !"frame-pointer", i32 2}
74+
!2 = !{i32 1, i32 2}
75+
!3 = !{i32 4, i32 100000}
76+
!4 = !{!""}
77+
!5 = !{i32 1091}
78+
!6 = !{}
79+
!7 = !{!"matrix_type::fp32,use::a,8,16;matrix_type::fp32,use::accumulator,8,16;matrix_type::fp32,use::b,16,16"}
80+
!8 = !{!"matrix_type::fp32,matrix_type::fp32,matrix_type::fp32,matrix_type::fp32,8,16,16"}
81+
!9 = !{i32 1529}
82+
!10 = !{!"matrix_type::fp64,use::a,16,16"}

‎llvm/tools/sycl-post-link/SYCLDeviceRequirements.cpp

Lines changed: 26 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -9,6 +9,7 @@
99
#include "SYCLDeviceRequirements.h"
1010
#include "ModuleSplitter.h"
1111

12+
#include "llvm/ADT/SmallString.h"
1213
#include "llvm/ADT/StringRef.h"
1314
#include "llvm/IR/Module.h"
1415
#include "llvm/Support/PropertySetIO.h"
@@ -88,6 +89,31 @@ void llvm::getSYCLDeviceRequirements(
8889
if (ReqdWorkGroupSize)
8990
Requirements["reqd_work_group_size_uint64_t"] = *ReqdWorkGroupSize;
9091

92+
auto ExtractStringFromMDNodeOperand =
93+
[=](const MDNode *N, unsigned OpNo) -> llvm::SmallString<256> {
94+
MDString *S = cast<llvm::MDString>(N->getOperand(OpNo).get());
95+
return S->getString();
96+
};
97+
98+
// { LLVM-IR metadata name , [SYCL/Device requirements] property name }, see:
99+
// https://github.com/intel/llvm/blob/sycl/sycl/doc/design/OptionalDeviceFeatures.md#create-the-sycldevice-requirements-property-set
100+
// Scan the module and if the metadata is present fill the corresponing
101+
// property with metadata's aspects
102+
constexpr std::pair<const char *, const char *> MatrixMDs[] = {
103+
{"sycl_joint_matrix", "joint_matrix"},
104+
{"sycl_joint_matrix_mad", "joint_matrix_mad"}};
105+
106+
for (const auto &[MDName, MappedName] : MatrixMDs) {
107+
llvm::SmallString<256> Val;
108+
for (const Function &F : MD.getModule())
109+
if (const MDNode *MDN = F.getMetadata(MDName))
110+
Val = ExtractStringFromMDNodeOperand(
111+
MDN, 0); // there is always only one operand
112+
if (Val.empty())
113+
continue;
114+
Requirements[MappedName] = Val;
115+
}
116+
91117
// There should only be at most one function with
92118
// intel_reqd_sub_group_size metadata when considering the entry
93119
// points of a module, but not necessarily when considering all the

0 commit comments

Comments
 (0)