Skip to content

Commit 6b0848c

Browse files
committed
SYCL ggml-sycl: pool2D use sycl::nan and remove if-else block
1 parent b828f4a commit 6b0848c

File tree

1 file changed

+3
-9
lines changed

1 file changed

+3
-9
lines changed

ggml/src/ggml-sycl/ggml-sycl.cpp

Lines changed: 3 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -1789,14 +1789,12 @@ static void pool2d_nchw_kernel(
17891789
const int ew = sycl::min(iw, start_w + kw);
17901790

17911791
To res = 0;
1792-
bool op_valid = true;
17931792

17941793
switch (op) {
17951794
case GGML_OP_POOL_AVG: res = 0; break;
17961795
case GGML_OP_POOL_MAX: res = -FLT_MAX; break;
17971796
default:
1798-
res = NAN;
1799-
op_valid = false;
1797+
res = (To) sycl::nan(uint32_t(0));
18001798
break;
18011799
}
18021800

@@ -1817,16 +1815,12 @@ static void pool2d_nchw_kernel(
18171815
case GGML_OP_POOL_AVG: res += (cur / (kh * kw)); break;
18181816
case GGML_OP_POOL_MAX: res = sycl::max(res, (To)cur); break;
18191817
default:
1820-
op_valid = false;
1818+
res = (To) sycl::nan(uint32_t(0));
18211819
break;
18221820
}
18231821
}
18241822
}
1825-
if (op_valid) {
1826-
o_ptr[cur_oh * ow + cur_ow] = res;
1827-
} else {
1828-
o_ptr[cur_oh * ow + cur_ow] = NAN;
1829-
}
1823+
o_ptr[cur_oh * ow + cur_ow] = res;
18301824
}
18311825

18321826
template <int qk, int qr, dequantize_kernel_t dq>

0 commit comments

Comments
 (0)