Skip to content

Commit e4423ef

Browse files
AidanBeltonSAidanBeltonS
authored
[SYCL] Add local_accessor and deprecate target::local (#6341)
This PR creates the `local_accessor` class by aliasing the accessor class with `target::local`. The motivation behind this is that `target::local` has been deprecated in favour of `local_accessor` in SYCL2020. The approach of aliasing is taken as the spec states that local_access has the same semantics and restrictions as accessor with target::local. Related issue: #4713 llvm-test-suite: intel/llvm-test-suite#1063
1 parent b32dd41 commit e4423ef

20 files changed

+391
-91
lines changed

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 41 additions & 25 deletions
Original file line numberDiff line numberDiff line change
@@ -111,6 +111,11 @@ class Util {
111111
/// \param Tmpl whether the class is template instantiation or simple record
112112
static bool isSyclType(QualType Ty, StringRef Name, bool Tmpl = false);
113113

114+
/// Checks whether given clang type is a standard SYCL API accessor class,
115+
/// the check assumes the type is templated.
116+
/// \param Ty the clang type being checked
117+
static bool isSyclAccessorType(QualType Ty);
118+
114119
/// Checks whether given clang type is a full specialization of the SYCL
115120
/// specialization constant class.
116121
static bool isSyclSpecConstantType(QualType Ty);
@@ -1021,7 +1026,11 @@ static ParamDesc makeParamDesc(ASTContext &Ctx, StringRef Name, QualType Ty) {
10211026
}
10221027

10231028
/// \return the target of given SYCL accessor type
1024-
static target getAccessTarget(const ClassTemplateSpecializationDecl *AccTy) {
1029+
static target getAccessTarget(QualType FieldTy,
1030+
const ClassTemplateSpecializationDecl *AccTy) {
1031+
if (Util::isSyclType(FieldTy, "local_accessor", true /*Tmpl*/))
1032+
return local;
1033+
10251034
return static_cast<target>(
10261035
AccTy->getTemplateArgs()[3].getAsIntegral().getExtValue());
10271036
}
@@ -1615,7 +1624,7 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
16151624
assert(Util::isSyclSpecialType(Ty) &&
16161625
"Should only be called on sycl special class types.");
16171626
const RecordDecl *RecD = Ty->getAsRecordDecl();
1618-
if (IsSIMD && !Util::isSyclType(Ty, "accessor", true /*Tmp*/))
1627+
if (IsSIMD && !Util::isSyclAccessorType(Ty))
16191628
return SemaRef.Diag(Loc.getBegin(),
16201629
diag::err_sycl_esimd_not_supported_for_type)
16211630
<< RecD;
@@ -1927,19 +1936,24 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
19271936
}
19281937

19291938
// Additional processing is required for accessor type.
1930-
void handleAccessorType(const CXXRecordDecl *RecordDecl, SourceLocation Loc) {
1939+
void handleAccessorType(QualType FieldTy, const CXXRecordDecl *RecordDecl,
1940+
SourceLocation Loc) {
19311941
handleAccessorPropertyList(Params.back(), RecordDecl, Loc);
1932-
// Get access mode of accessor.
1933-
const auto *AccessorSpecializationDecl =
1934-
cast<ClassTemplateSpecializationDecl>(RecordDecl);
1935-
const TemplateArgument &AccessModeArg =
1936-
AccessorSpecializationDecl->getTemplateArgs().get(2);
1942+
1943+
// If "accessor" type check if read only
1944+
if (Util::isSyclType(FieldTy, "accessor", true /*Tmpl*/)) {
1945+
// Get access mode of accessor.
1946+
const auto *AccessorSpecializationDecl =
1947+
cast<ClassTemplateSpecializationDecl>(RecordDecl);
1948+
const TemplateArgument &AccessModeArg =
1949+
AccessorSpecializationDecl->getTemplateArgs().get(2);
1950+
if (isReadOnlyAccessor(AccessModeArg))
1951+
Params.back()->addAttr(
1952+
SYCLAccessorReadonlyAttr::CreateImplicit(SemaRef.getASTContext()));
1953+
}
19371954

19381955
// Add implicit attribute to parameter decl when it is a read only
19391956
// SYCL accessor.
1940-
if (isReadOnlyAccessor(AccessModeArg))
1941-
Params.back()->addAttr(
1942-
SYCLAccessorReadonlyAttr::CreateImplicit(SemaRef.getASTContext()));
19431957
Params.back()->addAttr(
19441958
SYCLAccessorPtrAttr::CreateImplicit(SemaRef.getASTContext()));
19451959
}
@@ -1952,8 +1966,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
19521966
const auto *RecordDecl = FieldTy->getAsCXXRecordDecl();
19531967
assert(RecordDecl && "The type must be a RecordDecl");
19541968
llvm::StringLiteral MethodName =
1955-
KernelDecl->hasAttr<SYCLSimdAttr>() &&
1956-
Util::isSyclType(FieldTy, "accessor", true /*Tmp*/)
1969+
KernelDecl->hasAttr<SYCLSimdAttr>() && Util::isSyclAccessorType(FieldTy)
19571970
? InitESIMDMethodName
19581971
: InitMethodName;
19591972
CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, MethodName);
@@ -1978,8 +1991,8 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
19781991
// added, this code needs to be refactored to call
19791992
// handleAccessorPropertyList for each class which requires it.
19801993
if (ParamTy.getTypePtr()->isPointerType() &&
1981-
Util::isSyclType(FieldTy, "accessor", true /*Tmp*/))
1982-
handleAccessorType(RecordDecl, FD->getBeginLoc());
1994+
Util::isSyclAccessorType(FieldTy))
1995+
handleAccessorType(FieldTy, RecordDecl, FD->getBeginLoc());
19831996
}
19841997
LastParamIndex = ParamIndex;
19851998
return true;
@@ -2073,8 +2086,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
20732086
const auto *RecordDecl = FieldTy->getAsCXXRecordDecl();
20742087
assert(RecordDecl && "The type must be a RecordDecl");
20752088
llvm::StringLiteral MethodName =
2076-
KernelDecl->hasAttr<SYCLSimdAttr>() &&
2077-
Util::isSyclType(FieldTy, "accessor", true /*Tmp*/)
2089+
KernelDecl->hasAttr<SYCLSimdAttr>() && Util::isSyclAccessorType(FieldTy)
20782090
? InitESIMDMethodName
20792091
: InitMethodName;
20802092
CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, MethodName);
@@ -2093,8 +2105,8 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
20932105
// added, this code needs to be refactored to call
20942106
// handleAccessorPropertyList for each class which requires it.
20952107
if (ParamTy.getTypePtr()->isPointerType() &&
2096-
Util::isSyclType(FieldTy, "accessor", true /*Tmp*/))
2097-
handleAccessorType(RecordDecl, BS.getBeginLoc());
2108+
Util::isSyclAccessorType(FieldTy))
2109+
handleAccessorType(FieldTy, RecordDecl, BS.getBeginLoc());
20982110
}
20992111
LastParamIndex = ParamIndex;
21002112
return true;
@@ -2215,9 +2227,8 @@ class SyclKernelArgsSizeChecker : public SyclKernelFieldHandler {
22152227
const CXXRecordDecl *RecordDecl = FieldTy->getAsCXXRecordDecl();
22162228
assert(RecordDecl && "The type must be a RecordDecl");
22172229
llvm::StringLiteral MethodName =
2218-
(IsSIMD && Util::isSyclType(FieldTy, "accessor", true /*Tmp*/))
2219-
? InitESIMDMethodName
2220-
: InitMethodName;
2230+
(IsSIMD && Util::isSyclAccessorType(FieldTy)) ? InitESIMDMethodName
2231+
: InitMethodName;
22212232
CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, MethodName);
22222233
assert(InitMethod && "The type must have the __init method");
22232234
for (const ParmVarDecl *Param : InitMethod->parameters())
@@ -3124,7 +3135,7 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
31243135
"Incorrect template args for Accessor Type");
31253136
int Dims = static_cast<int>(
31263137
AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue());
3127-
int Info = getAccessTarget(AccTy) | (Dims << 11);
3138+
int Info = getAccessTarget(FieldTy, AccTy) | (Dims << 11);
31283139
Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info,
31293140
CurOffset +
31303141
offsetOf(RD, BC.getType()->getAsCXXRecordDecl()));
@@ -3134,14 +3145,14 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
31343145
bool handleSyclSpecialType(FieldDecl *FD, QualType FieldTy) final {
31353146
const auto *ClassTy = FieldTy->getAsCXXRecordDecl();
31363147
assert(ClassTy && "Type must be a C++ record type");
3137-
if (Util::isSyclType(FieldTy, "accessor", true /*Tmp*/)) {
3148+
if (Util::isSyclAccessorType(FieldTy)) {
31383149
const auto *AccTy =
31393150
cast<ClassTemplateSpecializationDecl>(FieldTy->getAsRecordDecl());
31403151
assert(AccTy->getTemplateArgs().size() >= 2 &&
31413152
"Incorrect template args for Accessor Type");
31423153
int Dims = static_cast<int>(
31433154
AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue());
3144-
int Info = getAccessTarget(AccTy) | (Dims << 11);
3155+
int Info = getAccessTarget(FieldTy, AccTy) | (Dims << 11);
31453156

31463157
Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info,
31473158
CurOffset + offsetOf(FD, FieldTy));
@@ -5195,6 +5206,11 @@ bool Util::isSyclType(QualType Ty, StringRef Name, bool Tmpl) {
51955206
return matchQualifiedTypeName(Ty, Scopes);
51965207
}
51975208

5209+
bool Util::isSyclAccessorType(QualType Ty) {
5210+
return isSyclType(Ty, "accessor", true /* Tmpl */) ||
5211+
isSyclType(Ty, "local_accessor", true /* Tmpl */);
5212+
}
5213+
51985214
bool Util::isAccessorPropertyListType(QualType Ty) {
51995215
std::array<DeclContextDesc, 5> Scopes = {
52005216
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"),

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

Lines changed: 20 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -291,6 +291,26 @@ class accessor<dataT, dimensions, accessmode, access::target::host_image, access
291291
_ImageImplT<dimensions, accessmode, access::target::host_image> impl;
292292
};
293293

294+
template <typename dataT, int dimensions>
295+
class __attribute__((sycl_special_class))
296+
local_accessor: public accessor<dataT,
297+
dimensions, access::mode::read_write,
298+
access::target::local> {
299+
public:
300+
void use(void) const {}
301+
template <typename... T>
302+
void use(T... args) {}
303+
template <typename... T>
304+
void use(T... args) const {}
305+
_ImplT<dimensions> impl;
306+
307+
private:
308+
#ifdef __SYCL_DEVICE_ONLY__
309+
void __init(__attribute__((opencl_local)) dataT *Ptr, range<dimensions> AccessRange,
310+
range<dimensions> MemRange, id<dimensions> Offset) {}
311+
#endif
312+
};
313+
294314
// TODO: Add support for image_array accessor.
295315
// template <typename dataT, int dimensions, access::mode accessmode>
296316
//class accessor<dataT, dimensions, accessmode, access::target::image_array, access::placeholder::false_t>

‎clang/test/CodeGenSYCL/kernel-arg-accessor-pointer.cpp

Lines changed: 20 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -22,6 +22,8 @@ int main() {
2222
access::placeholder::true_t>
2323
acc3;
2424

25+
local_accessor<float, 2> acc4;
26+
2527
// kernel_A parameters : int*, sycl::range<1>, sycl::range<1>, sycl::id<1>,
2628
// int*, sycl::range<1>, sycl::range<1>,sycl::id<1>.
2729
q.submit([&](handler &h) {
@@ -67,11 +69,19 @@ int main() {
6769
// Using local accessor as a kernel parameter.
6870
// kernel_arg_runtime_aligned is generated for pointers from local accessors.
6971
q.submit([&](handler &h) {
70-
h.single_task<class localAccessor>([=]() {
72+
h.single_task<class localAccessorDep>([=]() {
7173
acc3.use();
7274
});
7375
});
7476

77+
// Using local_accessor as a kernel parameter.
78+
// kernel_arg_runtime_aligned is generated for pointers from local accessors.
79+
q.submit([&](handler &h) {
80+
h.single_task<class localAccessor>([=]() {
81+
acc4.use();
82+
});
83+
});
84+
7585
// kernel_acc_raw_ptr parameters : int*, sycl::range<1>, sycl::range<1>, sycl::id<1>, int*.
7686
int *rawPtr;
7787
q.submit([&](handler &h) {
@@ -130,14 +140,22 @@ int main() {
130140
// CHECK-NOT: kernel_arg_runtime_aligned
131141
// CHECK-NOT: kernel_arg_exclusive_ptr
132142

133-
// CHECK: define {{.*}}spir_kernel void @{{.*}}localAccessor
143+
// CHECK: define {{.*}}spir_kernel void @{{.*}}localAccessorDep
134144
// CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
135145
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
136146
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
137147
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]]
138148
// CHECK-SAME: !kernel_arg_runtime_aligned ![[#ACCESSORMD2]]
139149
// CHECK-SAME: !kernel_arg_exclusive_ptr ![[#ACCESSORMD2]]
140150

151+
// CHECK: define {{.*}}spir_kernel void @{{.*}}localAccessor
152+
// CHECK-SAME: ptr addrspace(3) noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
153+
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
154+
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
155+
// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]]
156+
// CHECK-SAME: !kernel_arg_runtime_aligned ![[#ACCESSORMD2]]
157+
// CHECK-SAME: !kernel_arg_exclusive_ptr ![[#ACCESSORMD2]]
158+
141159
// Check kernel_acc_raw_ptr parameters
142160
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_acc_raw_ptr
143161
// CHECK-SAME: ptr addrspace(1) noundef readonly align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],

‎clang/test/CodeGenSYCL/no_opaque_kernel-arg-accessor-pointer.cpp

Lines changed: 19 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -22,6 +22,8 @@ int main() {
2222
access::placeholder::true_t>
2323
acc3;
2424

25+
local_accessor<float, 2> acc4;
26+
2527
// kernel_A parameters : int*, sycl::range<1>, sycl::range<1>, sycl::id<1>,
2628
// int*, sycl::range<1>, sycl::range<1>,sycl::id<1>.
2729
q.submit([&](handler &h) {
@@ -67,11 +69,19 @@ int main() {
6769
// Using local accessor as a kernel parameter.
6870
// kernel_arg_runtime_aligned is generated for pointers from local accessors.
6971
q.submit([&](handler &h) {
70-
h.single_task<class localAccessor>([=]() {
72+
h.single_task<class localAccessorDep>([=]() {
7173
acc3.use();
7274
});
7375
});
7476

77+
// Using local accessor as a kernel parameter.
78+
// kernel_arg_runtime_aligned is generated for pointers from local accessors.
79+
q.submit([&](handler &h) {
80+
h.single_task<class localAccessor>([=]() {
81+
acc4.use();
82+
});
83+
});
84+
7585
// kernel_acc_raw_ptr parameters : int*, sycl::range<1>, sycl::range<1>, sycl::id<1>, int*.
7686
int *rawPtr;
7787
q.submit([&](handler &h) {
@@ -125,13 +135,20 @@ int main() {
125135
// CHECK-SAME: float addrspace(1)* noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]]
126136
// CHECK-NOT: kernel_arg_runtime_aligned
127137

128-
// CHECK: define {{.*}}spir_kernel void @{{.*}}localAccessor
138+
// CHECK: define {{.*}}spir_kernel void @{{.*}}localAccessorDep
129139
// CHECK-SAME: float addrspace(1)* noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
130140
// CHECK-SAME: %"struct.sycl::_V1::range.5"* noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
131141
// CHECK-SAME: %"struct.sycl::_V1::range.5"* noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
132142
// CHECK-SAME: %"struct.sycl::_V1::id.6"* noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]]
133143
// CHECK-SAME: !kernel_arg_runtime_aligned ![[#RTALIGNED2]]
134144

145+
// CHECK: define {{.*}}spir_kernel void @{{.*}}localAccessor
146+
// CHECK-SAME: float addrspace(3)* noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
147+
// CHECK-SAME: %"struct.sycl::_V1::range.5"* noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
148+
// CHECK-SAME: %"struct.sycl::_V1::range.5"* noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
149+
// CHECK-SAME: %"struct.sycl::_V1::id.6"* noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]]
150+
// CHECK-SAME: !kernel_arg_runtime_aligned ![[#RTALIGNED2]]
151+
135152
// Check kernel_acc_raw_ptr parameters
136153
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_acc_raw_ptr
137154
// CHECK-SAME: i32 addrspace(1)* noundef readonly align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],

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

Lines changed: 20 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -207,6 +207,26 @@ class __attribute__((sycl_special_class)) accessor<dataT, dimensions, accessmode
207207
#endif
208208
};
209209

210+
template <typename dataT, int dimensions>
211+
class __attribute__((sycl_special_class))
212+
local_accessor: public accessor<dataT,
213+
dimensions, access::mode::read_write,
214+
access::target::local> {
215+
public:
216+
void use(void) const {}
217+
template <typename... T>
218+
void use(T... args) {}
219+
template <typename... T>
220+
void use(T... args) const {}
221+
_ImplT<dimensions> impl;
222+
223+
private:
224+
#ifdef __SYCL_DEVICE_ONLY__
225+
void __init(__attribute__((opencl_local)) dataT *Ptr, range<dimensions> AccessRange,
226+
range<dimensions> MemRange, id<dimensions> Offset) {}
227+
#endif
228+
};
229+
210230
struct sampler_impl {
211231
#ifdef __SYCL_DEVICE_ONLY__
212232
__ocl_sampler_t m_Sampler;

‎clang/test/SemaSYCL/accessors-targets.cpp

Lines changed: 11 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -11,7 +11,9 @@ int main() {
1111
// Access work-group local memory with read and write access.
1212
sycl::accessor<int, 1, sycl::access::mode::read_write,
1313
sycl::access::target::local>
14-
local_acc;
14+
local_acc_dep;
15+
// Access work-group local memory with read and write access.
16+
sycl::local_accessor<int, 1> local_acc;
1517
// Access buffer via global memory with read and write access.
1618
sycl::accessor<int, 1, sycl::access::mode::read_write,
1719
sycl::access::target::global_buffer>
@@ -21,6 +23,13 @@ int main() {
2123
sycl::access::target::constant_buffer>
2224
constant_acc;
2325

26+
q.submit([&](sycl::handler &h) {
27+
h.single_task<class use_local_dep>(
28+
[=] {
29+
local_acc_dep.use();
30+
});
31+
});
32+
2433
q.submit([&](sycl::handler &h) {
2534
h.single_task<class use_local>(
2635
[=] {
@@ -42,6 +51,7 @@ int main() {
4251
});
4352
});
4453
}
54+
// CHECK: {{.*}}use_local_dep{{.*}} 'void (__local int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
4555
// CHECK: {{.*}}use_local{{.*}} 'void (__local int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
4656
// CHECK: {{.*}}use_global{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
4757
// CHECK: {{.*}}use_constant{{.*}} 'void (__constant int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'

‎sycl/include/sycl/access/access.hpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -17,7 +17,7 @@ namespace access {
1717
enum class target {
1818
global_buffer __SYCL2020_DEPRECATED("use 'target::device' instead") = 2014,
1919
constant_buffer = 2015,
20-
local = 2016,
20+
local __SYCL2020_DEPRECATED("use `local_accessor` instead") = 2016,
2121
image = 2017,
2222
host_buffer = 2018,
2323
host_image = 2019,

0 commit comments

Comments
 (0)