Skip to content

Commit 5725d40

Browse files
v-klochkovv-klochkov
authored
[ESIMD] Update the EMBARGO ats/pvc func names - remove esimd_ prefix (intel#383)
This patch also repleaces some obsolete APIs such as esimd_get_value() or flat_atomic() Signed-off-by: Vyacheslav N Klochkov <[email protected]>
1 parent 9d42a1f commit 5725d40

13 files changed

+33
-33
lines changed

‎SYCL_ESIMD_EMBARGO/ESIMD/EMBARGO/dpas_ats.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -51,9 +51,9 @@ int main(void) {
5151
mb.select<8, 1, 1, 1>(0, 0) = 4;
5252

5353
simd<int, Size> vc(0);
54-
vc = esimd_dpas<EsimdPrecisionType::S2, EsimdPrecisionType::S2, 8, 8, int,
55-
int, int, Size, 32, 32>(vc, ma.bit_cast_view<int>(),
56-
mb.bit_cast_view<int>());
54+
vc = dpas<EsimdPrecisionType::S2, EsimdPrecisionType::S2, 8, 8, int,
55+
int, int, Size, 32, 32>(vc, ma.bit_cast_view<int>(),
56+
mb.bit_cast_view<int>());
5757

5858
for (int i = 0; i < Size; i += VL) {
5959
simd<int, VL> output = vc.select<VL, 1>(i);

‎SYCL_ESIMD_EMBARGO/ESIMD/EMBARGO/dpas_pvc.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -51,9 +51,9 @@ int main(void) {
5151
mb.select<8, 1, 1, 1>(0, 0) = 4;
5252

5353
simd<int, Size> vc(0);
54-
vc = esimd_dpas<EsimdPrecisionType::S2, EsimdPrecisionType::S2, 8, 8, int,
55-
int, int, Size, 64, 32>(vc, ma.bit_cast_view<int>(),
56-
mb.bit_cast_view<int>());
54+
vc = dpas<EsimdPrecisionType::S2, EsimdPrecisionType::S2, 8, 8, int,
55+
int, int, Size, 64, 32>(vc, ma.bit_cast_view<int>(),
56+
mb.bit_cast_view<int>());
5757

5858
for (int i = 0; i < Size; i += VL) {
5959
simd<int, VL> output = vc.select<VL, 1>(i);

‎SYCL_ESIMD_EMBARGO/ESIMD/EMBARGO/dpasw_ats.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -51,9 +51,9 @@ int main(void) {
5151
mb.select<4, 2, 1, 1>(0, 0) = 4;
5252

5353
simd<int, Size> vc(0);
54-
vc = esimd_dpasw<EsimdPrecisionType::S2, EsimdPrecisionType::S2, 8, 8,
55-
int, int, int, Size, 32, 16>(vc, ma.bit_cast_view<int>(),
56-
mb.bit_cast_view<int>());
54+
vc = dpasw<EsimdPrecisionType::S2, EsimdPrecisionType::S2, 8, 8,
55+
int, int, int, Size, 32, 16>(vc, ma.bit_cast_view<int>(),
56+
mb.bit_cast_view<int>());
5757

5858
for (int i = 0; i < Size; i += VL) {
5959
simd<int, VL> output = vc.select<VL, 1>(i);

‎SYCL_ESIMD_EMBARGO/ESIMD/EMBARGO/histogram_2d_ats.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -190,7 +190,7 @@ int main(int argc, char *argv[]) {
190190
src = histogram.select<8, 1>(i);
191191

192192
#ifdef __SYCL_DEVICE_ONLY__
193-
flat_atomic<sycl::ext::intel::experimental::esimd::atomic_op::add, unsigned int, 8>(
193+
atomic_update<sycl::ext::intel::experimental::esimd::atomic_op::add, unsigned int, 8>(
194194
bins, offset, src, 1);
195195
offset += 8 * sizeof(unsigned int);
196196
#else

‎SYCL_ESIMD_EMBARGO/ESIMD/EMBARGO/histogram_raw_send_ats.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -79,8 +79,8 @@ ESIMD_INLINE void atomic_write(T *bins, simd<unsigned, n> offset,
7979
constexpr uint8_t isEOT = 0;
8080
constexpr uint8_t isSendc = 0;
8181

82-
esimd_raw_sends_load(oldDst, vAddr, src0, exDesc, desc, execSize, sfid,
83-
numSrc0, numSrc1, numDst, isEOT, isSendc, pred);
82+
raw_sends_load(oldDst, vAddr, src0, exDesc, desc, execSize, sfid,
83+
numSrc0, numSrc1, numDst, isEOT, isSendc, pred);
8484
}
8585

8686
int main(int argc, char *argv[]) {

‎SYCL_ESIMD_EMBARGO/ESIMD/EMBARGO/lsc_fence_pvc.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -20,7 +20,7 @@ implied warranties, other than those that are expressly stated in the License.
2020
// RUN: %clangxx -fsycl %s -DESIMD_GEN12_7 -o %t.out
2121
// RUNx: %GPU_RUN_PLACEHOLDER %t.out
2222

23-
#include "esimd_test_utils.hpp"
23+
#include "../esimd_test_utils.hpp"
2424

2525
#include <CL/sycl.hpp>
2626
#include <algorithm>

‎SYCL_ESIMD_EMBARGO/ESIMD/EMBARGO/lsc_flat_2d_pvc.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -17,7 +17,7 @@ implied warranties, other than those that are expressly stated in the License.
1717
// RUN: %clangxx -fsycl %s -DESIMD_GEN12_7 -o %t.out
1818
// RUN: %GPU_RUN_PLACEHOLDER %t.out
1919

20-
#include "esimd_test_utils.hpp"
20+
#include "../esimd_test_utils.hpp"
2121

2222
#include <CL/sycl.hpp>
2323
#include <algorithm>

‎SYCL_ESIMD_EMBARGO/ESIMD/EMBARGO/lsc_flat_atomic_cachehint_pvc.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -10,7 +10,7 @@
1010
// RUN: %clangxx -fsycl %s -DESIMD_GEN12_7 -o %t.out
1111
// RUN: %GPU_RUN_PLACEHOLDER %t.out
1212

13-
#include "esimd_test_utils.hpp"
13+
#include "../esimd_test_utils.hpp"
1414

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

‎SYCL_ESIMD_EMBARGO/ESIMD/EMBARGO/lsc_flat_pvc.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -17,7 +17,7 @@ implied warranties, other than those that are expressly stated in the License.
1717
// RUN: %clangxx -fsycl %s -DESIMD_GEN12_7 -o %t.out
1818
// RUN: %GPU_RUN_PLACEHOLDER %t.out
1919

20-
#include "esimd_test_utils.hpp"
20+
#include "../esimd_test_utils.hpp"
2121

2222
#include <CL/sycl.hpp>
2323
#include <algorithm>

‎SYCL_ESIMD_EMBARGO/ESIMD/EMBARGO/lsc_slm_pvc.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -17,7 +17,7 @@ implied warranties, other than those that are expressly stated in the License.
1717
// RUN: %clangxx -fsycl %s -DESIMD_GEN12_7 -o %t.out
1818
// RUN: %GPU_RUN_PLACEHOLDER %t.out
1919

20-
#include "esimd_test_utils.hpp"
20+
#include "../esimd_test_utils.hpp"
2121

2222
#include <CL/sycl.hpp>
2323
#include <algorithm>

‎SYCL_ESIMD_EMBARGO/ESIMD/EMBARGO/lsc_surf_pvc.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -17,7 +17,7 @@ implied warranties, other than those that are expressly stated in the License.
1717
// RUN: %clangxx -fsycl %s -DESIMD_GEN12_7 -o %t.out
1818
// RUN: %GPU_RUN_PLACEHOLDER %t.out
1919

20-
#include "esimd_test_utils.hpp"
20+
#include "../esimd_test_utils.hpp"
2121

2222
#include <CL/sycl.hpp>
2323
#include <algorithm>

‎SYCL_ESIMD_EMBARGO/ESIMD/EMBARGO/vadd_2d_stateless_pvc.cpp

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -61,19 +61,19 @@ int main(void) {
6161
constexpr unsigned yoff = 0;
6262

6363
simd<float, VL> va =
64-
esimd_2d_statelss_load<float, BLKX, BLKY, 1, false, false,
65-
CacheHint::Streaming, CacheHint::Uncached>(
64+
load_2d_stateless<float, BLKX, BLKY, 1, false, false,
65+
CacheHint::Streaming, CacheHint::Uncached>(
6666
A, width, height, pitch, xoff, yoff);
6767
simd<float, VL> vb =
68-
esimd_2d_statelss_load<float, BLKX, BLKY, 1, false, false,
69-
CacheHint::Streaming, CacheHint::Uncached>(
68+
load_2d_stateless<float, BLKX, BLKY, 1, false, false,
69+
CacheHint::Streaming, CacheHint::Uncached>(
7070
B, width, height, pitch, xoff, yoff);
7171

7272
simd<float, VL> vc = va + vb;
7373

74-
esimd_2d_statelss_store<float, BLKX, BLKY, CacheHint::Uncached,
75-
CacheHint::WriteBack>(C, width, height, pitch,
76-
xoff, yoff, vc);
74+
store_2d_stateless<float, BLKX, BLKY, CacheHint::Uncached,
75+
CacheHint::WriteBack>(C, width, height, pitch,
76+
xoff, yoff, vc);
7777
});
7878
});
7979
e.wait();

‎SYCL_ESIMD_EMBARGO/ESIMD/EMBARGO/vadd_raw_send_ats.cpp

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -28,15 +28,15 @@ ESIMD_INLINE simd<T, N> dwaligned_block_read(AccessorTy acc,
2828

2929
src0.select<1, 1>(2) = offset;
3030
uint32_t exDesc = 0xA;
31-
uint32_t desc = esimd_get_value(acc);
31+
SurfaceIndex desc = get_surface_index(acc);
3232
desc += 0x2284300;
3333
constexpr uint8_t execSize = 0x84;
3434
constexpr uint8_t sfid = 0xA;
3535
constexpr uint8_t numSrc0 = 0x1;
3636
constexpr uint8_t numDst = 0x2;
3737

38-
return esimd_raw_send_load(oldDst, src0, exDesc, desc, execSize, sfid,
39-
numSrc0, numDst);
38+
return raw_send_load(oldDst, src0, exDesc, desc, execSize, sfid,
39+
numSrc0, numDst);
4040
}
4141

4242
template <typename T, int N, typename AccessorTy>
@@ -46,15 +46,15 @@ ESIMD_INLINE void block_write1(AccessorTy acc, unsigned int offset,
4646

4747
src0.template select<1, 1>(2) = offset >> 4;
4848
uint32_t exDesc = 0x4A;
49-
uint32_t desc = esimd_get_value(acc);
49+
SurfaceIndex desc = get_surface_index(acc);
5050
desc += 0x20A0200;
5151
constexpr uint8_t execSize = 0x83;
5252
constexpr uint8_t sfid = 0xA;
5353
constexpr uint8_t numSrc0 = 0x1;
5454
constexpr uint8_t numSrc1 = 0x1;
5555

56-
return esimd_raw_sends_store(src0, data, exDesc, desc, execSize, sfid,
57-
numSrc0, numSrc1);
56+
return raw_sends_store(src0, data, exDesc, desc, execSize, sfid,
57+
numSrc0, numSrc1);
5858
}
5959

6060
template <typename T, int N, typename AccessorTy>
@@ -68,13 +68,13 @@ ESIMD_INLINE void block_write2(AccessorTy acc, unsigned int offset,
6868
src0_ref1.template select<1, 1>(2) = offset >> 4;
6969
src0_ref2 = data;
7070
uint32_t exDesc = 0xA;
71-
uint32_t desc = esimd_get_value(acc);
71+
SurfaceIndex desc = get_surface_index(acc);
7272
desc += 0x40A0200;
7373
constexpr uint8_t execSize = 0x83;
7474
constexpr uint8_t sfid = 0xA;
7575
constexpr uint8_t numSrc0 = 0x2;
7676

77-
return esimd_raw_send_store(src0, exDesc, desc, execSize, sfid, numSrc0);
77+
return raw_send_store(src0, exDesc, desc, execSize, sfid, numSrc0);
7878
}
7979

8080
int main(void) {

0 commit comments

Comments
 (0)