restore local workgroup size adjustments for large inputs

Author: qnixsynapse

ggml/src/ggml-sycl/im2col.cpp

@@ -58,11 +58,13 @@ static void im2col_sycl_internal(const float * x, T * dst, int64_t IW, int64_t I
                                  int64_t KH, int64_t IC, int64_t batch, int64_t batch_offset, int64_t offset_delta,
                                  int s0, int s1, int p0, int p1, int d0, int d1, queue_ptr stream) {
     const int64_t parallel_elements = OW * KW * KH;
-    const int64_t block_size_x      = SYCL_IM2COL_BLOCK_SIZE;
-    const int64_t num_groups_x      = (parallel_elements + block_size_x - 1) / block_size_x;
+    const int64_t num_blocks      = (parallel_elements + SYCL_IM2COL_BLOCK_SIZE - 1) / SYCL_IM2COL_BLOCK_SIZE;
 
-    sycl::range<3> block_nums(batch * IC, OH, num_groups_x);
-    sycl::range<3> local_range(1, 1, block_size_x);
+    // decrease global range when it exceeds the max int
+    int64_t local_size = downsample_sycl_global_range(batch * IC * OH * num_blocks, SYCL_IM2COL_BLOCK_SIZE);
+
+    sycl::range<3> block_nums(batch * IC, OH, num_blocks);
+    sycl::range<3> local_range(1, 1, local_size);
 
     const int64_t CHW = IC * KH * KW;