Skip to content

[SYCL][ESIMD] Support get_multi_ptr in ESIMD context #10638

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
Aug 2, 2023
Merged

[SYCL][ESIMD] Support get_multi_ptr in ESIMD context #10638

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
1270ffa
[SYCL][ESIMD] Support get_multi_ptr in ESIMD context
turinevgeny Aug 1, 2023
2f0fc2d
Merge branch 'sycl' into SupportGetMultiPtrESIMD
turinevgeny Aug 2, 2023
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
2 changes: 2 additions & 0 deletions llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -39,6 +39,7 @@ static const char *LegalSYCLFunctions[] = {
"^sycl::_V1::local_accessor<.+>::local_accessor",
"^sycl::_V1::local_accessor<.+>::__init_esimd",
"^sycl::_V1::local_accessor<.+>::get_pointer",
"^sycl::_V1::local_accessor<.+>::get_multi_ptr",
"^sycl::_V1::local_accessor_base<.+>::local_accessor_base",
"^sycl::_V1::local_accessor_base<.+>::__init_esimd",
"^sycl::_V1::local_accessor_base<.+>::getQualifiedPtr",
Expand Down Expand Up @@ -67,6 +68,7 @@ static const char *LegalSYCLFunctions[] = {

static const char *LegalSYCLFunctionsInStatelessMode[] = {
"^sycl::_V1::accessor<.+>::get_pointer.+",
"^sycl::_V1::accessor<.+>::get_multi_ptr.+",
"^sycl::_V1::accessor<.+>::getPointerAdjusted",
"^sycl::_V1::accessor<.+>::getTotalOffset",
"^sycl::_V1::accessor<.+>::getLinearIndex",
Expand Down
4 changes: 3 additions & 1 deletion sycl/test-e2e/InvokeSimd/Regression/nbarrier_exec_in_order.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -53,7 +53,9 @@ ESIMD_INLINE void ESIMD_CALLEE_nbarrier(local_accessor<int, 1> local_acc,

esimd::simd<int, VL> val(local_id);

unsigned int slm_base = static_cast<uint32_t>(reinterpret_cast<std::uintptr_t>(local_acc.get_pointer()));
unsigned int slm_base =
static_cast<uint32_t>(reinterpret_cast<std::uintptr_t>(
local_acc.get_multi_ptr<access::decorated::no>().get_raw()));

/* Each thread operates on a region of memory (global or slm) that overlaps
* with that of the previous and next threads.
Expand Down
5 changes: 3 additions & 2 deletions sycl/test-e2e/InvokeSimd/Regression/nbarrier_loop.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -60,8 +60,9 @@ ESIMD_INLINE void ESIMD_CALLEE_nbarrier(local_accessor<int, 1> local_acc,
unsigned int prods[2] = {2, 1}; // number of producers

// Producer writes to SLM, consumer reads what producer wrote.
unsigned int slm_base = static_cast<uint32_t>(
reinterpret_cast<std::uintptr_t>(local_acc.get_pointer()));
unsigned int slm_base =
static_cast<uint32_t>(reinterpret_cast<std::uintptr_t>(
local_acc.get_multi_ptr<access::decorated::no>().get_raw()));

esimd::barrier();

Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/InvokeSimd/Regression/nbarrier_multiple_wg.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -64,8 +64,8 @@ ESIMD_INLINE void ESIMD_CALLEE_nbarrier(local_accessor<int, 1> local_acc,
unsigned int global_off = VL * (is_producer ? global_id : (global_id - 1));

unsigned int slm_base =
static_cast<uint32_t>(
reinterpret_cast<std::uintptr_t>(local_acc.get_pointer()));
static_cast<uint32_t>(reinterpret_cast<std::uintptr_t>(
local_acc.get_multi_ptr<access::decorated::no>().get_raw()));
unsigned int slm_off = slm_base + global_off * sizeof(int);

esimd::barrier();
Expand Down
11 changes: 9 additions & 2 deletions sycl/test/esimd/esimd_verify.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,6 +10,7 @@ using namespace sycl;
using namespace sycl::ext::intel::esimd;

// CHECK-NEGATIVE-DAG: error: function 'sycl::_V1::multi_ptr<{{.+}}> sycl::_V1::accessor<{{.+}}>::get_pointer<{{.+}}>() const' is not supported in ESIMD context
// CHECK-NEGATIVE-DAG: error: function 'std::conditional<true, sycl::_V1::multi_ptr{{.+}}>::type sycl::_V1::accessor<{{.+}}>::get_multi_ptr<{{.+}}>() const' is not supported in ESIMD context
// CHECK-NEGATIVE-DAG: error: function '{{.+}} sycl::_V1::accessor<{{.+}}>::operator[]<{{.+}}>({{.+}}) const' is not supported in ESIMD context
// CHECK-NEGATIVE-DAG: error: function '{{.+}}combine(int const&)' is not supported in ESIMD context

Expand All @@ -21,13 +22,19 @@ test0(accessor<int, 1, access::mode::read_write, access::target::device> &acc)
return acc.get_pointer();
}

SYCL_EXTERNAL void
SYCL_EXTERNAL auto
test1(accessor<int, 1, access::mode::read_write, access::target::device> &acc)
SYCL_ESIMD_FUNCTION {
return acc.get_multi_ptr<access::decorated::no>();
}

SYCL_EXTERNAL void
test2(accessor<int, 1, access::mode::read_write, access::target::device> &acc)
SYCL_ESIMD_FUNCTION {
acc[0] = 0;
}

void test2(sycl::handler &cgh, int *buf) {
void test3(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)
Expand Down