Skip to content

[SYCL] Remove free function queries call detection from FE #5178

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

Merged
merged 4 commits into from
Dec 23, 2021
Merged

[SYCL] Remove free function queries call detection from FE #5178

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
0de9b86
[SYCL] Remove free function queries call detection from FE
alexbatashev Dec 19, 2021
82ae241
remove more code
alexbatashev Dec 22, 2021
22b39a2
more removals
alexbatashev Dec 22, 2021
fa65690
Merge remote-tracking branch 'upstream/sycl' into remove_this_func_fe
alexbatashev Dec 23, 2021
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
20 changes: 0 additions & 20 deletions clang/include/clang/Sema/Sema.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -359,13 +359,6 @@ class SYCLIntegrationHeader {
Itr->updateKernelNames(Name, StableName);
}

/// Note which free functions (this_id, this_item, etc) are called within the
/// kernel
void setCallsThisId(bool B);
void setCallsThisItem(bool B);
void setCallsThisNDItem(bool B);
void setCallsThisGroup(bool B);

private:
// Kernel actual parameter descriptor.
struct KernelParamDesc {
Expand All @@ -383,15 +376,6 @@ class SYCLIntegrationHeader {
KernelParamDesc() = default;
};

// there are four free functions the kernel may call (this_id, this_item,
// this_nd_item, this_group)
struct KernelCallsSYCLFreeFunction {
bool CallsThisId = false;
bool CallsThisItem = false;
bool CallsThisNDItem = false;
bool CallsThisGroup = false;
};

// Kernel invocation descriptor
struct KernelDesc {
/// sycl_kernel function associated with this kernel.
Expand All @@ -414,10 +398,6 @@ class SYCLIntegrationHeader {
/// Descriptor of kernel actual parameters.
SmallVector<KernelParamDesc, 8> Params;

// Whether kernel calls any of the SYCL free functions (this_item(),
// this_id(), etc)
KernelCallsSYCLFreeFunction FreeFunctionCalls;

// If we are in unnamed kernel/lambda mode AND this is one that the user
// hasn't provided an explicit name for.
bool IsUnnamedKernel;
Expand Down
117 changes: 0 additions & 117 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -114,12 +114,6 @@ class Util {
/// \param Tmpl whether the class is template instantiation or simple record
static bool isSyclType(QualType Ty, StringRef Name, bool Tmpl = false);

/// Checks whether given function is a standard SYCL API function with given
/// name.
/// \param FD the function being checked.
/// \param Name the function name to be checked against.
static bool isSyclFunction(const FunctionDecl *FD, StringRef Name);

/// Checks whether given clang type is a full specialization of the SYCL
/// specialization constant class.
static bool isSyclSpecConstantType(QualType Ty);
Expand Down Expand Up @@ -3143,60 +3137,6 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
return !SemaRef.getASTContext().hasSameType(FD->getType(), Ty);
}

// Sets a flag if the kernel is a parallel_for that calls the
// free function API "this_item".
void setThisItemIsCalled(FunctionDecl *KernelFunc) {
if (getKernelInvocationKind(KernelFunc) != InvokeParallelFor)
return;

// The call graph for this translation unit.
CallGraph SYCLCG;
SYCLCG.addToCallGraph(SemaRef.getASTContext().getTranslationUnitDecl());
using ChildParentPair =
std::pair<const FunctionDecl *, const FunctionDecl *>;
llvm::SmallPtrSet<const FunctionDecl *, 16> Visited;
llvm::SmallVector<ChildParentPair, 16> WorkList;
WorkList.push_back({KernelFunc, nullptr});

while (!WorkList.empty()) {
const FunctionDecl *FD = WorkList.back().first;
WorkList.pop_back();
if (!Visited.insert(FD).second)
continue; // We've already seen this Decl

// Check whether this call is to free functions (sycl::this_item(),
// this_id, etc.).
if (Util::isSyclFunction(FD, "this_id")) {
Header.setCallsThisId(true);
return;
}
if (Util::isSyclFunction(FD, "this_item")) {
Header.setCallsThisItem(true);
return;
}
if (Util::isSyclFunction(FD, "this_nd_item")) {
Header.setCallsThisNDItem(true);
return;
}
if (Util::isSyclFunction(FD, "this_group")) {
Header.setCallsThisGroup(true);
return;
}

CallGraphNode *N = SYCLCG.getNode(FD);
if (!N)
continue;

for (const CallGraphNode *CI : *N) {
if (auto *Callee = dyn_cast<FunctionDecl>(CI->getDecl())) {
Callee = Callee->getMostRecentDecl();
if (!Visited.count(Callee))
WorkList.push_back({Callee, FD});
}
}
}
}

public:
static constexpr const bool VisitInsideSimpleContainers = false;
SyclKernelIntHeaderCreator(Sema &S, SYCLIntegrationHeader &H,
Expand All @@ -3206,7 +3146,6 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
bool IsSIMDKernel = isESIMDKernelType(KernelObj);
Header.startKernel(KernelFunc, NameType, KernelObj->getLocation(),
IsSIMDKernel, IsSYCLUnnamedKernel(S, KernelFunc));
setThisItemIsCalled(KernelFunc);
}

bool handleSyclSpecialType(const CXXRecordDecl *RD,
Expand Down Expand Up @@ -4687,16 +4626,6 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
O << " __SYCL_DLL_LOCAL\n";
O << " static constexpr bool isESIMD() { return " << K.IsESIMDKernel
<< "; }\n";
O << " __SYCL_DLL_LOCAL\n";
O << " static constexpr bool callsThisItem() { return ";
O << K.FreeFunctionCalls.CallsThisItem << "; }\n";
O << " __SYCL_DLL_LOCAL\n";
O << " static constexpr bool callsAnyThisFreeFunction() { return ";
O << (K.FreeFunctionCalls.CallsThisId ||
K.FreeFunctionCalls.CallsThisItem ||
K.FreeFunctionCalls.CallsThisNDItem ||
K.FreeFunctionCalls.CallsThisGroup)
<< "; }\n";
O << "};\n";
CurStart += N;
}
Expand Down Expand Up @@ -4751,30 +4680,6 @@ void SYCLIntegrationHeader::addSpecConstant(StringRef IDName, QualType IDType) {
SpecConsts.emplace_back(std::make_pair(IDType, IDName.str()));
}

void SYCLIntegrationHeader::setCallsThisId(bool B) {
KernelDesc *K = getCurKernelDesc();
assert(K && "no kernel");
K->FreeFunctionCalls.CallsThisId = B;
}

void SYCLIntegrationHeader::setCallsThisItem(bool B) {
KernelDesc *K = getCurKernelDesc();
assert(K && "no kernel");
K->FreeFunctionCalls.CallsThisItem = B;
}

void SYCLIntegrationHeader::setCallsThisNDItem(bool B) {
KernelDesc *K = getCurKernelDesc();
assert(K && "no kernel");
K->FreeFunctionCalls.CallsThisNDItem = B;
}

void SYCLIntegrationHeader::setCallsThisGroup(bool B) {
KernelDesc *K = getCurKernelDesc();
assert(K && "no kernel");
K->FreeFunctionCalls.CallsThisGroup = B;
}

SYCLIntegrationHeader::SYCLIntegrationHeader(Sema &S) : S(S) {}

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

bool Util::isSyclFunction(const FunctionDecl *FD, StringRef Name) {
if (!FD->isFunctionOrMethod() || !FD->getIdentifier() ||
FD->getName().empty() || Name != FD->getName())
return false;

const DeclContext *DC = FD->getDeclContext();
if (DC->isTranslationUnit())
return false;

std::array<DeclContextDesc, 2> ScopesSycl = {
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl")};
std::array<DeclContextDesc, 5> ScopesOneapiExp = {
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "ext"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "oneapi"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "experimental")};

return matchContext(DC, ScopesSycl) || matchContext(DC, ScopesOneapiExp);
}

bool Util::isAccessorPropertyListType(QualType Ty) {
std::array<DeclContextDesc, 5> Scopes = {
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
Expand Down
13 changes: 0 additions & 13 deletions clang/test/CodeGenSYCL/Inputs/sycl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -143,22 +143,9 @@ namespace ext {
namespace oneapi {
template <typename... properties>
class accessor_property_list {};
namespace experimental {
template <int Dims> item<Dims>
this_item() { return item<Dims>{}; }

template <int Dims> id<Dims>
this_id() { return id<Dims>{}; }
} // namespace experimental
} // namespace oneapi
} // namespace ext

template <int Dims> item<Dims>
this_item() { return item<Dims>{}; }

template <int Dims> id<Dims>
this_id() { return id<Dims>{}; }

template <int dim>
struct range {
template <typename... T>
Expand Down
Loading