opencl: check that the compiler accepts the Adreno large-buffer option

Advertising cl_qcom_large_buffer and being able to build for it are different
questions. The Adreno 850 (compiler E17.51) reports the extension and then
rejects -qcom-enable-large-buffer with "invalid option", so every kernel build
fails and the process exits. Setting GGML_OPENCL_ADRENO_USE_LARGE_BUFFER=1
there does not cost throughput, it stops the backend working at all.

Build a trivial program once with the option and keep large-buffer mode only
if that succeeds, logging the downgrade so the fallback is visible. This sits
beside the existing check for whether the driver advertises the extension at
all, which has the same shape.

Costs one small program build, and only when the variable is set. On the 850,
test-backend-ops -o CPY with the variable set goes from a compile error and
exit 1 to 138/138 passing.

Author: wanghqc

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

@@ -5752,6 +5752,21 @@ static std::vector<ggml_backend_device> ggml_opencl_probe_devices(ggml_backend_r
     return found_devices;
 }
 
+// Advertising an extension and being able to build for it are different questions: a driver can
+// report cl_qcom_large_buffer and then reject -qcom-enable-large-buffer as an invalid option,
+// which fails every kernel build and takes the process down. Build a trivial program and ask.
+static bool ggml_opencl_compiler_accepts(cl_context context, cl_device_id device, const char * opt) {
+    const char * src = "__kernel void ggml_cl_probe(void) {}\n";
+    cl_int err = CL_SUCCESS;
+    cl_program p = clCreateProgramWithSource(context, 1, &src, NULL, &err);
+    if (err != CL_SUCCESS || p == NULL) {
+        return false;
+    }
+    err = clBuildProgram(p, 1, &device, opt, NULL, NULL);
+    clReleaseProgram(p);
+    return err == CL_SUCCESS;
+}
+
 static void ggml_opencl_print_backend_info(ggml_backend_opencl_device_context * dev_ctx) {
     GGML_ASSERT(dev_ctx);
     GGML_ASSERT(dev_ctx->backend_ctx);
@@ -5807,6 +5822,10 @@ static void ggml_opencl_print_backend_info(ggml_backend_opencl_device_context *
         if (!backend_ctx->adreno_has_large_buffer) {
             GGML_LOG_INFO("ggml_opencl: Adreno large buffer requested but not supported by driver, will use regular buffer\n");
             backend_ctx->adreno_use_large_buffer = false;
+        } else if (!ggml_opencl_compiler_accepts(backend_ctx->context, backend_ctx->device,
+                                                 " -qcom-enable-large-buffer ")) {
+            GGML_LOG_INFO("ggml_opencl: Adreno large buffer advertised but rejected by the compiler, will use regular buffer\n");
+            backend_ctx->adreno_use_large_buffer = false;
         } else {
             GGML_LOG_INFO("ggml_opencl: Adreno large buffer enabled\n");
         }