Skip to content

Commit b17089d

Browse files
author
Alberto Cabrera
committed
Added smaller reproducer, changed verbosity as printing hid the issue
Signed-off-by: Alberto Cabrera <[email protected]>
1 parent c063ed8 commit b17089d

File tree

3 files changed

+102
-28
lines changed

3 files changed

+102
-28
lines changed

‎.github/workflows/sycl-windows-precommit.yml‎

Lines changed: 22 additions & 22 deletions
Original file line numberDiff line numberDiff line change
@@ -34,25 +34,25 @@ jobs:
3434
detect_changes:
3535
uses: ./.github/workflows/sycl-detect-changes.yml
3636

37-
build:
38-
needs: [detect_changes]
39-
if: |
40-
always() && success()
41-
&& github.repository == 'intel/llvm'
42-
uses: ./.github/workflows/sycl-windows-build.yml
43-
with:
44-
changes: ${{ needs.detect_changes.outputs.filters }}
45-
46-
e2e:
47-
needs: build
48-
# Continue if build was successful.
49-
if: |
50-
always()
51-
&& !cancelled()
52-
&& needs.build.outputs.build_conclusion == 'success'
53-
uses: ./.github/workflows/sycl-windows-run-tests.yml
54-
with:
55-
name: Intel GEN12 Graphics with Level Zero
56-
runner: '["Windows","gen12"]'
57-
sycl_toolchain_archive: ${{ needs.build.outputs.artifact_archive_name }}
58-
extra_lit_opts: --param gpu-intel-gen12=True
37+
# build:
38+
# needs: [detect_changes]
39+
# if: |
40+
# always() && success()
41+
# && github.repository == 'intel/llvm'
42+
# uses: ./.github/workflows/sycl-windows-build.yml
43+
# with:
44+
# changes: ${{ needs.detect_changes.outputs.filters }}
45+
#
46+
# e2e:
47+
# needs: build
48+
# # Continue if build was successful.
49+
# if: |
50+
# always()
51+
# && !cancelled()
52+
# && needs.build.outputs.build_conclusion == 'success'
53+
# uses: ./.github/workflows/sycl-windows-run-tests.yml
54+
# with:
55+
# name: Intel GEN12 Graphics with Level Zero
56+
# runner: '["Windows","gen12"]'
57+
# sycl_toolchain_archive: ${{ needs.build.outputs.artifact_archive_name }}
58+
# extra_lit_opts: --param gpu-intel-gen12=True

‎sycl/test-e2e/syclcompat/math/math_fixt.hpp‎

Lines changed: 8 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -142,7 +142,9 @@ template <typename... ValueT> struct should_skip {
142142
if (syclcompat::detail::isnan(RESULT[i])) { \
143143
assert(syclcompat::detail::isnan(EXPECTED[i])); \
144144
} else { \
145-
assert(fabs(RESULT[i] - EXPECTED[i]) < ERROR_TOLERANCE); \
145+
assert((fabs(RESULT[i] - EXPECTED[i]) < ERROR_TOLERANCE) || \
146+
!(std::cerr << "-- " << RESULT[i] << " - " << EXPECTED[i] \
147+
<< " < " << ERROR_TOLERANCE << "-- ")); \
146148
} \
147149
} \
148150
} else { \
@@ -209,7 +211,7 @@ class BinaryOpTestLauncher : OpTestLauncher {
209211
syclcompat::wait();
210212
syclcompat::memcpy<ResultT>(&res_h_, res_, data_size_);
211213

212-
CHECK_PRINT(ResultT, res_h_, expected);
214+
// CHECK_PRINT(ResultT, res_h_, expected);
213215
CHECK(ResultT, res_h_, expected);
214216
};
215217

@@ -223,7 +225,7 @@ class BinaryOpTestLauncher : OpTestLauncher {
223225
syclcompat::wait();
224226
syclcompat::memcpy<ResultT>(&res_h_, res_, data_size_);
225227

226-
CHECK_PRINT(ResultT, res_h_, expected);
228+
// CHECK_PRINT(ResultT, res_h_, expected);
227229
CHECK(ResultT, res_h_, expected);
228230
};
229231
template <auto Kernel>
@@ -241,7 +243,7 @@ class BinaryOpTestLauncher : OpTestLauncher {
241243
syclcompat::memcpy<bool>(&res_hi_h_, res_hi_, 1);
242244
syclcompat::memcpy<bool>(&res_lo_h_, res_lo_, 1);
243245

244-
CHECK_PRINT(ResultT, res_h_, expected);
246+
// CHECK_PRINT(ResultT, res_h_, expected);
245247
CHECK(ResultT, res_h_, expected);
246248
assert(res_hi_h_ == expected_hi);
247249
assert(res_lo_h_ == expected_lo);
@@ -282,7 +284,7 @@ class UnaryOpTestLauncher : OpTestLauncher {
282284
syclcompat::wait();
283285
syclcompat::memcpy<ResultT>(&res_h_, res_, data_size_);
284286

285-
CHECK_PRINT(ResultT, res_h_, expected);
287+
// CHECK_PRINT(ResultT, res_h_, expected);
286288
CHECK(ResultT, res_h_, expected);
287289
}
288290
};
@@ -334,7 +336,7 @@ class TernaryOpTestLauncher : OpTestLauncher {
334336
syclcompat::wait();
335337
syclcompat::memcpy<ResultT>(&res_h_, res_, data_size_);
336338

337-
CHECK_PRINT(ResultT, res_h_, expected);
339+
// CHECK_PRINT(ResultT, res_h_, expected);
338340
CHECK(ResultT, res_h_, expected);
339341
};
340342
};

‎sycl/test-e2e/syclcompat/math/repro.cpp‎

Lines changed: 72 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,72 @@
1+
// DEFINE: %{mathflags} = %if cl_options %{/clang:-fno-fast-math%} %else %{-fno-fast-math%}
2+
3+
// RUN: %{build} %{mathflags} -o %t.out
4+
// RUN: %{run} %t.out
5+
6+
#include "sycl/detail/builtins/builtins.hpp"
7+
#include <sycl/usm.hpp>
8+
9+
#include <syclcompat/dims.hpp>
10+
#include <syclcompat/math.hpp>
11+
12+
inline void fmax_nan_kernel(float *a, float *b, sycl::vec<float, 2> *r) {
13+
*r = syclcompat::fmax_nan(*a, *b);
14+
}
15+
16+
void test_container_syclcompat_fmax_nan() {
17+
std::cout << __PRETTY_FUNCTION__ << std::endl;
18+
sycl::queue q;
19+
20+
sycl::range global{1};
21+
sycl::range local{1};
22+
sycl::nd_range ndr{global, local};
23+
24+
const sycl::vec<float, 2> op1 = {5.0f, 10.0f};
25+
const sycl::vec<float, 2> op2 = {10.0f, 5.0f};
26+
const sycl::vec<float, 2> expected{static_cast<float>(10),
27+
static_cast<float>(10)};
28+
sycl::vec<float, 2> res;
29+
30+
sycl::vec<float, 2> *op1_d = sycl::malloc_device<sycl::vec<float, 2>>(1, q);
31+
sycl::vec<float, 2> *op2_d = sycl::malloc_device<sycl::vec<float, 2>>(1, q);
32+
sycl::vec<float, 2> *res_d = sycl::malloc_device<sycl::vec<float, 2>>(1, q);
33+
34+
q.memcpy(op1_d, &op1, sizeof(sycl::vec<float, 2>));
35+
q.memcpy(op2_d, &op2, sizeof(sycl::vec<float, 2>));
36+
q.submit([&](sycl::handler &cgh) {
37+
cgh.parallel_for(ndr, [=](sycl::nd_item<1> nd_item) {
38+
*res_d = syclcompat::fmax_nan(*op1_d, *op2_d);
39+
});
40+
}).wait_and_throw();
41+
q.memcpy(&res, res_d, sizeof(sycl::vec<float, 2>)).wait();
42+
43+
constexpr float ERROR_TOLERANCE = 1e-6;
44+
for (size_t i = 0; i < 2; i++) {
45+
assert((res[i] - expected[i]) < ERROR_TOLERANCE);
46+
}
47+
48+
const sycl::vec<float, 2> op3 = {sycl::nan(static_cast<unsigned int>(0)),
49+
sycl::nan(static_cast<unsigned int>(0))};
50+
51+
q.memcpy(op2_d, &op3, sizeof(sycl::vec<float, 2>));
52+
q.submit([&](sycl::handler &cgh) {
53+
cgh.parallel_for(ndr, [=](sycl::nd_item<1> nd_item) {
54+
*res_d = syclcompat::fmax_nan(*op1_d, *op2_d);
55+
});
56+
}).wait_and_throw();
57+
q.memcpy(&res, res_d, sizeof(sycl::vec<float, 2>)).wait();
58+
59+
for (size_t i = 0; i < 2; i++) {
60+
assert(sycl::isnan(res[i]));
61+
}
62+
63+
sycl::free(op1_d, q);
64+
sycl::free(op2_d, q);
65+
sycl::free(res_d, q);
66+
}
67+
68+
int main() {
69+
test_container_syclcompat_fmax_nan();
70+
71+
return 0;
72+
}

0 commit comments

Comments
 (0)