Skip to content

Commit bc0579a

Browse files
dongkyunahn-inteldongkyunahn-intel
authored
[SYCL][ESIMD][EMU] PI_API debug for esimd_emulator plug-in (#5606)
* [SYCL][ESIMD][EMU] Fix bugs in and implement some PI API implementation in esimd_emulator plug-in - Enabling piEnqueueMemBufferMap/Unmap - Serializing acces to Addr2CmBufferSVM - Replacing sycl::detail::SpinLock with std::mutex - Unused functions are removed - sycl_get_cm_buffer/image_params - Interface functions for getting surface info are renamed to not have 'cm' as they are used for surfaces generated by both CM and Host - As there is no legacy from previous productization, interface functions can be revised for now while keeping interface version as v1 - Share-malloc size adjusting to power-of-2 for piextUSMSharedAlloc() - fixes memory corruption error from '$TEST_SUITE/ESIMD/SYCL/api/simd_any_all.cpp' - Changes in atomic operations - Removing unused surface index generator - Applying structured cm_buf type for regular and user-provided memory
1 parent a9be4a2 commit bc0579a

File tree

4 files changed

+393
-198
lines changed

4 files changed

+393
-198
lines changed

‎sycl/include/sycl/ext/intel/esimd/detail/memory_intrin.hpp

Lines changed: 21 additions & 29 deletions
Original file line numberDiff line numberDiff line change
@@ -225,10 +225,9 @@ __esimd_oword_ld_unaligned(SurfIndAliasTy surf_ind, uint32_t offset)
225225
uint32_t width;
226226
std::mutex *mutexLock;
227227

228-
I->sycl_get_cm_buffer_params_index_ptr(surf_ind, &readBase, &width,
229-
&mutexLock);
228+
I->sycl_get_cm_buffer_params_ptr(surf_ind, &readBase, &width, &mutexLock);
230229

231-
std::unique_lock<std::mutex> lock(*mutexLock);
230+
std::lock_guard<std::mutex> lock(*mutexLock);
232231

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

273-
I->sycl_get_cm_buffer_params_index_ptr(surf_ind, &writeBase, &width,
274-
&mutexLock);
272+
I->sycl_get_cm_buffer_params_ptr(surf_ind, &writeBase, &width, &mutexLock);
275273

276-
std::unique_lock<std::mutex> lock(*mutexLock);
274+
std::lock_guard<std::mutex> lock(*mutexLock);
277275

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

461-
I->sycl_get_cm_buffer_params_index_ptr(surf_ind, &writeBase, &width,
462-
&mutexLock);
459+
I->sycl_get_cm_buffer_params_ptr(surf_ind, &writeBase, &width, &mutexLock);
463460
writeBase += global_offset;
464461

465-
std::unique_lock<std::mutex> lock(*mutexLock);
462+
std::lock_guard<std::mutex> lock(*mutexLock);
466463

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

597-
I->sycl_get_cm_buffer_params_index_ptr(surf_ind, &readBase, &width,
598-
&mutexLock);
594+
I->sycl_get_cm_buffer_params_ptr(surf_ind, &readBase, &width, &mutexLock);
599595
readBase += global_offset;
600596

601-
std::unique_lock<std::mutex> lock(*mutexLock);
597+
std::lock_guard<std::mutex> lock(*mutexLock);
602598

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

675-
I->sycl_get_cm_buffer_params_index_ptr(surf_ind, &readBase, &width,
676-
&mutexLock);
671+
I->sycl_get_cm_buffer_params_ptr(surf_ind, &readBase, &width, &mutexLock);
677672

678673
readBase += global_offset;
679-
std::unique_lock<std::mutex> lock(*mutexLock);
674+
std::lock_guard<std::mutex> lock(*mutexLock);
680675
for (int idx = 0; idx < N; idx++) {
681676
if (pred[idx]) {
682677
RestoredTy *addr =
@@ -727,10 +722,9 @@ __esimd_oword_ld(SurfIndAliasTy surf_ind, uint32_t addr)
727722
uint32_t width;
728723
std::mutex *mutexLock;
729724

730-
I->sycl_get_cm_buffer_params_index_ptr(surf_ind, &readBase, &width,
731-
&mutexLock);
725+
I->sycl_get_cm_buffer_params_ptr(surf_ind, &readBase, &width, &mutexLock);
732726

733-
std::unique_lock<std::mutex> lock(*mutexLock);
727+
std::lock_guard<std::mutex> lock(*mutexLock);
734728

735729
for (int idx = 0; idx < N; idx++) {
736730
if (addr >= width) {
@@ -768,9 +762,8 @@ __ESIMD_INTRIN
768762
} else {
769763
uint32_t width;
770764
std::mutex *mutexLock;
771-
I->sycl_get_cm_buffer_params_index_ptr(surf_ind, &ReadBase, &width,
772-
&mutexLock);
773-
std::unique_lock<std::mutex> lock(*mutexLock);
765+
I->sycl_get_cm_buffer_params_ptr(surf_ind, &ReadBase, &width, &mutexLock);
766+
std::lock_guard<std::mutex> lock(*mutexLock);
774767
}
775768

776769
ReadBase += global_offset;
@@ -812,9 +805,8 @@ __ESIMD_INTRIN void __esimd_scatter4_scaled(
812805
} else {
813806
uint32_t width;
814807
std::mutex *mutexLock;
815-
I->sycl_get_cm_buffer_params_index_ptr(surf_ind, &WriteBase, &width,
816-
&mutexLock);
817-
std::unique_lock<std::mutex> lock(*mutexLock);
808+
I->sycl_get_cm_buffer_params_ptr(surf_ind, &WriteBase, &width, &mutexLock);
809+
std::lock_guard<std::mutex> lock(*mutexLock);
818810
}
819811

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

934-
sycl::detail::getESIMDDeviceInterface()->sycl_get_cm_image_params_index_ptr(
926+
sycl::detail::getESIMDDeviceInterface()->sycl_get_cm_image_params_ptr(
935927
handle, &readBase, &imgWidth, &imgHeight, &bpp, &mutexLock);
936928

937-
std::unique_lock<std::mutex> lock(*mutexLock);
929+
std::lock_guard<std::mutex> lock(*mutexLock);
938930

939931
int x_pos_a, y_pos_a, offset, index;
940932

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

1064-
I->sycl_get_cm_image_params_index_ptr(handle, &writeBase, &imgWidth,
1065-
&imgHeight, &bpp, &mutexLock);
1056+
I->sycl_get_cm_image_params_ptr(handle, &writeBase, &imgWidth, &imgHeight,
1057+
&bpp, &mutexLock);
10661058

10671059
int x_pos_a, y_pos_a, offset;
10681060

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

1075-
std::unique_lock<std::mutex> lock(*mutexLock);
1067+
std::lock_guard<std::mutex> lock(*mutexLock);
10761068

10771069
for (int i = 0, k = 0; i < M; i++) {
10781070
for (int j = 0; j < N; j++) {

‎sycl/include/sycl/ext/intel/esimd/emu/detail/esimd_emulator_functions_v1.h

Lines changed: 4 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -46,15 +46,10 @@ void (*cm_fence_ptr)(void);
4646
char *(*sycl_get_surface_base_addr_ptr)(int);
4747
char *(*__cm_emu_get_slm_ptr)(void);
4848
void (*cm_slm_init_ptr)(size_t);
49-
void (*sycl_get_cm_buffer_params_ptr)(void *, char **, uint32_t *,
50-
std::mutex **);
51-
void (*sycl_get_cm_image_params_ptr)(void *, char **, uint32_t *, uint32_t *,
52-
uint32_t *, std::mutex **);
5349

5450
unsigned int (*sycl_get_cm_surface_index_ptr)(void *);
55-
void (*sycl_get_cm_buffer_params_index_ptr)(unsigned int, char **, uint32_t *,
56-
std::mutex **);
57-
void (*sycl_get_cm_image_params_index_ptr)(unsigned int, char **, uint32_t *,
58-
uint32_t *, uint32_t *,
59-
std::mutex **);
51+
void (*sycl_get_cm_buffer_params_ptr)(unsigned int, char **, uint32_t *,
52+
std::mutex **);
53+
void (*sycl_get_cm_image_params_ptr)(unsigned int, char **, uint32_t *,
54+
uint32_t *, uint32_t *, std::mutex **);
6055
/// @endcond ESIMD_EMU

0 commit comments

Comments
 (0)