Skip to content

Commit e4791d1

Browse files
alexbatashevalexbatashev
authored
[SYCL] Remove free function queries call detection from FE (#5178)
Removing free function call detection from clang front-end dramatically improves compile times of device-side code (up to 3 times in some cases). This has the following implications: 1. Range rounding is no longer disabled based on free-function queries. Users are encouraged to disable range rounding manually. 2. Free function queries are no longer supported on host device.
1 parent ecb568e commit e4791d1

File tree

13 files changed

+46
-493
lines changed

13 files changed

+46
-493
lines changed

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

Lines changed: 0 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -359,13 +359,6 @@ class SYCLIntegrationHeader {
359359
Itr->updateKernelNames(Name, StableName);
360360
}
361361

362-
/// Note which free functions (this_id, this_item, etc) are called within the
363-
/// kernel
364-
void setCallsThisId(bool B);
365-
void setCallsThisItem(bool B);
366-
void setCallsThisNDItem(bool B);
367-
void setCallsThisGroup(bool B);
368-
369362
private:
370363
// Kernel actual parameter descriptor.
371364
struct KernelParamDesc {
@@ -383,15 +376,6 @@ class SYCLIntegrationHeader {
383376
KernelParamDesc() = default;
384377
};
385378

386-
// there are four free functions the kernel may call (this_id, this_item,
387-
// this_nd_item, this_group)
388-
struct KernelCallsSYCLFreeFunction {
389-
bool CallsThisId = false;
390-
bool CallsThisItem = false;
391-
bool CallsThisNDItem = false;
392-
bool CallsThisGroup = false;
393-
};
394-
395379
// Kernel invocation descriptor
396380
struct KernelDesc {
397381
/// sycl_kernel function associated with this kernel.
@@ -414,10 +398,6 @@ class SYCLIntegrationHeader {
414398
/// Descriptor of kernel actual parameters.
415399
SmallVector<KernelParamDesc, 8> Params;
416400

417-
// Whether kernel calls any of the SYCL free functions (this_item(),
418-
// this_id(), etc)
419-
KernelCallsSYCLFreeFunction FreeFunctionCalls;
420-
421401
// If we are in unnamed kernel/lambda mode AND this is one that the user
422402
// hasn't provided an explicit name for.
423403
bool IsUnnamedKernel;

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 0 additions & 117 deletions
Original file line numberDiff line numberDiff line change
@@ -114,12 +114,6 @@ class Util {
114114
/// \param Tmpl whether the class is template instantiation or simple record
115115
static bool isSyclType(QualType Ty, StringRef Name, bool Tmpl = false);
116116

117-
/// Checks whether given function is a standard SYCL API function with given
118-
/// name.
119-
/// \param FD the function being checked.
120-
/// \param Name the function name to be checked against.
121-
static bool isSyclFunction(const FunctionDecl *FD, StringRef Name);
122-
123117
/// Checks whether given clang type is a full specialization of the SYCL
124118
/// specialization constant class.
125119
static bool isSyclSpecConstantType(QualType Ty);
@@ -3143,60 +3137,6 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
31433137
return !SemaRef.getASTContext().hasSameType(FD->getType(), Ty);
31443138
}
31453139

3146-
// Sets a flag if the kernel is a parallel_for that calls the
3147-
// free function API "this_item".
3148-
void setThisItemIsCalled(FunctionDecl *KernelFunc) {
3149-
if (getKernelInvocationKind(KernelFunc) != InvokeParallelFor)
3150-
return;
3151-
3152-
// The call graph for this translation unit.
3153-
CallGraph SYCLCG;
3154-
SYCLCG.addToCallGraph(SemaRef.getASTContext().getTranslationUnitDecl());
3155-
using ChildParentPair =
3156-
std::pair<const FunctionDecl *, const FunctionDecl *>;
3157-
llvm::SmallPtrSet<const FunctionDecl *, 16> Visited;
3158-
llvm::SmallVector<ChildParentPair, 16> WorkList;
3159-
WorkList.push_back({KernelFunc, nullptr});
3160-
3161-
while (!WorkList.empty()) {
3162-
const FunctionDecl *FD = WorkList.back().first;
3163-
WorkList.pop_back();
3164-
if (!Visited.insert(FD).second)
3165-
continue; // We've already seen this Decl
3166-
3167-
// Check whether this call is to free functions (sycl::this_item(),
3168-
// this_id, etc.).
3169-
if (Util::isSyclFunction(FD, "this_id")) {
3170-
Header.setCallsThisId(true);
3171-
return;
3172-
}
3173-
if (Util::isSyclFunction(FD, "this_item")) {
3174-
Header.setCallsThisItem(true);
3175-
return;
3176-
}
3177-
if (Util::isSyclFunction(FD, "this_nd_item")) {
3178-
Header.setCallsThisNDItem(true);
3179-
return;
3180-
}
3181-
if (Util::isSyclFunction(FD, "this_group")) {
3182-
Header.setCallsThisGroup(true);
3183-
return;
3184-
}
3185-
3186-
CallGraphNode *N = SYCLCG.getNode(FD);
3187-
if (!N)
3188-
continue;
3189-
3190-
for (const CallGraphNode *CI : *N) {
3191-
if (auto *Callee = dyn_cast<FunctionDecl>(CI->getDecl())) {
3192-
Callee = Callee->getMostRecentDecl();
3193-
if (!Visited.count(Callee))
3194-
WorkList.push_back({Callee, FD});
3195-
}
3196-
}
3197-
}
3198-
}
3199-
32003140
public:
32013141
static constexpr const bool VisitInsideSimpleContainers = false;
32023142
SyclKernelIntHeaderCreator(Sema &S, SYCLIntegrationHeader &H,
@@ -3206,7 +3146,6 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
32063146
bool IsSIMDKernel = isESIMDKernelType(KernelObj);
32073147
Header.startKernel(KernelFunc, NameType, KernelObj->getLocation(),
32083148
IsSIMDKernel, IsSYCLUnnamedKernel(S, KernelFunc));
3209-
setThisItemIsCalled(KernelFunc);
32103149
}
32113150

32123151
bool handleSyclSpecialType(const CXXRecordDecl *RD,
@@ -4687,16 +4626,6 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
46874626
O << " __SYCL_DLL_LOCAL\n";
46884627
O << " static constexpr bool isESIMD() { return " << K.IsESIMDKernel
46894628
<< "; }\n";
4690-
O << " __SYCL_DLL_LOCAL\n";
4691-
O << " static constexpr bool callsThisItem() { return ";
4692-
O << K.FreeFunctionCalls.CallsThisItem << "; }\n";
4693-
O << " __SYCL_DLL_LOCAL\n";
4694-
O << " static constexpr bool callsAnyThisFreeFunction() { return ";
4695-
O << (K.FreeFunctionCalls.CallsThisId ||
4696-
K.FreeFunctionCalls.CallsThisItem ||
4697-
K.FreeFunctionCalls.CallsThisNDItem ||
4698-
K.FreeFunctionCalls.CallsThisGroup)
4699-
<< "; }\n";
47004629
O << "};\n";
47014630
CurStart += N;
47024631
}
@@ -4751,30 +4680,6 @@ void SYCLIntegrationHeader::addSpecConstant(StringRef IDName, QualType IDType) {
47514680
SpecConsts.emplace_back(std::make_pair(IDType, IDName.str()));
47524681
}
47534682

4754-
void SYCLIntegrationHeader::setCallsThisId(bool B) {
4755-
KernelDesc *K = getCurKernelDesc();
4756-
assert(K && "no kernel");
4757-
K->FreeFunctionCalls.CallsThisId = B;
4758-
}
4759-
4760-
void SYCLIntegrationHeader::setCallsThisItem(bool B) {
4761-
KernelDesc *K = getCurKernelDesc();
4762-
assert(K && "no kernel");
4763-
K->FreeFunctionCalls.CallsThisItem = B;
4764-
}
4765-
4766-
void SYCLIntegrationHeader::setCallsThisNDItem(bool B) {
4767-
KernelDesc *K = getCurKernelDesc();
4768-
assert(K && "no kernel");
4769-
K->FreeFunctionCalls.CallsThisNDItem = B;
4770-
}
4771-
4772-
void SYCLIntegrationHeader::setCallsThisGroup(bool B) {
4773-
KernelDesc *K = getCurKernelDesc();
4774-
assert(K && "no kernel");
4775-
K->FreeFunctionCalls.CallsThisGroup = B;
4776-
}
4777-
47784683
SYCLIntegrationHeader::SYCLIntegrationHeader(Sema &S) : S(S) {}
47794684

47804685
void SYCLIntegrationFooter::addVarDecl(const VarDecl *VD) {
@@ -5099,28 +5004,6 @@ bool Util::isSyclType(QualType Ty, StringRef Name, bool Tmpl) {
50995004
return matchQualifiedTypeName(Ty, Scopes);
51005005
}
51015006

5102-
bool Util::isSyclFunction(const FunctionDecl *FD, StringRef Name) {
5103-
if (!FD->isFunctionOrMethod() || !FD->getIdentifier() ||
5104-
FD->getName().empty() || Name != FD->getName())
5105-
return false;
5106-
5107-
const DeclContext *DC = FD->getDeclContext();
5108-
if (DC->isTranslationUnit())
5109-
return false;
5110-
5111-
std::array<DeclContextDesc, 2> ScopesSycl = {
5112-
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
5113-
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl")};
5114-
std::array<DeclContextDesc, 5> ScopesOneapiExp = {
5115-
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
5116-
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"),
5117-
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "ext"),
5118-
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "oneapi"),
5119-
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "experimental")};
5120-
5121-
return matchContext(DC, ScopesSycl) || matchContext(DC, ScopesOneapiExp);
5122-
}
5123-
51245007
bool Util::isAccessorPropertyListType(QualType Ty) {
51255008
std::array<DeclContextDesc, 5> Scopes = {
51265009
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),

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

Lines changed: 0 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -143,22 +143,9 @@ namespace ext {
143143
namespace oneapi {
144144
template <typename... properties>
145145
class accessor_property_list {};
146-
namespace experimental {
147-
template <int Dims> item<Dims>
148-
this_item() { return item<Dims>{}; }
149-
150-
template <int Dims> id<Dims>
151-
this_id() { return id<Dims>{}; }
152-
} // namespace experimental
153146
} // namespace oneapi
154147
} // namespace ext
155148

156-
template <int Dims> item<Dims>
157-
this_item() { return item<Dims>{}; }
158-
159-
template <int Dims> id<Dims>
160-
this_id() { return id<Dims>{}; }
161-
162149
template <int dim>
163150
struct range {
164151
template <typename... T>

0 commit comments

Comments
 (0)