use ciel_div helper for num_blocks

ggml-ci

Author: qnixsynapse

ggml/src/ggml-sycl/common.hpp

@@ -494,5 +494,9 @@ static __dpct_inline__ Tp* get_pointer(sycl::local_accessor<Tp, dim> acc) {
 
 int64_t downsample_sycl_global_range(int64_t accumulate_block_num, int64_t block_size);
 
+constexpr size_t ceil_div(const size_t m, const size_t n) {
+    return (m + n - 1) / n;
+}
+
 bool gpu_has_xmx(sycl::device &dev);
 #endif // GGML_SYCL_COMMON_HPP

ggml/src/ggml-sycl/element_wise.cpp

@@ -363,7 +363,7 @@ static void silu_sycl(const T *x, T *dst, const int k,
 template<typename T>
 static void sgn_sycl(const T * x, T * dst, const int k, queue_ptr stream) {
     // hard code for now
-    const int num_blocks = (k + 256 - 1) / 256;
+    const int num_blocks = ceil_div(k, 256);
     stream->parallel_for(
             sycl::nd_range<3>((sycl::range<3>(1, 1, num_blocks) * sycl::range(1, 1, 256)), sycl::range(1, 1, 256)), [=](sycl::nd_item<3> item_ct1) {
             sgn(x, dst, k, item_ct1);
@@ -373,7 +373,7 @@ static void sgn_sycl(const T * x, T * dst, const int k, queue_ptr stream) {
 template<typename T>
 static void abs_sycl(const T * x, T * dst, const int k, queue_ptr stream) {
     // hard code for now
-    const int num_blocks = (k + 256 - 1) / 256;
+    const int num_blocks = ceil_div(k, 256);
     stream->parallel_for(
             sycl::nd_range<3>((sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, 256)), sycl::range<3>(1, 1, 256)), [=](sycl::nd_item<3> item_ct1) {
             abs_op(x, dst, k, item_ct1);
@@ -384,7 +384,7 @@ static void abs_sycl(const T * x, T * dst, const int k, queue_ptr stream) {
 template<typename T>
 static void elu_sycl(const T * x, T * dst, const int k, queue_ptr stream) {
     // hard code for now
-    const int num_blocks = (k + 256 - 1) / 256;
+    const int num_blocks = ceil_div(k, 256);
     stream->parallel_for(
             sycl::nd_range<3>((sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, 256)), sycl::range<3>(1, 1, 256)), [=](sycl::nd_item<3> item_ct1) {
             elu_op(x, dst, k, item_ct1);