Skip to content

Commit 69d8d7a

Browse files
kbobrovskbobrovs
committed
[SYCL][ESIMD] LowerESIMD: add 'buffer_t' MD for accessor kernel arguments.
Signed-off-by: Konstantin S Bobrovsky <[email protected]>
1 parent f768b3d commit 69d8d7a

File tree

5 files changed

+104
-6
lines changed

5 files changed

+104
-6
lines changed

‎clang/test/CodeGenSYCL/Inputs/sycl.hpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -156,6 +156,7 @@ class accessor {
156156
private:
157157
void __init(__attribute__((opencl_global)) dataT *Ptr, range<dimensions> AccessRange,
158158
range<dimensions> MemRange, id<dimensions> Offset) {}
159+
void __init_esimd(__attribute__((opencl_global)) dataT *Ptr) {}
159160
};
160161

161162
template <int dimensions, access::mode accessmode, access::target accesstarget>

‎clang/test/CodeGenSYCL/esimd-accessor-ptr-md.cpp

Lines changed: 44 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,44 @@
1+
// RUN: %clang_cc1 -fsycl -fsycl-explicit-simd -fsycl-is-device \
2+
// RUN: -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice \
3+
// RUN: -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
4+
5+
// This test checks
6+
// 1) proper metadata generation for accessors used in ESIMD
7+
// kernels:
8+
// - Proper 'kernel_arg_accessor_ptr' metadata is generated by the FE for
9+
// ESIMD kernels
10+
// - Pointers originating from accessors are marked with 'buffer_t' and proper
11+
// argument kind.
12+
// 2) __init_esimd function is used to initialize the accessor rather than
13+
// __init.
14+
15+
#include "sycl.hpp"
16+
17+
using namespace cl::sycl;
18+
19+
void test(int val) {
20+
queue q;
21+
q.submit([&](handler &h) {
22+
cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write> accessorA;
23+
cl::sycl::accessor<int, 1, cl::sycl::access::mode::read> accessorB;
24+
25+
h.single_task<class esimd_kernel>(
26+
[=]() __attribute__((sycl_explicit_simd)) {
27+
accessorA.use(val);
28+
accessorB.use();
29+
});
30+
});
31+
32+
// --- Name
33+
// CHECK-LABEL: define spir_kernel void @"_ZTSZZ4testiENK3$_0clERN2cl4sycl7handlerEE12esimd_kernel"(
34+
// --- Signature
35+
// CHECK: i32 addrspace(1)* "VCArgumentDesc"="buffer_t" "VCArgumentIOKind"="0" "VCArgumentKind"="2" %_arg_,
36+
// CHECK: i32 "VCArgumentDesc" "VCArgumentIOKind"="0" "VCArgumentKind"="0" %_arg_1,
37+
// CHECK: i32 addrspace(1)* "VCArgumentDesc"="buffer_t" "VCArgumentIOKind"="0" "VCArgumentKind"="2" %_arg_3)
38+
// --- Attributes
39+
// CHECK: {{.*}} !kernel_arg_accessor_ptr ![[ACC_PTR_ATTR:[0-9]+]] !sycl_explicit_simd !{{[0-9]+}} {{.*}}{
40+
// --- init_esimd call is expected instead of __init:
41+
// CHECK: call spir_func void @{{.*}}__init_esimd{{.*}}(%"{{.*}}sycl::accessor" addrspace(4)* %{{[0-9]+}}, i32 addrspace(1)* %{{[0-9]+}})
42+
// CHECK-LABEL: }
43+
// CHECK: ![[ACC_PTR_ATTR]] = !{i1 true, i1 false, i1 true}
44+
}

‎clang/test/CodeGenSYCL/kernel-by-reference.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -15,15 +15,15 @@ int simple_add(int i) {
1515
int main() {
1616
queue q;
1717
#if defined(SYCL2020)
18-
// expected-warning@Inputs/sycl.hpp:285 {{Passing kernel functions by value is deprecated in SYCL 2020}}
18+
// expected-warning@Inputs/sycl.hpp:286 {{Passing kernel functions by value is deprecated in SYCL 2020}}
1919
// expected-note@+3 {{in instantiation of function template specialization}}
2020
#endif
2121
q.submit([&](handler &h) {
2222
h.single_task_2017<class sycl2017>([]() { simple_add(10); });
2323
});
2424

2525
#if defined(SYCL2017)
26-
// expected-warning@Inputs/sycl.hpp:280 {{Passing of kernel functions by reference is a SYCL 2020 extension}}
26+
// expected-warning@Inputs/sycl.hpp:281 {{Passing of kernel functions by reference is a SYCL 2020 extension}}
2727
// expected-note@+3 {{in instantiation of function template specialization}}
2828
#endif
2929
q.submit([&](handler &h) {

‎llvm/lib/SYCLLowerIR/LowerESIMD.cpp

Lines changed: 20 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1136,6 +1136,7 @@ void SYCLLowerESIMDLegacyPass::generateKernelMetadata(Module &M) {
11361136
SmallVector<Metadata *, 8> ArgTypeDescs;
11371137

11381138
auto *KernelArgTypes = F.getMetadata("kernel_arg_type");
1139+
auto *KernelArgAccPtrs = F.getMetadata("kernel_arg_accessor_ptr");
11391140
unsigned Idx = 0;
11401141

11411142
// Iterate argument list to gather argument kinds and generate argument
@@ -1148,14 +1149,29 @@ void SYCLLowerESIMDLegacyPass::generateKernelMetadata(Module &M) {
11481149

11491150
if (ArgType.find("image1d_t") != std::string::npos ||
11501151
ArgType.find("image2d_t") != std::string::npos ||
1151-
ArgType.find("image3d_t") != std::string::npos ||
1152-
ArgType.find("image1d_buffer_t") != std::string::npos) {
1152+
ArgType.find("image3d_t") != std::string::npos) {
11531153
Kind = AK_SURFACE;
11541154
ArgTypeDescs.push_back(MDString::get(Ctx, ArgType));
11551155
} else {
11561156
StringRef ArgDesc = "";
1157-
if (Arg.getType()->isPointerTy())
1158-
ArgDesc = "svmptr_t";
1157+
1158+
if (Arg.getType()->isPointerTy()) {
1159+
const auto *IsAccMD =
1160+
KernelArgAccPtrs
1161+
? cast<ConstantAsMetadata>(KernelArgAccPtrs->getOperand(Idx))
1162+
: nullptr;
1163+
unsigned IsAcc =
1164+
IsAccMD
1165+
? static_cast<unsigned>(cast<ConstantInt>(IsAccMD->getValue())
1166+
->getValue()
1167+
.getZExtValue())
1168+
: 0;
1169+
if (IsAcc) {
1170+
ArgDesc = "buffer_t";
1171+
Kind = AK_SURFACE;
1172+
} else
1173+
ArgDesc = "svmptr_t";
1174+
}
11591175
ArgTypeDescs.push_back(MDString::get(Ctx, ArgDesc));
11601176
}
11611177

‎llvm/test/SYCLLowerIR/esimd_acc_ptr.ll

Lines changed: 37 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,37 @@
1+
; RUN: opt -LowerESIMD -S < %s | FileCheck %s
2+
3+
; This test checks that LowerESIMD pass correctly interpretes the
4+
; 'kernel_arg_accessor_ptr' metadata. Particularly, that it generates additional
5+
; vector of per-argument metadata (accessible from "genx.kernels" top-level
6+
; metadata node):
7+
; - for those arguments having non-zero in the corresponding
8+
; 'kernel_arg_accessor_ptr' position:
9+
; * "argument kind" metadata element is set to '2' - 'surface'
10+
; * "argument descriptor" metadata element is set to 'buffer_t'
11+
; - for those pointer arguments having '0' in the corresponding
12+
; 'kernel_arg_accessor_ptr' position, the kind/descriptor is set to
13+
; '0'/'svmptr_t'
14+
15+
define weak_odr dso_local spir_kernel void @ESIMDKernel(i32 %_arg_, float addrspace(1)* %_arg_1, float addrspace(1)* %_arg_3, i32 %_arg_5, float addrspace(1)* %_arg_7) !kernel_arg_accessor_ptr !0 !sycl_explicit_simd !1 !intel_reqd_sub_group_size !2 {
16+
; CHECK: {{.*}} spir_kernel void @ESIMDKernel({{.*}}) #[[GENX_MAIN:[0-9]+]]
17+
ret void
18+
}
19+
20+
; kernel_arg_accessor_ptr:
21+
; arg0=<scalar>
22+
; arg1=<ptr from accessor>
23+
; arg2=<ptr from accessor>
24+
; arg3=<scalar>
25+
; arg4=<ptr>
26+
; buffer_t and argument kind 2 (surface) metadata must be added for args 1 and 2
27+
!0 = !{i32 0, i32 1, i32 1, i32 0, i32 0}
28+
!1 = !{}
29+
!2 = !{i32 1}
30+
31+
; CHECK: attributes #[[GENX_MAIN]] = { "CMGenxMain" "oclrt"="1" }
32+
; CHECK: !genx.kernels = !{![[GENX_KERNELS:[0-9]+]]}
33+
; CHECK: ![[GENX_KERNELS]] = !{void (i32, float addrspace(1)*, float addrspace(1)*, i32, float addrspace(1)*)* @ESIMDKernel, !"ESIMDKernel", ![[ARG_KINDS:[0-9]+]], i32 0, i32 0, ![[ARG_IO_KINDS:[0-9]+]], ![[ARG_DESCS:[0-9]+]]}
34+
; CHECK: ![[ARG_KINDS]] = !{i32 0, i32 2, i32 2, i32 0, i32 0}
35+
; CHECK: ![[ARG_IO_KINDS]] = !{i32 0, i32 0, i32 0, i32 0, i32 0}
36+
; CHECK: ![[ARG_DESCS]] = !{!"", !"buffer_t", !"buffer_t", !"", !"svmptr_t"}
37+

0 commit comments

Comments
 (0)