Skip to content

Commit 819a973

Browse files
MrSidimsMrSidims
committed
[SYCL][FPGA] Add clang support for buffer_location property
This is a compiler-time known accessor property which serves as an optimization hint for a compiler on where exactly buffer was allocated. This is needed when a board has multiple disjoint global memories that must be managed explicitly by a programmer. When the property is added as a template parameter of an accessor - SemaSYCL will implicitly add ``intelfpga::kernel_arg_buffer_location`` attribute to an OpenCL kernel generated from SYCL kernel object. It is not allowed to use the attribute explicitly in SYCL code. When the attribute is applied, clang generates metadata attached to OpenCL kernel. Number of values stored in the metadata is the same as number of kernel parameters. Order of metadata values is following the order of pointer kernel parameters. Metadata values are of an integer type and is being set accordingly values passed through accessor property ``buffer_location``. This values are mapped in hardware backend to the actual locations of buffers (DDR, QDR etc). Default value passed in the metadata is '-1'. Signed-off-by: Dmitry Sidorov <dmitry.sidorov@intel.com>
1 parent 0fdeb61 commit 819a973

14 files changed

+223
-34
lines changed

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

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1188,6 +1188,28 @@ def SYCLDeviceIndirectlyCallable : InheritableAttr {
11881188
let PragmaAttributeSupport = 0;
11891189
}
11901190

1191+
def SYCLIntelBufferLocation : InheritableAttr {
1192+
let Spellings = [CXX11<"intelfpga","kernel_arg_buffer_location">];
1193+
let LangOpts = [SYCLIsDevice, SYCLIsHost];
1194+
let Subjects = SubjectList<[Function], ErrorDiag>;
1195+
1196+
let AdditionalMembers = [{
1197+
std::vector<size_t> ActualArgs;
1198+
1199+
void setActualArgs(std::vector<size_t> ArgVec) {
1200+
ActualArgs = ArgVec;
1201+
}
1202+
1203+
std::vector<size_t> getActualArgs() const {
1204+
return ActualArgs;
1205+
}
1206+
}];
1207+
1208+
let Documentation = [SYCLIntelBufferLocationAttrDocs];
1209+
let HasCustomParsing = 1;
1210+
let PragmaAttributeSupport = 0;
1211+
}
1212+
11911213
def SYCLIntelKernelArgsRestrict : InheritableAttr {
11921214
let Spellings = [ CXX11<"intel", "kernel_args_restrict"> ];
11931215
let Subjects = SubjectList<[Function], ErrorDiag>;

‎clang/include/clang/Basic/AttrDocs.td

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1994,6 +1994,24 @@ can be lowered.
19941994
}];
19951995
}
19961996

1997+
def SYCLIntelBufferLocationAttrDocs : Documentation {
1998+
let Category = DocCatFunction;
1999+
let Heading = "kernel_args_buffer_location";
2000+
let Content = [{
2001+
The attribute ``intelfpga::kernel_arg_buffer_location`` is being implicitly
2002+
applied to an OpenCL kernel generated from SYCL kernel object. It is not allowed
2003+
to use the attribute explicitly in SYCL code.
2004+
2005+
When the attribute is applied, clang generates metadata attached to OpenCL
2006+
kernel. Number of values stored in the metadata is the same as number of kernel
2007+
parameters. Order of metadata values is following the order of pointer
2008+
kernel parameters. Metadata values are of an integer type and is being set
2009+
accordingly values passed through accessor property ``buffer_location``. This
2010+
values are mapped in hardware backend to the actual locations of buffers
2011+
(DDR, QDR etc). Default value passed in the metadata is '-1'.
2012+
}];
2013+
}
2014+
19972015
def SYCLIntelKernelArgsRestrictDocs : Documentation {
19982016
let Category = DocCatVariable;
19992017
let Heading = "kernel_args_restrict";

‎clang/include/clang/Basic/DiagnosticSemaKinds.td

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10982,6 +10982,8 @@ def warn_sycl_implicit_decl
1098210982
def warn_sycl_restrict_recursion
1098310983
: Warning<"SYCL kernel cannot call a recursive function">,
1098410984
InGroup<SyclStrict>, DefaultError;
10985+
def warn_sycl_implicit_attr_usage : Warning <
10986+
"%0 attribute cannot be used explicitly">, InGroup<IgnoredAttributes>;
1098510987
def err_ivdep_duplicate_arg : Error<
1098610988
"duplicate argument to 'ivdep'. attribute requires one or both of a safelen "
1098710989
"and array">;

‎clang/lib/CodeGen/CodeGenFunction.cpp

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -660,6 +660,16 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD,
660660
if (A->getEnabled())
661661
Fn->setMetadata("no_global_work_offset", llvm::MDNode::get(Context, {}));
662662
}
663+
664+
if (const SYCLIntelBufferLocationAttr *A =
665+
FD->getAttr<SYCLIntelBufferLocationAttr>()) {
666+
std::vector<size_t> Args = A->getActualArgs();
667+
std::vector<llvm::Metadata *> AttrMDArgs;
668+
for (auto A : Args)
669+
AttrMDArgs.push_back(llvm::ConstantAsMetadata::get(Builder.getInt32(A)));
670+
Fn->setMetadata("kernel_arg_buffer_location",
671+
llvm::MDNode::get(Context, AttrMDArgs));
672+
}
663673
}
664674

665675
/// Determine whether the function F ends with a return stmt.

‎clang/lib/Sema/SemaDeclAttr.cpp

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3082,6 +3082,12 @@ static void handleMaxGlobalWorkDimAttr(Sema &S, Decl *D,
30823082
S.Context, Attr, MaxGlobalWorkDim));
30833083
}
30843084

3085+
// Handles kernel_arg_buffer_location attr.
3086+
static void handleBufferLocationAttr(Sema &S, Decl *D, const ParsedAttr &Attr) {
3087+
S.Diag(Attr.getLoc(), diag::warn_sycl_implicit_attr_usage)
3088+
<< Attr;
3089+
}
3090+
30853091
static void handleVecTypeHint(Sema &S, Decl *D, const ParsedAttr &AL) {
30863092
if (!AL.hasParsedType()) {
30873093
S.Diag(AL.getLoc(), diag::err_attribute_wrong_number_arguments) << AL << 1;
@@ -7789,6 +7795,9 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
77897795
case ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset:
77907796
handleNoGlobalWorkOffsetAttr(S, D, AL);
77917797
break;
7798+
case ParsedAttr::AT_SYCLIntelBufferLocation:
7799+
handleBufferLocationAttr(S, D, AL);
7800+
break;
77927801
case ParsedAttr::AT_VecTypeHint:
77937802
handleVecTypeHint(S, D, AL);
77947803
break;

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 74 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -80,6 +80,10 @@ class Util {
8080
/// half class.
8181
static bool isSyclHalfType(const QualType &Ty);
8282

83+
/// Checks whether given clang type is a full specialization of the SYCL
84+
/// property_list class.
85+
static bool isSyclBufferLocation(const QualType &Ty);
86+
8387
/// Checks whether given clang type is a standard SYCL API class with given
8488
/// name.
8589
/// \param Ty the clang type being checked
@@ -1171,23 +1175,28 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
11711175
// Holds the last handled field's first parameter. This doesn't store an
11721176
// iterator as push_back invalidates iterators.
11731177
size_t LastParamIndex = 0;
1178+
// This vector stores information about buffer location. If no buffer_location
1179+
// property of an accessor is set - the appropriate value stored in the
1180+
// vector = -1.
1181+
std::vector<size_t> BufferLocationMD;
11741182

1175-
void addParam(const FieldDecl *FD, QualType FieldTy) {
1183+
void addParam(const FieldDecl *FD, QualType FieldTy, size_t LocationID = -1) {
11761184
const ConstantArrayType *CAT =
11771185
SemaRef.getASTContext().getAsConstantArrayType(FieldTy);
11781186
if (CAT)
11791187
FieldTy = CAT->getElementType();
11801188
ParamDesc newParamDesc = makeParamDesc(FD, FieldTy);
1181-
addParam(newParamDesc, FieldTy);
1189+
addParam(newParamDesc, FieldTy, LocationID);
11821190
}
11831191

1184-
void addParam(const CXXBaseSpecifier &BS, QualType FieldTy) {
1192+
void addParam(const CXXBaseSpecifier &BS, QualType FieldTy,
1193+
size_t LocationID = -1) {
11851194
ParamDesc newParamDesc =
11861195
makeParamDesc(SemaRef.getASTContext(), BS, FieldTy);
1187-
addParam(newParamDesc, FieldTy);
1196+
addParam(newParamDesc, FieldTy, LocationID);
11881197
}
11891198

1190-
void addParam(ParamDesc newParamDesc, QualType FieldTy) {
1199+
void addParam(ParamDesc newParamDesc, QualType FieldTy, size_t LocationID) {
11911200
// Create a new ParmVarDecl based on the new info.
11921201
auto *NewParam = ParmVarDecl::Create(
11931202
SemaRef.getASTContext(), KernelDecl, SourceLocation(), SourceLocation(),
@@ -1198,13 +1207,41 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
11981207

11991208
LastParamIndex = Params.size();
12001209
Params.push_back(NewParam);
1210+
BufferLocationMD.push_back(LocationID);
1211+
}
1212+
1213+
// Obtain an integer value stored in a template parameter of buffer_location
1214+
// property to pass it to buffer_location kernel attribute
1215+
size_t handleBufferLocationProperty(QualType FieldTy) {
1216+
const auto *AccTy =
1217+
cast<ClassTemplateSpecializationDecl>(FieldTy->getAsRecordDecl());
1218+
1219+
// TODO: when SYCL headers' part is ready - replace this 'if' with an assert
1220+
if (AccTy->getTemplateArgs().size() < 6)
1221+
return -1;
1222+
1223+
// TODO: at this point of time it's unclear, what representation in LLVM IR
1224+
// is going to be for other compile time known accessor properties, hence
1225+
// it's not clear, how handle them in SemaSYCL. But in general property_list
1226+
// is a parameter pack and shall be handled appropriately.
1227+
const auto Prop =
1228+
cast<TemplateArgument>(AccTy->getTemplateArgs()[5]);
1229+
QualType PropTy = Prop.getAsType();
1230+
if (!Util::isSyclBufferLocation(PropTy))
1231+
return -1;
1232+
1233+
const auto *PropDecl = cast<ClassTemplateSpecializationDecl>(
1234+
PropTy->getAsRecordDecl());
1235+
return static_cast<int>(
1236+
PropDecl->getTemplateArgs()[0].getAsIntegral().getExtValue());
12011237
}
12021238

12031239
// All special SYCL objects must have __init method. We extract types for
12041240
// kernel parameters from __init method parameters. We will use __init method
12051241
// and kernel parameters which we build here to initialize special objects in
12061242
// the kernel body.
1207-
bool handleSpecialType(FieldDecl *FD, QualType FieldTy) {
1243+
bool handleSpecialType(FieldDecl *FD, QualType FieldTy,
1244+
bool isAccessorType = false) {
12081245
const auto *RecordDecl = FieldTy->getAsCXXRecordDecl();
12091246
assert(RecordDecl && "The accessor/sampler must be a RecordDecl");
12101247
CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, InitMethodName);
@@ -1213,8 +1250,17 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
12131250
// Don't do -1 here because we count on this to be the first parameter added
12141251
// (if any).
12151252
size_t ParamIndex = Params.size();
1216-
for (const ParmVarDecl *Param : InitMethod->parameters())
1217-
addParam(FD, Param->getType().getCanonicalType());
1253+
auto ParamIt = InitMethod->parameters().begin();
1254+
if (*ParamIt) {
1255+
// Add meaningful argument (not '-1') to buffer_location attribute only
1256+
// for an accessor pointer
1257+
size_t BufferLocAttrArg =
1258+
isAccessorType ? handleBufferLocationProperty(FieldTy) : -1;
1259+
addParam(FD, (*ParamIt)->getType().getCanonicalType(), BufferLocAttrArg);
1260+
++ParamIt;
1261+
for (; ParamIt != InitMethod->parameters().end(); ++ParamIt)
1262+
addParam(FD, (*ParamIt)->getType().getCanonicalType(), -1);
1263+
}
12181264
LastParamIndex = ParamIndex;
12191265
return true;
12201266
}
@@ -1270,6 +1316,11 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
12701316
KernelDecl->setType(FuncType);
12711317
KernelDecl->setParams(Params);
12721318

1319+
// Add SYCLIntelBufferLocationAttr to the kernel declaration
1320+
auto *BufferLocAttr = SYCLIntelBufferLocationAttr::CreateImplicit(Ctx);
1321+
BufferLocAttr->setActualArgs(BufferLocationMD);
1322+
KernelDecl->addAttr(BufferLocAttr);
1323+
12731324
if (ArgChecker.isValid())
12741325
SemaRef.addSyclDeviceDecl(KernelDecl);
12751326
}
@@ -1285,13 +1336,13 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
12851336
// (if any).
12861337
size_t ParamIndex = Params.size();
12871338
for (const ParmVarDecl *Param : InitMethod->parameters())
1288-
addParam(BS, Param->getType().getCanonicalType());
1339+
addParam(BS, Param->getType().getCanonicalType(), 42);
12891340
LastParamIndex = ParamIndex;
12901341
return true;
12911342
}
12921343

12931344
bool handleSyclAccessorType(FieldDecl *FD, QualType FieldTy) final {
1294-
return handleSpecialType(FD, FieldTy);
1345+
return handleSpecialType(FD, FieldTy, /*isAccessorType*/ true);
12951346
}
12961347

12971348
bool handleSyclSamplerType(FieldDecl *FD, QualType FieldTy) final {
@@ -2820,6 +2871,19 @@ bool Util::isSyclSpecConstantType(const QualType &Ty) {
28202871
return matchQualifiedTypeName(Ty, Scopes);
28212872
}
28222873

2874+
bool Util::isSyclBufferLocation(const QualType &Ty) {
2875+
const StringRef &Name = "buffer_location";
2876+
std::array<DeclContextDesc, 4> Scopes = {
2877+
Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"},
2878+
Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"},
2879+
// TODO: this doesn't belong to property namespace, instead it shall be
2880+
// in its own namespace. Change it, when the actual implementation in SYCL
2881+
// headers is ready
2882+
Util::DeclContextDesc{clang::Decl::Kind::Namespace, "property"},
2883+
Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization, Name}};
2884+
return matchQualifiedTypeName(Ty, Scopes);
2885+
}
2886+
28232887
bool Util::isSyclType(const QualType &Ty, StringRef Name, bool Tmpl) {
28242888
Decl::Kind ClassDeclKind =
28252889
Tmpl ? Decl::Kind::ClassTemplateSpecialization : Decl::Kind::CXXRecord;

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

Lines changed: 21 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -71,15 +71,22 @@ enum prop_type {
7171
base_prop
7272
};
7373

74+
// Compile time known accessor property
75+
// TODO: this doesn't belong to property namespace, instead it shall be in its
76+
// own namespace. Change it, when the actual implementation in SYCL headers is
77+
// ready
78+
template<int>
79+
class buffer_location {};
80+
7481
struct property_base {
7582
virtual prop_type type() const = 0;
7683
};
7784
} // namespace property
7885

86+
template<typename... properties>
7987
class property_list {
8088
public:
81-
template <typename... propertyTN>
82-
property_list(propertyTN... props) {}
89+
property_list(properties... props) {}
8390

8491
template <typename propertyT>
8592
bool has_property() const { return true; }
@@ -127,7 +134,8 @@ struct _ImplT {
127134

128135
template <typename dataT, int dimensions, access::mode accessmode,
129136
access::target accessTarget = access::target::global_buffer,
130-
access::placeholder isPlaceholder = access::placeholder::false_t>
137+
access::placeholder isPlaceholder = access::placeholder::false_t,
138+
typename propertyListT = property_list<>>
131139
class accessor {
132140

133141
public:
@@ -141,6 +149,8 @@ class accessor {
141149
private:
142150
void __init(__attribute__((opencl_global)) dataT *Ptr, range<dimensions> AccessRange,
143151
range<dimensions> MemRange, id<dimensions> Offset) {}
152+
153+
propertyListT prop_list;
144154
};
145155

146156
template <int dimensions, access::mode accessmode, access::target accesstarget>
@@ -326,7 +336,8 @@ const stream& operator<<(const stream &S, T&&) {
326336
}
327337

328338
template <typename T, int dimensions = 1,
329-
typename AllocatorT = int /*fake type as AllocatorT is not used*/>
339+
typename AllocatorT = int /*fake type as AllocatorT is not used*/,
340+
typename... properties>
330341
class buffer {
331342
public:
332343
using value_type = T;
@@ -338,13 +349,13 @@ class buffer {
338349
buffer(ParamTypes... args) {} // fake constructor
339350

340351
buffer(const range<dimensions> &bufferRange,
341-
const property_list &propList = {}) {}
352+
const property_list<properties...> &propList = {}) {}
342353

343354
buffer(T *hostData, const range<dimensions> &bufferRange,
344-
const property_list &propList = {}) {}
355+
const property_list<properties...> &propList = {}) {}
345356

346357
buffer(const T *hostData, const range<dimensions> &bufferRange,
347-
const property_list &propList = {}) {}
358+
const property_list<properties...> &propList = {}) {}
348359

349360
buffer(const buffer &rhs) = default;
350361

@@ -412,11 +423,12 @@ enum class image_channel_type : unsigned int {
412423
fp32
413424
};
414425

415-
template <int dimensions = 1, typename AllocatorT = int>
426+
template <int dimensions = 1, typename AllocatorT = int, typename... properties>
416427
class image {
417428
public:
418429
image(image_channel_order Order, image_channel_type Type,
419-
const range<dimensions> &Range, const property_list &PropList = {}) {}
430+
const range<dimensions> &Range,
431+
const property_list<properties...> &PropList = {}) {}
420432

421433
/* -- common interface members -- */
422434

‎clang/test/CodeGenSYCL/accessor_inheritance.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -67,13 +67,13 @@ int main() {
6767
// CHECK: [[ACC_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base* [[BITCAST]], i32 0, i32 2
6868
// CHECK: [[ACC1_AS_CAST:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}cl::sycl::accessor"* [[ACC_FIELD]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)*
6969
// Default constructor call
70-
// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0EEC1Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC1_AS_CAST]])
70+
// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_13property_listIJEEEEC1Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC1_AS_CAST]])
7171
// CHECK: [[BITCAST1:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured* [[GEP]] to i8*
72-
// CHECK: [[GEP1:%[a-zA-Z0-9_]+]] = getelementptr inbounds i8, i8* [[BITCAST1]], i64 20
72+
// CHECK: [[GEP1:%[a-zA-Z0-9_]+]] = getelementptr inbounds i8, i8* [[BITCAST1]], i64 24
7373
// CHECK: [[BITCAST2:%[a-zA-Z0-9_]+]] = bitcast i8* [[GEP1]] to %"class{{.*}}cl::sycl::accessor"*
7474
// CHECK: [[ACC2_AS_CAST:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}cl::sycl::accessor"* [[BITCAST2]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)*
7575
// Default constructor call
76-
// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0EEC2Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC2_AS_CAST]])
76+
// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_13property_listIJEEEEC2Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC2_AS_CAST]])
7777

7878
// CHECK C field initialization
7979
// CHECK: [[FIELD_C:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Captured, %struct{{.*}}Captured* [[GEP]], i32 0, i32 2

‎clang/test/CodeGenSYCL/buffer_location.cpp

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,18 @@
1+
// RUN: %clang_cc1 -fsycl -fsycl-is-device -I %S/Inputs -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s
2+
3+
// CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE15kernel_function{{.*}} !kernel_arg_buffer_location ![[MDBL:[0-9]+]]
4+
// CHECK: ![[MDBL]] = !{i32 3, i32 -1, i32 -1, i32 -1}
5+
6+
#include "sycl.hpp"
7+
8+
int main() {
9+
cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
10+
cl::sycl::access::target::global_buffer,
11+
cl::sycl::access::placeholder::false_t,
12+
cl::sycl::property::buffer_location<3>> accessorA;
13+
cl::sycl::kernel_single_task<class kernel_function>(
14+
[=]() {
15+
accessorA.use();
16+
});
17+
return 0;
18+
}

0 commit comments

Comments
 (0)