5
5
// Check that full compilation works:
6
6
// RUN: %clangxx -fsycl -fno-sycl-device-code-split-esimd -Xclang -fsycl-allow-func-ptr %s -o %t.out
7
7
// RUN: env IGC_VCSaveStackCallLinkage=1 IGC_VCDirectCallsOnly=1 %GPU_RUN_PLACEHOLDER %t.out
8
- //
9
- // VISALTO enable run
10
- // RUN: env IGC_VISALTO=63 IGC_VCSaveStackCallLinkage=1 IGC_VCDirectCallsOnly=1 %GPU_RUN_PLACEHOLDER %t.out
11
8
#include < sycl/detail/boost/mp11.hpp>
12
9
#include < sycl/ext/intel/esimd.hpp>
13
10
#include < sycl/ext/oneapi/experimental/invoke_simd.hpp>
@@ -23,7 +20,7 @@ constexpr int VL = 16;
23
20
24
21
[[intel::device_indirectly_callable]] simd<float , VL>
25
22
SIMD_CALLEE (simd<float , VL> va, simd_mask<float , VL> mask) SYCL_ESIMD_FUNCTION {
26
- esimd::simd<float , VL> ret;
23
+ esimd::simd<float , VL> ret ( 0 ) ;
27
24
esimd::simd_mask<VL> emask;
28
25
for (int i = 0 ; i < VL; i++)
29
26
emask[i] = static_cast <bool >(mask[i]);
@@ -50,7 +47,6 @@ int main() {
50
47
M[i] = i % 2 ;
51
48
}
52
49
53
- auto ctxt = q.get_context ();
54
50
sycl::buffer<float > ABuf (A);
55
51
sycl::buffer<float > CBuf (C);
56
52
sycl::buffer<bool > MBuf (M);
@@ -68,7 +64,7 @@ int main() {
68
64
sycl::accessor M_acc{MBuf, cgh, sycl::read_only};
69
65
cgh.parallel_for (Range, [=](nd_item<1 > ndi) {
70
66
sub_group sg = ndi.get_sub_group ();
71
- uint32_t wi_id = ndi.get_global_linear_id () + sg. get_local_id () ;
67
+ uint32_t wi_id = ndi.get_global_linear_id ();
72
68
float res = invoke_simd (sg, SIMD_CALLEE, A_acc[wi_id], M_acc[wi_id]);
73
69
C_acc[wi_id] = res;
74
70
});
0 commit comments