Skip to content

Commit 32945a2

Browse files
fabiomestrefabiomestre
committed
Implement dynamic local accessors
1 parent 6f9292e commit 32945a2

File tree

22 files changed

+1028
-232
lines changed

22 files changed

+1028
-232
lines changed

‎clang/include/clang/Basic/Attr.td

+4-2
Original file line numberDiff line numberDiff line change
@@ -1621,12 +1621,14 @@ def SYCLType: InheritableAttr {
16211621
let Subjects = SubjectList<[CXXRecord, Enum], ErrorDiag>;
16221622
let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
16231623
let Args = [EnumArgument<"Type", "SYCLType", /*is_string=*/true,
1624-
["accessor", "local_accessor", "work_group_memory", "dynamic_work_group_memory",
1624+
["accessor", "local_accessor", "dynamic_local_accessor",
1625+
"work_group_memory", "dynamic_work_group_memory",
16251626
"specialization_id", "kernel_handler", "buffer_location",
16261627
"no_alias", "accessor_property_list", "group",
16271628
"private_memory", "aspect", "annotated_ptr", "annotated_arg",
16281629
"stream", "sampler", "host_pipe", "multi_ptr"],
1629-
["accessor", "local_accessor", "work_group_memory", "dynamic_work_group_memory",
1630+
["accessor", "local_accessor", "dynamic_local_accessor",
1631+
"work_group_memory", "dynamic_work_group_memory",
16301632
"specialization_id", "kernel_handler", "buffer_location",
16311633
"no_alias", "accessor_property_list", "group",
16321634
"private_memory", "aspect", "annotated_ptr", "annotated_arg",

‎clang/include/clang/Sema/SemaSYCL.h

+2-1
Original file line numberDiff line numberDiff line change
@@ -64,7 +64,8 @@ class SYCLIntegrationHeader {
6464
kind_stream,
6565
kind_work_group_memory,
6666
kind_dynamic_work_group_memory,
67-
kind_last = kind_dynamic_work_group_memory
67+
kind_dynamic_accessor,
68+
kind_last = kind_dynamic_accessor
6869
};
6970

7071
public:

‎clang/lib/Sema/SemaSYCL.cpp

+22-5
Original file line numberDiff line numberDiff line change
@@ -94,7 +94,8 @@ bool SemaSYCL::isSyclType(QualType Ty, SYCLTypeAttr::SYCLType TypeName) {
9494

9595
static bool isSyclAccessorType(QualType Ty) {
9696
return SemaSYCL::isSyclType(Ty, SYCLTypeAttr::accessor) ||
97-
SemaSYCL::isSyclType(Ty, SYCLTypeAttr::local_accessor);
97+
SemaSYCL::isSyclType(Ty, SYCLTypeAttr::local_accessor) ||
98+
SemaSYCL::isSyclType(Ty, SYCLTypeAttr::dynamic_local_accessor);
9899
}
99100

100101
// FIXME: Accessor property lists should be modified to use compile-time
@@ -1152,7 +1153,8 @@ static QualType GetSYCLKernelObjectType(const FunctionDecl *KernelCaller) {
11521153
/// \return the target of given SYCL accessor type
11531154
static target getAccessTarget(QualType FieldTy,
11541155
const ClassTemplateSpecializationDecl *AccTy) {
1155-
if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::local_accessor))
1156+
if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::local_accessor) ||
1157+
SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::dynamic_local_accessor))
11561158
return local;
11571159

11581160
return static_cast<target>(
@@ -4815,7 +4817,15 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
48154817
int Dims = static_cast<int>(
48164818
AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue());
48174819
int Info = getAccessTarget(FieldTy, AccTy) | (Dims << 11);
4818-
Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info,
4820+
4821+
SYCLIntegrationHeader::kernel_param_kind_t ParamKind =
4822+
SYCLIntegrationHeader::kind_accessor;
4823+
4824+
if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::dynamic_local_accessor)) {
4825+
ParamKind = SYCLIntegrationHeader::kind_dynamic_accessor;
4826+
}
4827+
4828+
Header.addParamDesc(ParamKind, Info,
48194829
CurOffset +
48204830
offsetOf(RD, BC.getType()->getAsCXXRecordDecl()));
48214831
} else if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::work_group_memory)) {
@@ -4841,8 +4851,14 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
48414851
AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue());
48424852
int Info = getAccessTarget(FieldTy, AccTy) | (Dims << 11);
48434853

4844-
Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info,
4845-
CurOffset + offsetOf(FD, FieldTy));
4854+
SYCLIntegrationHeader::kernel_param_kind_t ParamKind =
4855+
SYCLIntegrationHeader::kind_accessor;
4856+
4857+
if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::dynamic_local_accessor)) {
4858+
ParamKind = SYCLIntegrationHeader::kind_dynamic_accessor;
4859+
}
4860+
4861+
Header.addParamDesc(ParamKind, Info, CurOffset + offsetOf(FD, FieldTy));
48464862
} else if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::stream)) {
48474863
addParam(FD, FieldTy, SYCLIntegrationHeader::kind_stream);
48484864
} else if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::work_group_memory)) {
@@ -6025,6 +6041,7 @@ static const char *paramKind2Str(KernelParamKind K) {
60256041
CASE(pointer);
60266042
CASE(work_group_memory);
60276043
CASE(dynamic_work_group_memory);
6044+
CASE(dynamic_accessor);
60286045
}
60296046
return "<ERROR>";
60306047

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

+13
Original file line numberDiff line numberDiff line change
@@ -680,6 +680,19 @@ __SYCL_TYPE(dynamic_work_group_memory) dynamic_work_group_memory {
680680
work_group_memory<DataT> LocalMem;
681681
};
682682

683+
template <typename DataT, int Dimensions>
684+
class __attribute__((sycl_special_class))
685+
__SYCL_TYPE(dynamic_local_accessor) dynamic_local_accessor {
686+
public:
687+
dynamic_local_accessor() = default;
688+
689+
void __init(__attribute((opencl_local)) DataT *Ptr) { this->LocalMem.__init(Ptr); }
690+
local_accessor<DataT, Dimensions> get() const { return LocalMem; }
691+
692+
private:
693+
local_accessor<DataT, Dimensions> LocalMem;
694+
};
695+
683696
template <typename T, int dimensions = 1,
684697
typename AllocatorT = int /*fake type as AllocatorT is not used*/>
685698
class buffer {

‎clang/test/CodeGenSYCL/dynamic_local_accessor.cpp

+39
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,39 @@
1+
// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o %t.ll
2+
// RUN: FileCheck < %t.ll %s --check-prefix CHECK-IR
3+
// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -fsycl-int-header=%t.h %s
4+
// RUN: FileCheck < %t.h %s --check-prefix CHECK-INT-HEADER
5+
//
6+
// Tests for dynamic_local_accessor kernel parameter using the dummy implementation in Inputs/sycl.hpp.
7+
// The first two RUN commands verify that the init call is generated with the correct arguments in LLVM IR
8+
// and the second two RUN commands verify the contents of the integration header produced by the frontend.
9+
//
10+
// CHECK-IR: define dso_local spir_kernel void @
11+
// CHECK-IR-SAME: ptr addrspace(3) noundef align 4 [[PTR:%[a-zA-Z0-9_]+]]
12+
//
13+
// CHECK-IR: [[PTR]].addr = alloca ptr addrspace(3), align 8
14+
// CHECK-IR: [[PTR]].addr.ascast = addrspacecast ptr [[PTR]].addr to ptr addrspace(4)
15+
// CHECK-IR: store ptr addrspace(3) [[PTR]], ptr addrspace(4) [[PTR]].addr.ascast, align 8
16+
// CHECK-IR: [[PTR_LOAD:%[a-zA-Z0-9_]+]] = load ptr addrspace(3), ptr addrspace(4) [[PTR]].addr.ascast, align 8
17+
//
18+
// CHECK-IR: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) noundef align 8 dereferenceable_or_null(8) %{{[a-zA-Z0-9_]+}}, ptr addrspace(3) noundef [[PTR_LOAD]])
19+
//
20+
// CHECK-INT-HEADER: const kernel_param_desc_t kernel_signatures[] = {
21+
// CHECK-INT-HEADER-NEXT: //--- _ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlNS0_4itemILi1EEEE_
22+
// CHECK-INT-HEADER-NEXT: { kernel_param_kind_t::kind_dynamic_local_accessor, {{[4,8]}}, 0 },
23+
// CHECK-INT-HEADER-EMPTY:
24+
// CHECK-INT-HEADER-NEXT: { kernel_param_kind_t::kind_invalid, -987654321, -987654321 },
25+
// CHECK-INT-HEADER-NEXT: };
26+
27+
#include "Inputs/sycl.hpp"
28+
29+
int main() {
30+
sycl::queue Q;
31+
sycl::dynamic_local_accessor<int, 1> dynLocalAcc;
32+
Q.submit([&](sycl::handler &CGH) {
33+
sycl::range<1> ndr;
34+
CGH.parallel_for(ndr, [=](sycl::item<1> it) {
35+
auto localAcc = dynLocalAcc.get();
36+
int *ptr = &localAcc; });
37+
});
38+
return 0;
39+
}

‎sycl-jit/common/include/Kernel.h

Whitespace-only changes.

‎sycl/include/sycl/accessor.hpp

+8
Original file line numberDiff line numberDiff line change
@@ -215,6 +215,7 @@
215215
namespace sycl {
216216
inline namespace _V1 {
217217
class stream;
218+
218219
namespace ext::intel::esimd::detail {
219220
// Forward declare a "back-door" access class to support ESIMD.
220221
class AccessorPrivateProxy;
@@ -227,6 +228,10 @@ template <typename DataT, int Dimensions = 1,
227228
typename PropertyListT = ext::oneapi::accessor_property_list<>>
228229
class accessor;
229230

231+
namespace ext::oneapi::experimental {
232+
template <typename, int> class dynamic_local_accessor;
233+
}
234+
230235
namespace detail {
231236

232237
template <typename... Ts>
@@ -344,6 +349,7 @@ class accessor_common {
344349
typename AccType =
345350
accessor<DataT, Dimensions, AccessMode, AccessTarget,
346351
IsPlaceholder, PropertyListT>>
352+
347353
class AccessorSubscript {
348354
static constexpr int Dims = Dimensions;
349355

@@ -2148,6 +2154,7 @@ class __SYCL_SPECIAL_CLASS local_accessor_base :
21482154
#endif
21492155
public detail::accessor_common<DataT, Dimensions, AccessMode,
21502156
access::target::local, IsPlaceholder> {
2157+
21512158
protected:
21522159
constexpr static int AdjustedDim = Dimensions == 0 ? 1 : Dimensions;
21532160

@@ -2638,6 +2645,7 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor
26382645

26392646
private:
26402647
friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy;
2648+
template<typename, int> friend class ext::oneapi::experimental::dynamic_local_accessor;
26412649
};
26422650

26432651
template <typename DataT, int Dimensions = 1,

‎sycl/include/sycl/detail/kernel_desc.hpp

+1
Original file line numberDiff line numberDiff line change
@@ -60,6 +60,7 @@ enum class kernel_param_kind_t {
6060
kind_stream = 5,
6161
kind_work_group_memory = 6,
6262
kind_dynamic_work_group_memory = 7,
63+
kind_dynamic_accessor = 8,
6364
kind_invalid = 0xf, // not a valid kernel kind
6465
};
6566

0 commit comments

Comments
 (0)