Skip to content

[SYCL][ESIMD] Change esimd-verifier logic for detecting valid SYCL calls #5914

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 2 commits into from
Mar 31, 2022
Merged

[SYCL][ESIMD] Change esimd-verifier logic for detecting valid SYCL calls #5914

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
0119913
[SYCL][ESIMD] Change esimd-verifier logic for detecting valid SYCL calls
sndmitriev Mar 29, 2022
36cd085
Addressed comments from code review.
sndmitriev Mar 30, 2022
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
111 changes: 93 additions & 18 deletions llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,6 +13,7 @@

#include "llvm/SYCLLowerIR/ESIMD/ESIMDVerifier.h"
#include "llvm/Demangle/Demangle.h"
#include "llvm/Demangle/ItaniumDemangle.h"
#include "llvm/IR/InstIterator.h"
#include "llvm/IR/Instructions.h"
#include "llvm/IR/Module.h"
Expand All @@ -22,17 +23,64 @@
#include "llvm/Support/Regex.h"

using namespace llvm;
namespace id = itanium_demangle;

#define DEBUG_TYPE "esimd-verifier"

// A list of unsupported functions in ESIMD context.
static const char *IllegalFunctions[] = {
"^cl::sycl::multi_ptr<.+> cl::sycl::accessor<.+>::get_pointer<.+>\\(\\) "
"const",
" cl::sycl::accessor<.+>::operator\\[\\]<.+>\\(.+\\) const"};
// A list of SYCL functions (regexps) allowed for use in ESIMD context.
static const char *LegalSYCLFunctions[] = {
"^cl::sycl::accessor<.+>::accessor",
"^cl::sycl::accessor<.+>::~accessor",
"^cl::sycl::accessor<.+>::getNativeImageObj",
"^cl::sycl::accessor<.+>::__init_esimd",
"^cl::sycl::id<.+>::.+",
"^cl::sycl::item<.+>::.+",
"^cl::sycl::nd_item<.+>::.+",
"^cl::sycl::group<.+>::.+",
"^cl::sycl::sub_group<.+>::.+",
"^cl::sycl::range<.+>::.+",
"^cl::sycl::kernel_handler::.+",
"^cl::sycl::cos<.+>",
"^cl::sycl::sin<.+>",
"^cl::sycl::log<.+>",
"^cl::sycl::exp<.+>",
"^cl::sycl::operator.+<.+>",
"^cl::sycl::ext::oneapi::sub_group::.+",
"^cl::sycl::ext::oneapi::experimental::spec_constant<.+>::.+",
"^cl::sycl::ext::oneapi::experimental::this_sub_group"};

namespace {

// Simplest possible implementation of an allocator for the Itanium demangler
class SimpleAllocator {
protected:
SmallVector<void *, 128> Ptrs;

public:
void reset() {
for (void *Ptr : Ptrs) {
// Destructors are not called, but that is OK for the
// itanium_demangle::Node subclasses
std::free(Ptr);
}
Ptrs.resize(0);
}

template <typename T, typename... Args> T *makeNode(Args &&...args) {
void *Ptr = std::calloc(1, sizeof(T));
Ptrs.push_back(Ptr);
return new (Ptr) T(std::forward<Args>(args)...);
}

void *allocateNodeArray(size_t sz) {
void *Ptr = std::calloc(sz, sizeof(id::Node *));
Ptrs.push_back(Ptr);
return Ptr;
}

~SimpleAllocator() { reset(); }
};

class ESIMDVerifierImpl {
const Module &M;

Expand Down Expand Up @@ -63,22 +111,49 @@ class ESIMDVerifierImpl {
if (!Callee)
continue;

// Demangle called function name and check if it matches any illegal
// function name. Report an error if there is a match.
std::string DemangledName = demangle(Callee->getName().str());
for (const char *Name : IllegalFunctions) {
Regex NameRE(Name);
assert(NameRE.isValid() && "invalid function name regex");
if (NameRE.match(DemangledName)) {
std::string ErrorMsg = std::string("function '") + DemangledName +
"' is not supported in ESIMD context";
F->getContext().emitError(&I, ErrorMsg);
}
}

// Add callee to the list to be analyzed if it is not a declaration.
if (!Callee->isDeclaration())
Add2Worklist(Callee);

// Demangle called function name and check if it is legal to use this
// function in ESIMD context.
StringRef MangledName = Callee->getName();
id::ManglingParser<SimpleAllocator> Parser(MangledName.begin(),
MangledName.end());
id::Node *AST = Parser.parse();
if (!AST || AST->getKind() != id::Node::KFunctionEncoding)
continue;

auto *FE = static_cast<id::FunctionEncoding *>(AST);
const id::Node *NameNode = FE->getName();
if (!NameNode) // Can it be null?
continue;

id::OutputBuffer NameBuf;
NameNode->print(NameBuf);
StringRef Name(NameBuf.getBuffer(), NameBuf.getCurrentPosition());

// We are interested in functions defined in SYCL namespace, but
// outside of ESIMD namespaces.
if (!Name.startswith("cl::sycl::") ||
Name.startswith("cl::sycl::detail::") ||
Name.startswith("cl::sycl::ext::intel::esimd::") ||
Name.startswith("cl::sycl::ext::intel::experimental::esimd::"))
continue;

// Check if function name matches any allowed SYCL function name.
if (any_of(LegalSYCLFunctions, [Name](const char *LegalName) {
Regex LegalNameRE(LegalName);
assert(LegalNameRE.isValid() && "invalid function name regex");
return LegalNameRE.match(Name);
}))
continue;

// If not, report an error.
std::string ErrorMsg = std::string("function '") +
demangle(MangledName.str()) +
"' is not supported in ESIMD context";
F->getContext().emitError(&I, ErrorMsg);
}
}
}
Expand Down
9 changes: 9 additions & 0 deletions sycl/test/esimd/esimd_verify.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,13 +3,15 @@
// RUN: not %clangxx -fsycl -fsycl-device-only -flegacy-pass-manager -O0 -S %s -o /dev/null 2>&1 | FileCheck %s
// RUN: not %clangxx -fsycl -fsycl-device-only -fno-legacy-pass-manager -O0 -S %s -o /dev/null 2>&1 | FileCheck %s

#include <CL/sycl.hpp>
#include <sycl/ext/intel/esimd.hpp>

using namespace cl::sycl;
using namespace sycl::ext::intel::esimd;

// CHECK-DAG: error: function 'cl::sycl::multi_ptr<{{.+}}> cl::sycl::accessor<{{.+}}>::get_pointer<{{.+}}>() const' is not supported in ESIMD context
// CHECK-DAG: error: function '{{.+}} cl::sycl::accessor<{{.+}}>::operator[]<{{.+}}>({{.+}}) const' is not supported in ESIMD context
// CHECK-DAG: error: function 'cl::sycl::ext::oneapi::detail::reducer<int, std::plus<int>, void>::combine(int const&)' is not supported in ESIMD context

SYCL_EXTERNAL auto
test(accessor<int, 1, access::mode::read_write, access::target::device> &acc)
Expand All @@ -22,3 +24,10 @@ test1(accessor<int, 1, access::mode::read_write, access::target::device> &acc)
SYCL_ESIMD_FUNCTION {
acc[0] = 0;
}

void test2(sycl::handler &cgh, int *buf) {
auto reduction = sycl::reduction(buf, sycl::plus<int>());
cgh.parallel_for<class Test2>(sycl::range<1>(1), reduction,
[=](sycl::id<1>, auto &reducer)
SYCL_ESIMD_KERNEL { reducer.combine(15); });
}