Skip to content

[SYCL][ESIMD][EMU] PI_API debug for esimd_emulator plug-in #5606

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

[SYCL][ESIMD][EMU] PI_API debug for esimd_emulator plug-in #5606

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
29 commits
Select commit Hold shift + click to select a range
d47d035
[SYCL][ESIMD][EMU] PI_API debug for esimd_emulator plug-in
dongkyunahn-intel Feb 17, 2022
34f4868
HostSurfaceIndex is generated for surface mapping
dongkyunahn-intel Feb 18, 2022
527762a
Minor changes
dongkyunahn-intel Feb 22, 2022
e8be377
Replacing sycl::detail::SpinLock with std::mutex
dongkyunahn-intel Feb 23, 2022
0206b19
ESIMD_EMU generates and manages its own surface index
dongkyunahn-intel Feb 23, 2022
3d163a4
Updating ESIMD Emulator Device interface
dongkyunahn-intel Feb 24, 2022
582d922
Missing renaming
dongkyunahn-intel Feb 24, 2022
31219e9
Share-malloc size adjusting to power-of-2 for piextUSMSharedAlloc()
dongkyunahn-intel Feb 25, 2022
0c2c01d
Using common 'getNextPowerOfTwo' from sycl::detail instead of local
dongkyunahn-intel Feb 28, 2022
5ec625e
Addressing Vlad's comments
dongkyunahn-intel Mar 1, 2022
afe99a8
Merge branch 'sycl' of https://github.com/intel/llvm into esimd_emu_p…
dongkyunahn-intel Mar 4, 2022
4ed7177
Preparing pulling origin/sycl branch
dongkyunahn-intel Mar 9, 2022
227358e
Merge branch 'sycl' of https://github.com/intel/llvm into esimd_emu_p…
dongkyunahn-intel Mar 9, 2022
6e3d768
Re-applying changes in header files to relocated files
dongkyunahn-intel Mar 9, 2022
c8ee300
Merge branch 'sycl' of https://github.com/intel/llvm into esimd_emu_p…
dongkyunahn-intel Mar 15, 2022
85edf0e
For Host-memory buffer/image, Use CM's UserProvided surface creation
dongkyunahn-intel Mar 15, 2022
5307e74
Function pointer renaming aligned with naming convention
dongkyunahn-intel Mar 15, 2022
72c05c6
clang-format fix
dongkyunahn-intel Mar 15, 2022
9a96610
Applying structured cm_buf type for regular and user-provided memory
dongkyunahn-intel Mar 16, 2022
0ea135c
structure cm_image type
dongkyunahn-intel Mar 16, 2022
10c2e19
Miscellaneous changes
dongkyunahn-intel Mar 17, 2022
a9fc46d
Addressing comments
dongkyunahn-intel Mar 17, 2022
829124e
Addressing comments (2)
dongkyunahn-intel Mar 22, 2022
78c5132
Addressing comments (3)
dongkyunahn-intel Mar 24, 2022
0a3da02
Merge branch 'sycl' of https://github.com/intel/llvm into esimd_emu_p…
dongkyunahn-intel Mar 28, 2022
a233331
Adding comments regarding CM Runtime function call
dongkyunahn-intel Mar 28, 2022
7e4aa48
Addressing Leonid.S's comments
dongkyunahn-intel Mar 28, 2022
6ad6d2a
Reverting virtual destructor change
dongkyunahn-intel Mar 28, 2022
0bb5d63
Recovering comment removed by mistake
dongkyunahn-intel Mar 28, 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
50 changes: 21 additions & 29 deletions sycl/include/sycl/ext/intel/esimd/detail/memory_intrin.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -225,10 +225,9 @@ __esimd_oword_ld_unaligned(SurfIndAliasTy surf_ind, uint32_t offset)
uint32_t width;
std::mutex *mutexLock;

I->sycl_get_cm_buffer_params_index_ptr(surf_ind, &readBase, &width,
&mutexLock);
I->sycl_get_cm_buffer_params_ptr(surf_ind, &readBase, &width, &mutexLock);

std::unique_lock<std::mutex> lock(*mutexLock);
std::lock_guard<std::mutex> lock(*mutexLock);

for (int idx = 0; idx < N; idx++) {
if (offset >= width) {
Expand Down Expand Up @@ -270,10 +269,9 @@ __ESIMD_INTRIN void __esimd_oword_st(SurfIndAliasTy surf_ind, uint32_t offset,
uint32_t width;
std::mutex *mutexLock;

I->sycl_get_cm_buffer_params_index_ptr(surf_ind, &writeBase, &width,
&mutexLock);
I->sycl_get_cm_buffer_params_ptr(surf_ind, &writeBase, &width, &mutexLock);

std::unique_lock<std::mutex> lock(*mutexLock);
std::lock_guard<std::mutex> lock(*mutexLock);

for (int idx = 0; idx < N; idx++) {
if (offset < width) {
Expand Down Expand Up @@ -458,11 +456,10 @@ __esimd_scatter_scaled(__ESIMD_DNS::simd_mask_storage_t<N> pred,
uint32_t width;
std::mutex *mutexLock;

I->sycl_get_cm_buffer_params_index_ptr(surf_ind, &writeBase, &width,
&mutexLock);
I->sycl_get_cm_buffer_params_ptr(surf_ind, &writeBase, &width, &mutexLock);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Why passing by pointer and not reference?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Pointer is more clear & explicit on the caller side.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Plus sycl_get_cm_buffer_params_ptr is a part of C interface, not C++.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

If it's C, then ok.

writeBase += global_offset;

std::unique_lock<std::mutex> lock(*mutexLock);
std::lock_guard<std::mutex> lock(*mutexLock);

for (int idx = 0; idx < N; idx++) {
if (pred[idx]) {
Expand Down Expand Up @@ -594,11 +591,10 @@ __esimd_gather_scaled(__ESIMD_DNS::simd_mask_storage_t<N> pred,
uint32_t width;
std::mutex *mutexLock;

I->sycl_get_cm_buffer_params_index_ptr(surf_ind, &readBase, &width,
&mutexLock);
I->sycl_get_cm_buffer_params_ptr(surf_ind, &readBase, &width, &mutexLock);
readBase += global_offset;

std::unique_lock<std::mutex> lock(*mutexLock);
std::lock_guard<std::mutex> lock(*mutexLock);

for (int idx = 0; idx < N; idx++) {
if (pred[idx]) {
Expand Down Expand Up @@ -672,11 +668,10 @@ __esimd_gather_masked_scaled2(SurfIndAliasTy surf_ind, uint32_t global_offset,
uint32_t width;
std::mutex *mutexLock;

I->sycl_get_cm_buffer_params_index_ptr(surf_ind, &readBase, &width,
&mutexLock);
I->sycl_get_cm_buffer_params_ptr(surf_ind, &readBase, &width, &mutexLock);

readBase += global_offset;
std::unique_lock<std::mutex> lock(*mutexLock);
std::lock_guard<std::mutex> lock(*mutexLock);
for (int idx = 0; idx < N; idx++) {
if (pred[idx]) {
RestoredTy *addr =
Expand Down Expand Up @@ -727,10 +722,9 @@ __esimd_oword_ld(SurfIndAliasTy surf_ind, uint32_t addr)
uint32_t width;
std::mutex *mutexLock;

I->sycl_get_cm_buffer_params_index_ptr(surf_ind, &readBase, &width,
&mutexLock);
I->sycl_get_cm_buffer_params_ptr(surf_ind, &readBase, &width, &mutexLock);

std::unique_lock<std::mutex> lock(*mutexLock);
std::lock_guard<std::mutex> lock(*mutexLock);

for (int idx = 0; idx < N; idx++) {
if (addr >= width) {
Expand Down Expand Up @@ -768,9 +762,8 @@ __ESIMD_INTRIN
} else {
uint32_t width;
std::mutex *mutexLock;
I->sycl_get_cm_buffer_params_index_ptr(surf_ind, &ReadBase, &width,
&mutexLock);
std::unique_lock<std::mutex> lock(*mutexLock);
I->sycl_get_cm_buffer_params_ptr(surf_ind, &ReadBase, &width, &mutexLock);
std::lock_guard<std::mutex> lock(*mutexLock);
}

ReadBase += global_offset;
Expand Down Expand Up @@ -812,9 +805,8 @@ __ESIMD_INTRIN void __esimd_scatter4_scaled(
} else {
uint32_t width;
std::mutex *mutexLock;
I->sycl_get_cm_buffer_params_index_ptr(surf_ind, &WriteBase, &width,
&mutexLock);
std::unique_lock<std::mutex> lock(*mutexLock);
I->sycl_get_cm_buffer_params_ptr(surf_ind, &WriteBase, &width, &mutexLock);
std::lock_guard<std::mutex> lock(*mutexLock);
}

WriteBase += global_offset;
Expand Down Expand Up @@ -931,10 +923,10 @@ __esimd_media_ld(TACC handle, unsigned x, unsigned y)
assert((handle != __ESIMD_NS::detail::SLM_BTI) &&
"__esimd_media_ld cannot access SLM");

sycl::detail::getESIMDDeviceInterface()->sycl_get_cm_image_params_index_ptr(
sycl::detail::getESIMDDeviceInterface()->sycl_get_cm_image_params_ptr(
handle, &readBase, &imgWidth, &imgHeight, &bpp, &mutexLock);

std::unique_lock<std::mutex> lock(*mutexLock);
std::lock_guard<std::mutex> lock(*mutexLock);

int x_pos_a, y_pos_a, offset, index;

Expand Down Expand Up @@ -1061,8 +1053,8 @@ __ESIMD_INTRIN void __esimd_media_st(TACC handle, unsigned x, unsigned y,
assert((handle != __ESIMD_NS::detail::SLM_BTI) &&
"__esimd_media_ld cannot access SLM");

I->sycl_get_cm_image_params_index_ptr(handle, &writeBase, &imgWidth,
&imgHeight, &bpp, &mutexLock);
I->sycl_get_cm_image_params_ptr(handle, &writeBase, &imgWidth, &imgHeight,
&bpp, &mutexLock);

int x_pos_a, y_pos_a, offset;

Expand All @@ -1072,7 +1064,7 @@ __ESIMD_INTRIN void __esimd_media_st(TACC handle, unsigned x, unsigned y,
// TODO : Remove intermediate 'out' matrix
std::vector<std::vector<Ty>> out(M, std::vector<Ty>(N));

std::unique_lock<std::mutex> lock(*mutexLock);
std::lock_guard<std::mutex> lock(*mutexLock);

for (int i = 0, k = 0; i < M; i++) {
for (int j = 0; j < N; j++) {
Expand Down
13 changes: 4 additions & 9 deletions sycl/include/sycl/ext/intel/esimd/emu/detail/esimd_emulator_functions_v1.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -46,15 +46,10 @@ void (*cm_fence_ptr)(void);
char *(*sycl_get_surface_base_addr_ptr)(int);
char *(*__cm_emu_get_slm_ptr)(void);
void (*cm_slm_init_ptr)(size_t);
void (*sycl_get_cm_buffer_params_ptr)(void *, char **, uint32_t *,
std::mutex **);
void (*sycl_get_cm_image_params_ptr)(void *, char **, uint32_t *, uint32_t *,
uint32_t *, std::mutex **);

unsigned int (*sycl_get_cm_surface_index_ptr)(void *);
void (*sycl_get_cm_buffer_params_index_ptr)(unsigned int, char **, uint32_t *,
std::mutex **);
void (*sycl_get_cm_image_params_index_ptr)(unsigned int, char **, uint32_t *,
uint32_t *, uint32_t *,
std::mutex **);
void (*sycl_get_cm_buffer_params_ptr)(unsigned int, char **, uint32_t *,
std::mutex **);
void (*sycl_get_cm_image_params_ptr)(unsigned int, char **, uint32_t *,
uint32_t *, uint32_t *, std::mutex **);
/// @endcond ESIMD_EMU
Loading