vbatts/llama.cpp
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

SYCL ggml-sycl: pool2D use sycl::nan and remove if-else block

Browse source
This commit is contained in:
Akarshan Biswas 2024-12-12 19:44:06 +05:30
parent b828f4aa5f
commit 6b0848ceaf
No known key found for this signature in database
GPG key ID: 52A578A14B32134D
1 changed files with 3 additions and 9 deletions
Download patch file Download diff file Expand all files Collapse all files

10
ggml/src/ggml-sycl/ggml-sycl.cpp
View file

@ -1789,14 +1789,12 @@ static void pool2d_nchw_kernel(
const int ew = sycl::min(iw, start_w + kw); const int ew = sycl::min(iw, start_w + kw);
To res = 0; To res = 0;
bool op_valid = true;
switch (op) { switch (op) {
case GGML_OP_POOL_AVG: res = 0; break; case GGML_OP_POOL_AVG: res = 0; break;
case GGML_OP_POOL_MAX: res = -FLT_MAX; break; case GGML_OP_POOL_MAX: res = -FLT_MAX; break;
default: default:
res = NAN; res = (To) sycl::nan(uint32_t(0));
op_valid = false;
break; break;
} }
@ -1817,16 +1815,12 @@ static void pool2d_nchw_kernel(
case GGML_OP_POOL_AVG: res += (cur / (kh * kw)); break; case GGML_OP_POOL_AVG: res += (cur / (kh * kw)); break;
case GGML_OP_POOL_MAX: res = sycl::max(res, (To)cur); break; case GGML_OP_POOL_MAX: res = sycl::max(res, (To)cur); break;
default: default:
op_valid = false; res = (To) sycl::nan(uint32_t(0));
break; break;
} }
} }
} }
if (op_valid) {
o_ptr[cur_oh * ow + cur_ow] = res; o_ptr[cur_oh * ow + cur_ow] = res;
} else {
o_ptr[cur_oh * ow + cur_ow] = NAN;
}
} }
template <int qk, int qr, dequantize_kernel_t dq> template <int qk, int qr, dequantize_kernel_t dq>