Skip to content

[SYCL][Graph] Implement dynamic local accessors #18437

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 2 commits into
base: sycl
Choose a base branch
from
+1,082 −227
Open

[SYCL][Graph] Implement dynamic local accessors #18437

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
cb78e9d
Implement dynamic local accessors
fabiomestre Apr 11, 2025
d9fc8f7
Address review comment
fabiomestre May 14, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 4 additions & 2 deletions clang/include/clang/Basic/Attr.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1621,12 +1621,14 @@ def SYCLType: InheritableAttr {
let Subjects = SubjectList<[CXXRecord, Enum], ErrorDiag>;
let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
let Args = [EnumArgument<"Type", "SYCLType", /*is_string=*/true,
["accessor", "local_accessor", "work_group_memory", "dynamic_work_group_memory",
["accessor", "local_accessor", "dynamic_local_accessor",
"work_group_memory", "dynamic_work_group_memory",
"specialization_id", "kernel_handler", "buffer_location",
"no_alias", "accessor_property_list", "group",
"private_memory", "aspect", "annotated_ptr", "annotated_arg",
"stream", "sampler", "host_pipe", "multi_ptr"],
["accessor", "local_accessor", "work_group_memory", "dynamic_work_group_memory",
["accessor", "local_accessor", "dynamic_local_accessor",
"work_group_memory", "dynamic_work_group_memory",
"specialization_id", "kernel_handler", "buffer_location",
"no_alias", "accessor_property_list", "group",
"private_memory", "aspect", "annotated_ptr", "annotated_arg",
Expand Down
3 changes: 2 additions & 1 deletion clang/include/clang/Sema/SemaSYCL.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -64,7 +64,8 @@ class SYCLIntegrationHeader {
kind_stream,
kind_work_group_memory,
kind_dynamic_work_group_memory,
kind_last = kind_dynamic_work_group_memory
kind_dynamic_accessor,
kind_last = kind_dynamic_accessor
};

public:
Expand Down
27 changes: 22 additions & 5 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -94,7 +94,8 @@ bool SemaSYCL::isSyclType(QualType Ty, SYCLTypeAttr::SYCLType TypeName) {

static bool isSyclAccessorType(QualType Ty) {
return SemaSYCL::isSyclType(Ty, SYCLTypeAttr::accessor) ||
SemaSYCL::isSyclType(Ty, SYCLTypeAttr::local_accessor);
SemaSYCL::isSyclType(Ty, SYCLTypeAttr::local_accessor) ||
SemaSYCL::isSyclType(Ty, SYCLTypeAttr::dynamic_local_accessor);
}

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

return static_cast<target>(
Expand Down Expand Up @@ -4815,7 +4817,15 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
int Dims = static_cast<int>(
AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue());
int Info = getAccessTarget(FieldTy, AccTy) | (Dims << 11);
Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info,

SYCLIntegrationHeader::kernel_param_kind_t ParamKind =
SYCLIntegrationHeader::kind_accessor;

if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::dynamic_local_accessor)) {
ParamKind = SYCLIntegrationHeader::kind_dynamic_accessor;
}

Header.addParamDesc(ParamKind, Info,
Comment on lines +4820 to +4828
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Nit, this could be made into ternary operator. Also same below.

Suggested change
SYCLIntegrationHeader::kernel_param_kind_t ParamKind =
SYCLIntegrationHeader::kind_accessor;
if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::dynamic_local_accessor)) {
ParamKind = SYCLIntegrationHeader::kind_dynamic_accessor;
}
Header.addParamDesc(ParamKind, Info,
SYCLIntegrationHeader::kernel_param_kind_t ParamKind =
SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::dynamic_local_accessor)
? SYCLIntegrationHeader::kind_dynamic_accessor
: SYCLIntegrationHeader::kind_accessor;
Header.addParamDesc(ParamKind, Info,

CurOffset +
offsetOf(RD, BC.getType()->getAsCXXRecordDecl()));
} else if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::work_group_memory)) {
Expand All @@ -4841,8 +4851,14 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue());
int Info = getAccessTarget(FieldTy, AccTy) | (Dims << 11);

Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info,
CurOffset + offsetOf(FD, FieldTy));
SYCLIntegrationHeader::kernel_param_kind_t ParamKind =
SYCLIntegrationHeader::kind_accessor;

if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::dynamic_local_accessor)) {
ParamKind = SYCLIntegrationHeader::kind_dynamic_accessor;
}

Header.addParamDesc(ParamKind, Info, CurOffset + offsetOf(FD, FieldTy));
} else if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::stream)) {
addParam(FD, FieldTy, SYCLIntegrationHeader::kind_stream);
} else if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::work_group_memory)) {
Expand Down Expand Up @@ -6025,6 +6041,7 @@ static const char *paramKind2Str(KernelParamKind K) {
CASE(pointer);
CASE(work_group_memory);
CASE(dynamic_work_group_memory);
CASE(dynamic_accessor);
}
return "<ERROR>";

Expand Down
20 changes: 20 additions & 0 deletions clang/test/CodeGenSYCL/Inputs/sycl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -431,6 +431,9 @@ local_accessor: public accessor<dataT,
#ifdef __SYCL_DEVICE_ONLY__
void __init(__attribute__((opencl_local)) dataT *Ptr, range<dimensions> AccessRange,
range<dimensions> MemRange, id<dimensions> Offset) {}

template <typename, int>
friend class dynamic_local_accessor;
#endif
};

Expand Down Expand Up @@ -680,6 +683,23 @@ __SYCL_TYPE(dynamic_work_group_memory) dynamic_work_group_memory {
work_group_memory<DataT> LocalMem;
};

template <typename DataT, int Dimensions>
class __attribute__((sycl_special_class))
__SYCL_TYPE(dynamic_local_accessor) dynamic_local_accessor {
public:
dynamic_local_accessor() = default;

void __init(__attribute__((opencl_local)) DataT *Ptr,
range<Dimensions> AccessRange, range<Dimensions> range,
id<Dimensions> id) {
this->LocalMem.__init(Ptr, AccessRange, range, id);
}
local_accessor<DataT, Dimensions> get() const { return LocalMem; }

private:
local_accessor<DataT, Dimensions> LocalMem;
};

template <typename T, int dimensions = 1,
typename AllocatorT = int /*fake type as AllocatorT is not used*/>
class buffer {
Expand Down
40 changes: 40 additions & 0 deletions clang/test/CodeGenSYCL/dynamic_local_accessor.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,40 @@
// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o %t.ll
// RUN: FileCheck < %t.ll %s --check-prefix CHECK-IR
// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -fsycl-int-header=%t.h %s
// RUN: FileCheck < %t.h %s --check-prefix CHECK-INT-HEADER
//
// Tests for dynamic_local_accessor kernel parameter using the dummy implementation in Inputs/sycl.hpp.
// The first two RUN commands verify that the init call is generated with the correct arguments in LLVM IR
// and the second two RUN commands verify the contents of the integration header produced by the frontend.
//
// CHECK-IR: define dso_local spir_kernel void @
// CHECK-IR-SAME: ptr addrspace(3) noundef align 4 [[PTR:%[a-zA-Z0-9_]+]]
//
// CHECK-IR: [[PTR]].addr = alloca ptr addrspace(3), align 8
// CHECK-IR: [[PTR]].addr.ascast = addrspacecast ptr [[PTR]].addr to ptr addrspace(4)
// CHECK-IR: store ptr addrspace(3) [[PTR]], ptr addrspace(4) [[PTR]].addr.ascast, align 8
// CHECK-IR: [[PTR_LOAD:%[a-zA-Z0-9_]+]] = load ptr addrspace(3), ptr addrspace(4) [[PTR]].addr.ascast, align 8
//
// CHECK-IR: call spir_func void @{{.*}}__init{{.*}}(ptr addrspace(4) noundef align 4 dereferenceable_or_null(24) %{{[a-zA-Z0-9_]+}}, ptr addrspace(3) noundef %{{[a-zA-Z0-9_]+}}, ptr noundef byval(%"struct.sycl::_V1::range") align 4 %agg.{{[a-zA-Z0-9_]+}}.ascast.ascast, ptr noundef byval(%"struct.sycl::_V1::range") align 4 %agg.{{[a-zA-Z0-9_]+}}.ascast.ascast, ptr noundef byval(%"struct.sycl::_V1::id") align 4 %agg.{{[a-zA-Z0-9_]+}}.ascast.ascast) #{{[0-9_]+}}
//
// CHECK-INT-HEADER: const kernel_param_desc_t kernel_signatures[] = {
// CHECK-INT-HEADER-NEXT: //--- _ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlNS0_4itemILi1EEEE_
// CHECK-INT-HEADER-NEXT: { kernel_param_kind_t::kind_dynamic_accessor, 4064, 0 },
// CHECK-INT-HEADER-EMPTY:
// CHECK-INT-HEADER-NEXT: { kernel_param_kind_t::kind_invalid, -987654321, -987654321 },
// CHECK-INT-HEADER-NEXT: };

#include "Inputs/sycl.hpp"

int main() {
sycl::queue Q;
sycl::dynamic_local_accessor<int, 1> dynLocalAcc;
Q.submit([&](sycl::handler &CGH) {
sycl::range<1> ndr;
CGH.parallel_for(ndr, [=](sycl::item<1> it) {
auto localAcc = dynLocalAcc.get();
auto* ptr = &localAcc;
});
});
return 0;
}
6 changes: 6 additions & 0 deletions sycl/include/sycl/accessor.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -227,6 +227,10 @@ template <typename DataT, int Dimensions = 1,
typename PropertyListT = ext::oneapi::accessor_property_list<>>
class accessor;

namespace ext::oneapi::experimental {
template <typename, int> class dynamic_local_accessor;
}

namespace detail {

template <typename... Ts>
Expand Down Expand Up @@ -2638,6 +2642,8 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor

private:
friend class sycl::ext::intel::esimd::detail::AccessorPrivateProxy;
template <typename, int>
friend class ext::oneapi::experimental::dynamic_local_accessor;
};

template <typename DataT, int Dimensions = 1,
Expand Down
1 change: 1 addition & 0 deletions sycl/include/sycl/detail/kernel_desc.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -60,6 +60,7 @@ enum class kernel_param_kind_t {
kind_stream = 5,
kind_work_group_memory = 6,
kind_dynamic_work_group_memory = 7,
kind_dynamic_accessor = 8,
kind_invalid = 0xf, // not a valid kernel kind
};

Expand Down
Loading
Loading