opencl: never enqueue a zero-size NDRange

clEnqueueNDRangeKernel returns CL_INVALID_GLOBAL_WORK_SIZE if any
global_work_size[i] is 0, and the CL_CHECK around it aborts the process. Most
drivers accept a zero-size range and do nothing, so the bug is invisible on
them; strict drivers enforce the spec and take the process down. This is a
spec violation on our side, not a device quirk.

Skip the enqueue when any global dimension is zero, at the single choke point
all dispatches funnel through. A dimension that is wrongly zero still surfaces
as a wrong result rather than being masked.

Scope was established by running the op suite on an UNGUARDED build: of the 74
ops in test-backend-ops, SET_ROWS is the only one that emits a zero-size
dispatch (a row-less tensor, ne01 == 0).

Verified on device:
  Adreno 619  aborted at SET_ROWS(f32->f32,i64,ne=[33,5,1,1],nr23=[2,3]) -> runs
  Adreno 642L same abort -> runs
  Adreno 840  op suite unchanged (lax driver, never hit the path)
  Adreno 740  op suite unchanged

Author: wanghqc

ggml/src/ggml-opencl/ggml-opencl.cpp

@@ -1066,6 +1066,17 @@ struct ggml_backend_opencl_context {
     }
 
     void enqueue_ndrange_kernel(cl_kernel kernel, cl_uint work_dim, size_t *global_work_size, size_t *local_work_size, const ggml_tensor * tensor) {
+        // An empty range is a no-op, but the spec says a zero global size is
+        // CL_INVALID_GLOBAL_WORK_SIZE, so enqueuing it is an error rather than nothing. Most
+        // drivers return CL_SUCCESS and do nothing; strict drivers (Adreno 642L, 619) return
+        // -63 and the CL_CHECK below aborts the process. Skip it, which is what every other
+        // driver effectively does. A dimension that is wrongly zero still shows up as a wrong
+        // result. Check first, before the profiling path, so no mode enqueues it.
+        for (cl_uint i = 0; i < work_dim; i++) {
+            if (global_work_size[i] == 0) {
+                return;
+            }
+        }
 #ifdef GGML_OPENCL_PROFILING
         cl_event evt;
         CL_CHECK(clEnqueueNDRangeKernel(queue, kernel, work_dim, NULL, global_work_size, local_work_size, 0, NULL, &evt));
@@ -11977,6 +11988,7 @@ static void ggml_cl_set_rows(ggml_backend_t backend, const ggml_tensor * src0, c
         (size_t)ne03};
     size_t local_work_size[] = {(size_t)nth0, (size_t)rows_per_workgroup, 1};
 
+    // ne01 == 0 makes global_work_size[0] zero here; enqueue_ndrange_kernel drops the empty range.
     backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst);
 }