Brunobkr/llama.cpp_AlgMor24_github

1#define CL_TARGET_OPENCL_VERSION GGML_OPENCL_TARGET_VERSION2#define CL_USE_DEPRECATED_OPENCL_1_2_APIS3 4// suppress warnings in CL headers for GCC and Clang5#pragma GCC diagnostic ignored "-Woverlength-strings"6#ifdef __clang__7#pragma GCC diagnostic ignored "-Wgnu-anonymous-struct"8#endif9 10#include "ggml-opencl.h"11#include "ggml-backend.h"12#include "ggml-impl.h"13#include "ggml-backend-impl.h"14#include "ggml.h"15 16#include "cl-program-cache.h"17 18#ifdef GGML_OPENCL_USE_ADRENO_BIN_KERNELS19#include "libdl.h"20#ifdef _WIN3221#define KERNEL_LIB_NAME "adreno-opencl-kernels.dll"22#else23#define KERNEL_LIB_NAME "libadreno-opencl-kernels.so"24#endif // _WIN3225#endif // GGML_OPENCL_USE_ADRENO_BIN_KERNELS26 27typedef const void * (*get_adreno_bin_kernel_func_t)(28    const char * name,29    const char * gpu_name,30    const char * compiler_ver,31    size_t     * out_size32);33 34#include <CL/cl.h>35 36#include <inttypes.h>37#include <string.h>38 39#include <cstddef>40#include <cstdint>41#include <fstream>42#include <vector>43#include <string>44#include <cmath>45#include <map>46#include <memory>47#include <charconv>48#include <mutex>49#include <regex>50#include <set>51#include <unordered_set>52 53#undef MIN54#undef MAX55#define MIN(a, b) ((a) < (b) ? (a) : (b))56#define MAX(a, b) ((a) > (b) ? (a) : (b))57#define CEIL_DIV(M, N) (((M) + (N)-1) / (N))58 59#define UNUSED(x) (void)(x)60 61#define CL_CHECK(err)                                               \62    do {                                                            \63        cl_int err_ = (err);                                        \64        if (err_ != CL_SUCCESS) {                                   \65            GGML_LOG_ERROR("ggml_opencl: %s error %d at %s:%d\n",  \66                #err, err_, __FILE__, __LINE__);                    \67            GGML_ASSERT(0);                                         \68        }                                                           \69    } while (0)70 71//------------------------------------------------------------------------------72// OpenCL73//------------------------------------------------------------------------------74 75bool ggml_cl_compute_forward(ggml_backend_t backend, struct ggml_tensor * tensor);76 77static bool ggml_cl_is_q4_0_soa(const ggml_tensor * tensor);78static bool ggml_cl_is_q8_0_soa(const ggml_tensor * tensor);79static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);80 81// See https://gmplib.org/~tege/divcnst-pldi94.pdf figure 4.1.82// Precompute mp (m' in the paper) and L such that division83// can be computed using a multiply (high 32b of 64b result)84// and a shift:85//86// n/d = (mulhi(n, mp) + n) >> L;87struct fastdiv_vals {88    uint32_t mp;89    uint32_t L;90    uint32_t d;91    uint32_t pad;92};93static_assert(sizeof(fastdiv_vals) == 16, "fastdiv_vals size incorrect");94 95static fastdiv_vals init_fastdiv_values(uint64_t d_64) {96    GGML_ASSERT(d_64 != 0);97    GGML_ASSERT(d_64 <= std::numeric_limits<uint32_t>::max());98 99    uint32_t d = (uint32_t)d_64;100 101    // compute L = ceil(log2(d));102    uint32_t L = 0;103    while (L < 32 && (uint32_t{ 1 } << L) < d) {104        L++;105    }106 107    uint32_t mp = (uint32_t) ((uint64_t{ 1 } << 32) * ((uint64_t{ 1 } << L) - d) / d + 1);108    // pack divisor as well to reduce error surface109    return { mp, L, d, 0 };110}111 112enum GPU_FAMILY {113    ADRENO,114    INTEL,115    UNKNOWN,116};117 118enum ADRENO_GPU_GEN {119    ADRENO_UNKNOWN,120    A6X,121    A7X,122    A8X,123    X1E,124    X2E,125};126 127enum ADRENO_CL_COMPILER_TYPE {128    E031,129    E17,130    DX,131};132 133struct ggml_cl_version {134    cl_uint major = 0;135    cl_uint minor = 0;136};137 138 139struct ggml_cl_compiler_version {140    ADRENO_CL_COMPILER_TYPE type;141    int major = -1;142    int minor = -1;143    int patch = -1;144 145    bool same(ADRENO_CL_COMPILER_TYPE t, int x, int y, int z) const {146        return major == x && minor == y && patch == z && type == t;147    }148    bool newer_than(ADRENO_CL_COMPILER_TYPE t, int x, int y, int z) const {149        return major*10000 + minor*100 + patch > x*10000 + y*100 + z && type == t;150    }151    bool newer_than_or_same(ADRENO_CL_COMPILER_TYPE t, int x, int y, int z) const {152        return same(t, x, y, z) || newer_than(t, x, y, z);153    }154};155 156static size_t align_to(size_t value, size_t to_alignment) {157    GGML_ASSERT(to_alignment && "Invalid alignment (must be non-zero)");158    GGML_ASSERT((to_alignment & (to_alignment - 1)) == 0 && "to_alignment must be power-of-two");159 160    return ((value + to_alignment - 1) / to_alignment) * to_alignment;161}162 163 164// Parses a version string of form "XX.YY ". On an error returns ggml_cl_version with all zeroes.165static ggml_cl_version parse_cl_version(std::string_view str) {166    size_t major_str_begin = 0;167    size_t major_str_end   = str.find(".", major_str_begin);168    if (major_str_end == std::string::npos) {169        return {};170    }171 172    size_t minor_str_begin = major_str_end + 1;173    size_t minor_str_end   = str.find(" ", minor_str_begin);174    if (minor_str_end == std::string::npos) {175        return {};176    }177 178    cl_uint version_major;179    if (std::from_chars(str.data() + major_str_begin, str.data() + major_str_end, version_major).ec != std::errc{}) {180        return {};181    }182 183    cl_uint version_minor;184    if (std::from_chars(str.data() + minor_str_begin, str.data() + minor_str_end, version_minor).ec != std::errc{}) {185        return {};186    }187    return { version_major, version_minor };188}189 190// Returns OpenCL platform's version. On an error returns ggml_cl_version with all zeroes.191static ggml_cl_version get_opencl_platform_version(cl_platform_id platform) {192    size_t param_size;193    CL_CHECK(clGetPlatformInfo(platform, CL_PLATFORM_VERSION, 0, nullptr, &param_size));194    std::unique_ptr<char[]> param_storage(new char[param_size]);195    CL_CHECK(clGetPlatformInfo(platform, CL_PLATFORM_VERSION, param_size, param_storage.get(), nullptr));196 197    auto              param_value    = std::string_view(param_storage.get(), param_size);198    const std::string version_prefix = "OpenCL ";  // Suffix: "XX.YY <platform-specific-info>"199    if (param_value.find(version_prefix) != 0) {200        return {};201    }202    param_value.remove_prefix(version_prefix.length());203    return parse_cl_version(param_value);204}205 206// Return a version to use in OpenCL C compilation. On an error returns ggml_cl_version with all zeroes.207static ggml_cl_version get_opencl_c_version(ggml_cl_version platform_version, cl_device_id device) {208    size_t param_size;209 210#if CL_TARGET_OPENCL_VERSION >= 300211    if (platform_version.major >= 3) {212        CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_ALL_VERSIONS, 0, nullptr, &param_size));213        if (!param_size) {214            return {};215        }216 217        std::unique_ptr<cl_name_version[]> versions(new cl_name_version[param_size]);218        CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_ALL_VERSIONS, param_size, versions.get(), nullptr));219        unsigned versions_count = param_size / sizeof(cl_name_version);220 221        cl_version version_max = 0;222        for (unsigned i = 0; i < versions_count; i++) {223            version_max = std::max<cl_version>(versions[i].version, version_max);224        }225 226        return { CL_VERSION_MAJOR(version_max), CL_VERSION_MINOR(version_max) };227    }228#else229    GGML_UNUSED(platform_version);230#endif  // CL_TARGET_OPENCL_VERSION >= 300231 232    CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_VERSION, 0, nullptr, &param_size));233    if (!param_size) {234        return {};235    }236 237    std::unique_ptr<char[]> param_storage(new char[param_size]);238    CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_VERSION, param_size, param_storage.get(), nullptr));239    auto param_value = std::string_view(param_storage.get(), param_size);240 241    const std::string version_prefix = "OpenCL C ";  // Suffix: "XX.YY <platform-specific-info>"242    if (param_value.find(version_prefix) != 0) {243        return {};244    }245    param_value.remove_prefix(version_prefix.length());246 247    return parse_cl_version(param_value);248}249 250static ADRENO_GPU_GEN get_adreno_gpu_gen(const char *device_name) {251    if (strstr(device_name, "610") || strstr(device_name, "612") ||252        strstr(device_name, "613") || strstr(device_name, "615") ||253        strstr(device_name, "616") || strstr(device_name, "618") ||254        strstr(device_name, "619") || strstr(device_name, "620") ||255        strstr(device_name, "630") || strstr(device_name, "640") ||256        strstr(device_name, "642") || strstr(device_name, "643") ||257        strstr(device_name, "644") || strstr(device_name, "650") ||258        strstr(device_name, "660") || strstr(device_name, "663") ||259        strstr(device_name, "680") || strstr(device_name, "685") ||260        strstr(device_name, "690")) {261        return ADRENO_GPU_GEN::A6X;262    }263 264    if (strstr(device_name, "730") ||265        strstr(device_name, "740") ||266        strstr(device_name, "750")) {267        return ADRENO_GPU_GEN::A7X;268    }269 270    if (strstr(device_name, "810") ||271        strstr(device_name, "830") ||272        strstr(device_name, "840") ||273        strstr(device_name, "850")) {274        return ADRENO_GPU_GEN::A8X;275    }276 277    if (strstr(device_name, "X1")) {278        return ADRENO_GPU_GEN::X1E;279    }280 281    if (strstr(device_name, "X2")) {282        return ADRENO_GPU_GEN::X2E;283    }284 285    return ADRENO_GPU_GEN::ADRENO_UNKNOWN;286}287 288static ggml_cl_compiler_version get_adreno_cl_compiler_version(const char *driver_version) {289    std::string driver_ver_str(driver_version);290    ADRENO_CL_COMPILER_TYPE type = ADRENO_CL_COMPILER_TYPE::E031;291    size_t compiler_ver_pos = driver_ver_str.find("E031");292    size_t compiler_ver_len = 13;293    size_t compiler_major_offset = 5;294    size_t compiler_minor_offset = 8;295    size_t compiler_patch_offset = 11;296 297    if (compiler_ver_pos == std::string::npos) {298        compiler_ver_pos = driver_ver_str.find("E17");299        if (compiler_ver_pos != std::string::npos) {300            type = ADRENO_CL_COMPILER_TYPE::E17;301            compiler_ver_len = 12;302            compiler_major_offset = 4;303            compiler_minor_offset = 7;304            compiler_patch_offset = 10;305        }306    }307 308    if (compiler_ver_pos == std::string::npos) {309        compiler_ver_pos = driver_ver_str.find("DX");310        if (compiler_ver_pos == std::string::npos) {311            return {};312        }313        type = ADRENO_CL_COMPILER_TYPE::DX;314        compiler_ver_len = 11;315        compiler_major_offset = 3;316        compiler_minor_offset = 6;317        compiler_patch_offset = 9;318    }319 320    std::string compiler_ver_str = driver_ver_str.substr(compiler_ver_pos, compiler_ver_len);321    int major = std::atoi(compiler_ver_str.substr(compiler_major_offset, 2).c_str());322    int minor = std::atoi(compiler_ver_str.substr(compiler_minor_offset, 2).c_str());323    int patch = std::atoi(compiler_ver_str.substr(compiler_patch_offset, 2).c_str());324    return { type, major, minor, patch };325}326 327// cl buffer wrapper328struct ggml_cl_buffer {329    cl_mem buffer;330    size_t size;331 332    ggml_cl_buffer()333        : buffer(nullptr), size(0) {}334 335    ~ggml_cl_buffer() {336        if (buffer) {337            CL_CHECK(clReleaseMemObject(buffer));338        }339    }340 341    void allocate(cl_context context, size_t new_size) {342        if (new_size > size) {343            size = new_size;344            if (buffer) {345                CL_CHECK(clReleaseMemObject(buffer));346            }347            cl_int err;348            CL_CHECK((buffer = clCreateBuffer(context, CL_MEM_READ_WRITE, size, NULL, &err), err));349        }350    }351};352 353// Profiling354struct ProfilingInfo {355    std::string op_name;356    std::string kernel_name;357 358    cl_kernel kernel;359    cl_event evt;360 361    cl_ulong cmd_queued;362    cl_ulong cmd_submit;363    cl_ulong cmd_start;364    cl_ulong cmd_end;365    cl_ulong overhead_start;366    cl_ulong overhead_end;367    // For the times below, see spec for clGetEventProfilingInfo368    // The time kernel spent in cmd queue - SUBMIT - QUEUED369    cl_ulong cmd_queued_duration_ns;370    // The time kernel spent for submission - START - SUBMIT371    cl_ulong cmd_submit_duration_ns;372    // Kernel execution time in nanoseconds - END - START373    cl_ulong cmd_duration_ns;374    // The time for the kernel to complete - COMPLETE - END375    cl_ulong cmd_complete_duration_ns;376    // Total time to finish the kernel - COMPLETE - QUEUED377    cl_ulong cmd_total_duration_ns;378    // Global and local work sizes.379    size_t global_size[3];380    size_t local_size[3];381    // Op output size.382    size_t output_size[4];383};384 385static void populateProfilingInfo(386        ProfilingInfo& info, cl_event evt, cl_kernel kernel, cl_uint work_dim,387        size_t global_size[3], size_t local_size[3],388        const ggml_tensor * tensor) {389    info.op_name     = tensor->name;390    info.kernel      = kernel;391    info.evt         = evt;392 393    // 0 means not specified, e.g., 2D workgroup, or NULL for driver to choose394    info.local_size[0] = 0;395    info.local_size[1] = 0;396    info.local_size[2] = 0;397 398    info.global_size[0] = 0;399    info.global_size[1] = 0;400    info.global_size[2] = 0;401 402    if (local_size) {403        for (cl_uint i = 0; i < work_dim; ++i) {404            info.local_size[i] = local_size[i];405        }406    }407 408    for (cl_uint i = 0; i < work_dim; ++i) {409        info.global_size[i] = global_size[i];410    }411 412    info.output_size[0] = tensor->ne[0];413    info.output_size[1] = tensor->ne[1];414    info.output_size[2] = tensor->ne[2];415    info.output_size[3] = tensor->ne[3];416}417 418struct ggml_backend_opencl_context;419 420// backend device context421struct ggml_backend_opencl_device_context {422    cl_platform_id platform;423    std::string platform_name;424 425    cl_device_id   device;426    std::string    device_name;427    cl_device_type device_type;428    std::string    device_version;429 430    // Initialized by ggml_cl_init().431    ggml_backend_opencl_context * backend_ctx = nullptr;432 433    // Initialized by ggml_backend_opencl_device_get_buffer_type()434    ggml_backend_buffer_type buffer_type;435 436    cl_context context = nullptr;437 438    GPU_FAMILY     gpu_family = GPU_FAMILY::UNKNOWN;439    ADRENO_GPU_GEN adreno_gen = ADRENO_GPU_GEN::ADRENO_UNKNOWN;440 441    std::regex *opfilter = nullptr; // regex of ops to not claim442    std::string opfilter_str = ""; // regex string for opfilter443    size_t global_mem_size = 0;444};445 446// Lazily-compiled flash-attention kernels and their per-(dk,dv) tile metadata.447// One map per (Q/KV dtype, decode/prefill, split) combination; the int maps448// hold tile dims (bm/bn), workgroup sizes and the n_kv split thresholds.449struct ggml_opencl_fa_kernels {450    // f16 Q / f16 KV451    std::map<std::pair<int, int>, cl_kernel> f16;452    std::map<std::pair<int, int>, cl_kernel> f16_q1;453    // f32 Q / f32 KV454    std::map<std::pair<int, int>, cl_kernel> f32;455    std::map<std::pair<int, int>, cl_kernel> f32_q1;456    // f32 Q / f16 KV (mixed)457    std::map<std::pair<int, int>, cl_kernel> f32_f16;458    std::map<std::pair<int, int>, cl_kernel> f32_f16_split;          // N_SPLIT>1 variant459    std::map<std::pair<int, int>, cl_kernel> f32_f16_split_k_img;    // DK=512 prefill split, K via image1d_buffer_t460    std::map<std::pair<int, int>, cl_kernel> f32_f16_q1;461    std::map<std::pair<int, int>, cl_kernel> f32_f16_q1_split;       // flash-decoding K-split462    // vec decode463    std::map<std::pair<int, int>, cl_kernel> f32_f16_q1_vec;464    // kv-head-coalesced vec decode465    std::map<std::pair<int, int>, cl_kernel> f32_f16_q1_vec_mq;466    // kv-head-coalesced + flash-decoding split467    std::map<std::pair<int, int>, cl_kernel> f32_f16_q1_vec_mq_split;468    // MQ_GQA=8 specializations469    std::map<std::pair<int, int>, cl_kernel> f32_f16_q1_vec_mq_g8;470    std::map<std::pair<int, int>, cl_kernel> f32_f16_q1_vec_mq_split_g8;471    // k-image variant of MQ_G8 vec_mq_split472    std::map<std::pair<int, int>, cl_kernel> f32_f16_q1_vec_mq_split_g8_k_img;473    // k-image variant of MQ_GQA=4 vec_mq_split474    std::map<std::pair<int, int>, cl_kernel> f32_f16_q1_vec_mq_split_k_img;475    // Cluster-parallel decode476    std::map<std::pair<int, int>, cl_kernel> f32_f16_q1_vec_mq_split_c8;477    std::map<std::pair<int, int>, cl_kernel> f32_f16_q1_vec_mq_split_g8_c8;478    // NSG_SPLIT=2 specializations (WG=128): the c8 kernel's register footprint479    // caps its per-kernel WG at 128 on X2, below the stock 256/192 requirement.480    // 2 subgroups × FA_CL_NCL streams still gives 16 in-flight rows per WG.481    std::map<std::pair<int, int>, cl_kernel> f32_f16_q1_vec_mq_split_c8_ns2;482    std::map<std::pair<int, int>, cl_kernel> f32_f16_q1_vec_mq_split_g8_c8_ns2;483    // FA_CL_C=32 / MQ_GQA=8 / NSG_SPLIT=2 specialization for the DK=DV=256484    // GQA=8 class (Qwen3.5/3.6-35B-A3B: 16 Q heads, 2 KV heads). o_acc =485    // DV_VEC/32 × 8 = 128B/lane (in budget); the baseline fa1 path for this486    // shape has NO MQ/FD at all and pays an 8× KV re-read per Q head.487    std::map<std::pair<int, int>, cl_kernel> f32_f16_q1_vec_mq_split_g8_c32;488    // alternative decode489    std::map<std::pair<int, int>, cl_kernel> f32_f16_q1_local_tile;490    // hybrid local-tile + MQ + FD-split kernel for DK=DV=128 only491    std::map<std::pair<int, int>, cl_kernel> f32_f16_q1_local_mq_split;492    std::map<std::pair<int, int>, cl_kernel> f32_f16_q1_local_mq_split_g8;493    std::map<std::pair<int, int>, int>       f32_f16_bm;494    std::map<std::pair<int, int>, int>       f32_f16_bn;495    std::map<std::pair<int, int>, int>       f32_f16_wg_size;496    std::map<std::pair<int, int>, int>       f32_f16_split_wg_size;497    std::map<std::pair<int, int>, int>       f32_f16_split_nkv_threshold;498    // f32 Q / native q8_0 KV499    std::map<std::pair<int, int>, cl_kernel> f32_q8_0_q1;            // decode500    std::map<std::pair<int, int>, cl_kernel> f32_q8_0_q1_vec;        // DV-split + multi-subgroup decode501    std::map<std::pair<int, int>, cl_kernel> f32_q8_0_q1_split;      // flash-decoding pass 1502    // KV-head-coalesced + flash-decoding split for q8_0 KV503    std::map<std::pair<int, int>, cl_kernel> f32_q8_0_q1_vec_mq_split;504    std::map<std::pair<int, int>, cl_kernel> f32_q8_0_q1_vec_mq_split_g8;505    // Cluster-parallel q8_0 decode506    std::map<std::pair<int, int>, cl_kernel> f32_q8_0_q1_vec_mq_split_c8;507    std::map<std::pair<int, int>, cl_kernel> f32_q8_0;               // prefill (baseline)508    std::map<std::pair<int, int>, cl_kernel> f32_q8_0_split;         // N_SPLIT>1 variant509    std::map<std::pair<int, int>, int>       f32_q8_0_split_wg_size;        // wg_size = bm*n_split510    std::map<std::pair<int, int>, int>       f32_q8_0_split_nkv_threshold;  // use split when n_kv >= this511    std::map<std::pair<int, int>, int>       f32_q8_0_split_bm;             // per-split BLOCK_M512    // f32 Q / native q4_0 KV513    std::map<std::pair<int, int>, cl_kernel> f32_q4_0_q1;514    std::map<std::pair<int, int>, cl_kernel> f32_q4_0_q1_vec;        // DV-split + multi-subgroup decode515    std::map<std::pair<int, int>, cl_kernel> f32_q4_0_q1_split;516    // kv-head-coalesced + flash-decoding split for q4_0 kv (dp4a K dot)517    std::map<std::pair<int, int>, cl_kernel> f32_q4_0_q1_vec_mq_split;518    std::map<std::pair<int, int>, cl_kernel> f32_q4_0_q1_vec_mq_split_g8;519    // Cluster-parallel q4_0 decode520    std::map<std::pair<int, int>, cl_kernel> f32_q4_0_q1_vec_mq_split_g8_c8;521    std::map<std::pair<int, int>, cl_kernel> f32_q4_0_q1_vec_mq_split_c8;522    std::map<std::pair<int, int>, cl_kernel> f32_q4_0;523    std::map<std::pair<int, int>, cl_kernel> f32_q4_0_split;524    std::map<std::pair<int, int>, int>       f32_q4_0_split_wg_size;525    std::map<std::pair<int, int>, int>       f32_q4_0_split_nkv_threshold;526    std::map<std::pair<int, int>, int>       f32_q4_0_split_bm;527    // shared: flash-decoding merge + prefill prepass (kv-pad, mask-pad, blk class)528    std::map<std::pair<int, int>, cl_kernel> f32_merge;529    std::map<std::pair<int, int>, cl_kernel> kv_pad_f16;530    std::map<std::pair<int, int>, cl_kernel> mask_pad_f16;531    std::map<std::pair<int, int>, cl_kernel> blk_f16;532    // generic prefill tile dims (f16 / f32 paths)533    std::map<std::pair<int, int>, int>       bm;534    std::map<std::pair<int, int>, int>       bn;535    // attempted (variant, (dk, dv))536    // all attempted FA kernels appear here, but those not registered failed compilation537    std::set<std::pair<int, std::pair<int, int>>> variant_attempted;538};539 540// backend context541struct ggml_backend_opencl_context {542    int ref_count;543 544    cl_device_id device;545    std::string device_name;546 547    ggml_cl_version platform_version;548    ggml_cl_version opencl_c_version;549 550    // argsort is loaded in supports_op because its availability depends on how551    // many workgroups are allowed, which requires kernel compilation.552    bool kernels_loaded_argsort = false;553    // rest of the kernels are currently always loaded in alloc_buffer.554    bool kernels_loaded = false;555 556    std::string driver_version;557 558    GPU_FAMILY gpu_family;559    ADRENO_GPU_GEN adreno_gen;560 561    cl_int alignment;562    size_t global_mem_size;563    size_t max_alloc_size;564    size_t max_workgroup_size;565    bool fp16_support;566    bool has_vector_subgroup_broadcast;567    bool has_subgroup_shuffle = false;       // cl_khr_subgroup_shuffle or cl_qcom_subgroup_shuffle568    bool has_integer_dot      = false;       // cl_khr_integer_dot_product or cl_qcom_dot_product8569    bool has_qcom_subgroup_shuffle = false;  // specifically cl_qcom_subgroup_shuffle570    bool disable_fusion;571 572    // ragged moe, use int to directly pass to kernel573    cl_uint  adreno_use_moe_ragged;574    cl_uint  adreno_moe_ragged_skip_gran;575    cl_uint  adreno_use_moe_ragged_dp4;576 577    // whether fuse moe combine578    cl_uint fuse_moe_combine;579 580    bool adreno_has_large_buffer;581    bool adreno_use_large_buffer;582    bool adreno_use_bin_kernels;583    get_adreno_bin_kernel_func_t get_adreno_bin_kernel_func = nullptr;584    ggml_cl_compiler_version adreno_cl_compiler_version;585 586    std::string kernel_compile_opts;  // cached for lazy-compiled kernels.587 588    int adreno_wave_size;589 590    cl_bool non_uniform_workgroups;591    size_t  image_max_buffer_size;592    size_t  image2d_max_width;593    size_t  image2d_max_height;594 595    cl_device_svm_capabilities svm_caps;596 597    cl_context context;598    cl_command_queue queue;599 600    // On-disk compiled-program cache (see GGML_OPENCL_KERNEL_CACHE_DIR).601    cl_program_cache_state program_cache;602    bool program_cache_initialized = false;603 604    // prealloc buffers for transposing weights and activations605    ggml_cl_buffer prealloc_quant_trans;606    ggml_cl_buffer prealloc_scales_trans;607    ggml_cl_buffer prealloc_act_trans;608    // q8_1-quantized reordered MoE activations for the dp4a prefill GEMM.609    ggml_cl_buffer prealloc_moe_qa;   // int8 quants  [tok_slots * ne00]610    ggml_cl_buffer prealloc_moe_da;   // per-block d  [tok_slots * ne00/32] (half)611    ggml_cl_buffer prealloc_moe_sa;   // per-block s  [tok_slots * ne00/32] (half)612    // scratch copy of the router weights to avoid dst aliasing613    ggml_cl_buffer prealloc_moe_combine_w;614 615    // pool of persistent image1d_buffer views over kv-cache layers, keyed by616    // (parent buffer, offset within parent)617    // used by the img-variant KQ/KQV dispatch paths to avoid per-call618    // clCreateSubBuffer + clCreateImage + pending-release-queue on long-context decode619    struct ImagePoolKey {620        uintptr_t buf;621        uint64_t  offset;622        bool operator<(const ImagePoolKey & o) const {623            if (buf != o.buf) return buf < o.buf;624            return offset < o.offset;625        }626    };627    struct ImagePoolEntry {628        cl_mem sub_buffer = nullptr;629        cl_mem image      = nullptr;630        size_t k_bytes    = 0;631        cl_channel_type channel_data_type = CL_FLOAT;632    };633    std::map<ImagePoolKey, ImagePoolEntry> kq_img_pool;634    std::map<ImagePoolKey, ImagePoolEntry> kqv_img_pool;635 636    // pool for the on-device f16 buffer for kv-cache with non-FA quantized-K (q8_0/q4_0)637    std::map<ImagePoolKey, ImagePoolEntry> dequant_f16_pool;638 639    // prealloc buffers for src0 and src1640    ggml_cl_buffer prealloc_src0;641    ggml_cl_buffer prealloc_src1;642 643#ifdef GGML_OPENCL_USE_ADRENO_KERNELS644    ggml_cl_buffer prealloc_adreno_xmem_const;645    bool adreno_xmem_gemm_enabled = false;646#endif647 648    // prealloc buffers for MoE router table preprocess649    bool toggle_reorder = false;650    ggml_cl_buffer prealloc_post_router;651    ggml_cl_buffer prealloc_emap;652    ggml_cl_buffer prealloc_hist;653    ggml_cl_buffer prealloc_tile_offset;654    ggml_cl_buffer prealloc_total_tiles;655    ggml_cl_buffer prealloc_slot_counter;656 657    cl_program program_add;658    cl_program program_add_id;659    cl_program program_clamp;660    cl_program program_cvt;661    cl_program program_diag_mask_inf;662    cl_program program_gelu;663    cl_program program_gemv_noshuffle_general;664    cl_program program_gemv_noshuffle;665    cl_program program_get_rows;666    cl_program program_set_rows;667    cl_program program_glu;668    cl_program program_im2col_f16;669    cl_program program_im2col_f32;670    cl_program program_mul_mat_Ab_Bi_8x4;671    cl_program program_mul_mv_q4_0_f32;672    cl_program program_mul_mv_q4_0_f32_v;673    cl_program program_mul_mv_q4_0_f32_8x_flat;674    cl_program program_mul_mv_q4_0_f32_1d_8x_flat;675    cl_program program_mul_mv_q4_0_f32_1d_16x_flat;676    cl_program program_mul_mv_q6_K;677    cl_program program_mul_mv_q8_0_f32, program_mul_mv_q8_0_f32_flat;678    cl_program program_mul_mv_mxfp4_f32;679    cl_program program_mul_mv_mxfp4_f32_flat;680    cl_program program_mul_mv_f16_f16;681    cl_program program_mul_mv_f16_f32_1row;682    cl_program program_mul_mv_f16_f32_l4;683    cl_program program_mul_mv_f16_f32;684    cl_program program_mul_mv_f32_f32;685    cl_program program_mul;686    cl_program program_mul_mat_f16_f32_tiled;687    cl_program program_mul_mm_f16_f32_kqv;688    cl_program program_mul_mm_f16_f32_kq;689    cl_program program_div;690    cl_program program_sub;691    cl_program program_norm;692    cl_program program_relu;693    cl_program program_rms_norm;694    cl_program program_group_norm;695    cl_program program_rope;696    cl_program program_silu;697    cl_program program_sigmoid;698    cl_program program_softmax_f32;699    cl_program program_softmax_f16;700    cl_program program_softmax_4_f32;701    cl_program program_softmax_4_f16;702    cl_program program_argsort_f32_i32;703    cl_program program_sum_rows_f32;704    cl_program program_pad;705    cl_program program_upscale;706    cl_program program_conv_2d_f16;707    cl_program program_conv_2d_f32;708    cl_program program_conv_2d_f16_f32;709    cl_program program_tsembd;710    cl_program program_gemv_moe_mxfp4_f32, program_gemm_moe_mxfp4_f32;711    cl_program program_mul_mv_id_q4_0_f32_8x_flat;712    cl_program program_mul_mv_id_q8_0_f32, program_mul_mv_id_q8_0_f32_flat;713    cl_program program_mul_mv_id_mxfp4_f32;714    cl_program program_mul_mv_id_mxfp4_f32_flat;715    cl_program program_mul_mm_f32_f32_l4_lm;716    cl_program program_mul_mm_f16_f32_l4_lm;717    cl_program program_mul_mm_q8_0_f32_l4_lm;718 719    cl_kernel kernel_add, kernel_add_row, kernel_add_f16, kernel_add_row_f16;720    cl_kernel kernel_mul, kernel_mul_row, kernel_mul_f16, kernel_mul_row_f16;721    cl_kernel kernel_div, kernel_div_row, kernel_div_f16, kernel_div_row_f16;722    cl_kernel kernel_sub, kernel_sub_row, kernel_sub_f16, kernel_sub_row_f16;723    cl_kernel kernel_add_id;724    cl_kernel kernel_scale_f32, kernel_scale_f32_4;725    cl_kernel kernel_sqr_cont_f32, kernel_sqr_cont_f32_4, kernel_sqr_cont_f16, kernel_sqr_cont_f16_4;726    cl_kernel kernel_sqrt_cont_f32, kernel_sqrt_cont_f32_4, kernel_sqrt_cont_f16, kernel_sqrt_cont_f16_4;727    cl_kernel kernel_mean_f32, kernel_mean_f32_4;728    cl_kernel kernel_silu, kernel_silu_4;729    cl_kernel kernel_gelu, kernel_gelu_4;730    cl_kernel kernel_gelu_erf, kernel_gelu_erf_4;731    cl_kernel kernel_gelu_quick, kernel_gelu_quick_4;732    cl_kernel kernel_relu;733    cl_kernel kernel_sigmoid_f32, kernel_sigmoid_f16;734    cl_kernel kernel_tri;735    cl_kernel kernel_fill;736    cl_kernel kernel_clamp;737    cl_kernel kernel_geglu, kernel_reglu, kernel_swiglu, kernel_swiglu_oai, kernel_geglu_erf, kernel_geglu_quick,738              kernel_geglu_f16, kernel_reglu_f16, kernel_swiglu_f16, kernel_geglu_erf_f16, kernel_geglu_quick_f16;739    cl_kernel kernel_norm, kernel_norm_mul_add;740    cl_kernel kernel_rms_norm, kernel_rms_norm_mul;741    cl_kernel kernel_l2_norm_f32;742    cl_kernel kernel_group_norm, kernel_group_norm_mul_add;743    cl_kernel kernel_diag_mask_inf, kernel_diag_mask_inf_8;744    cl_kernel kernel_diag_f32;745    cl_kernel kernel_soft_max, kernel_soft_max_4;746    cl_kernel kernel_soft_max_f16, kernel_soft_max_4_f16;747    ggml_opencl_fa_kernels fa;748    cl_kernel kernel_get_rows_f32, kernel_get_rows_f16, kernel_get_rows_q4_0;749    cl_kernel kernel_set_rows_f32_i64, kernel_set_rows_f32_i32, kernel_set_rows_f16_i64, kernel_set_rows_f16_i32;750    cl_kernel kernel_set_rows_q8_0_i64, kernel_set_rows_q8_0_i32;751    cl_kernel kernel_set_rows_q8_0_soa_i64, kernel_set_rows_q8_0_soa_i32;752    cl_kernel kernel_set_rows_q4_0_i64, kernel_set_rows_q4_0_i32;753    cl_kernel kernel_set_rows_q4_0_soa_i64, kernel_set_rows_q4_0_soa_i32;754    cl_kernel kernel_rope_norm_f32, kernel_rope_norm_f16, kernel_rope_neox_f32, kernel_rope_neox_f16;755    cl_kernel kernel_rope_multi_f32, kernel_rope_multi_f16, kernel_rope_vision_f32, kernel_rope_vision_f16;756    cl_kernel kernel_cpy_f16_f16, kernel_cpy_f16_f32, kernel_cpy_f32_f16, kernel_cpy_f32_f32, kernel_cpy_f32_f32_pack, kernel_cpy_i32_i32;757    cl_kernel kernel_mul_mat_f32_f32;758    cl_kernel kernel_mul_mat_f16_f16;759    cl_kernel kernel_mul_mat_f16_f32_1row;760    cl_kernel kernel_mul_mat_f16_f32;761    cl_kernel kernel_mul_mat_f16_f32_l4;762    cl_kernel kernel_mul_mat_f16_f32_l4_dr;763    cl_kernel kernel_mul_mat_f16_f32_l4_dr_ls;764    cl_kernel kernel_mul_mat_f16_f32_l4_dr_lq;765    cl_kernel kernel_mul_mat_f16_f32_l4_x8 = nullptr;766    cl_kernel kernel_mul_mat_f16_f32_l4_x8_pair = nullptr;767    cl_kernel kernel_mul_mat_f16_f32_l4_x8_gqa4 = nullptr;768    cl_kernel kernel_mul_mat_f16_f32_l4_x8_gqa4_img = nullptr;769    cl_kernel kernel_mul_mat_f16_f32_l4_x8_gqa_r4_img = nullptr;770    cl_kernel kernel_mul_mat_f16_f32_l4_x8_gqa_r2_dk256_img = nullptr;771    cl_kernel kernel_mul_mat_f16_f32_l4_y8 = nullptr;772    cl_kernel kernel_mul_mat_f16_f32_l4_y8_gqa = nullptr;773    cl_kernel kernel_mul_mat_f16_f32_l4_y8_gqa_img = nullptr;774    cl_kernel kernel_mul_mat_f16_f32_tiled;775    cl_kernel kernel_adreno_xmem_pack_src_f32;776    cl_kernel kernel_adreno_xmem_prepack_weight_f16;777    cl_kernel kernel_gemm_xmem_f16_f32_os8;778    cl_kernel kernel_adreno_xmem_store_dst_f32;779    cl_kernel kernel_mul_mm_f16_f32_kqv;780    cl_kernel kernel_mul_mm_f16_f32_kq;781    cl_kernel kernel_mul_mat_q4_0_f32, kernel_mul_mat_q4_0_f32_v;782    cl_kernel kernel_convert_block_q1_0, kernel_restore_block_q1_0;783    cl_kernel kernel_convert_block_q4_0, kernel_restore_block_q4_0;784    cl_kernel kernel_convert_block_q4_0_trans4_ns, kernel_restore_block_q4_0_trans4_ns;785    cl_kernel kernel_convert_block_q4_1, kernel_restore_block_q4_1;786    cl_kernel kernel_convert_block_q4_1_trans4_ns, kernel_restore_block_q4_1_trans4_ns;787    cl_kernel kernel_convert_block_q5_0, kernel_restore_block_q5_0;788    cl_kernel kernel_convert_block_q5_0_trans4_ns, kernel_restore_block_q5_0_trans4_ns;789    cl_kernel kernel_convert_block_q5_1, kernel_restore_block_q5_1;790    cl_kernel kernel_convert_block_q5_1_trans4_ns, kernel_restore_block_q5_1_trans4_ns;791    cl_kernel kernel_convert_block_q4_k_trans4_ns, kernel_restore_block_q4_k_trans4_ns;792    cl_kernel kernel_convert_block_q5_k_trans4_ns, kernel_restore_block_q5_k_trans4_ns;793    cl_kernel kernel_convert_block_q6_k_trans4_ns, kernel_restore_block_q6_k_trans4_ns;794    cl_kernel kernel_convert_block_mxfp4, kernel_convert_block_mxfp4_trans, kernel_restore_block_mxfp4, kernel_restore_block_mxfp4_trans;795    cl_kernel kernel_convert_block_mxfp4_trans4_ns, kernel_restore_block_mxfp4_trans4_ns;796    cl_kernel kernel_convert_block_q8_0, kernel_restore_block_q8_0, kernel_restore_block_q8_0_trans;797    cl_kernel kernel_dequant_q8_0_f16_view_aos;798    cl_kernel kernel_dequant_q8_0_f32_view_aos;799    cl_kernel kernel_dequant_q4_0_f16_view_aos;800    cl_kernel kernel_dequant_q4_0_f32_view_aos;801    cl_kernel kernel_convert_block_q6_K_noshuffle, kernel_restore_block_q6_K_noshuffle;802    cl_kernel kernel_convert_bf16_to_f16, kernel_convert_f16_to_bf16;803    cl_kernel kernel_mul_mat_q4_0_f32_8x_flat;804    cl_kernel kernel_convert_block_q4_0_noshuffle;805    cl_kernel kernel_restore_block_q4_0_noshuffle;806    cl_kernel kernel_convert_block_q4_1_noshuffle;807    cl_kernel kernel_restore_block_q4_1_noshuffle;808    cl_kernel kernel_convert_block_q5_0_noshuffle;809    cl_kernel kernel_restore_block_q5_0_noshuffle;810    cl_kernel kernel_convert_block_q5_1_noshuffle;811    cl_kernel kernel_restore_block_q5_1_noshuffle;812    cl_kernel kernel_convert_block_q4_K_noshuffle;813    cl_kernel kernel_restore_block_q4_K_noshuffle;814    cl_kernel kernel_convert_block_q4_K, kernel_restore_block_q4_K;815    cl_kernel kernel_convert_block_q5_K, kernel_restore_block_q5_K;816    cl_kernel kernel_convert_block_q5_K_noshuffle;817    cl_kernel kernel_restore_block_q5_K_noshuffle;818    cl_kernel kernel_convert_block_q6_K, kernel_restore_block_q6_K;819    cl_kernel kernel_convert_block_iq4_nl, kernel_restore_block_iq4_nl;820    cl_kernel kernel_convert_block_iq4_nl_noshuffle;821    cl_kernel kernel_restore_block_iq4_nl_noshuffle;822    cl_kernel kernel_mul_mv_q1_0_f32, kernel_mul_mv_q1_0_f32_flat;823    cl_kernel kernel_mul_mat_q4_0_f32_1d_8x_flat, kernel_mul_mat_q4_0_f32_1d_16x_flat;824    cl_kernel kernel_mul_mv_q4_1_f32;825    cl_kernel kernel_mul_mv_q4_1_f32_flat;826    cl_kernel kernel_mul_mv_q5_0_f32;827    cl_kernel kernel_mul_mv_q5_0_f32_flat;828    cl_kernel kernel_mul_mv_q5_1_f32;829    cl_kernel kernel_mul_mv_q5_1_f32_flat;830    cl_kernel kernel_mul_mv_q4_K_f32;831    cl_kernel kernel_mul_mv_q4_K_f32_flat;832    cl_kernel kernel_mul_mv_q5_K_f32;833    cl_kernel kernel_mul_mv_q5_K_f32_flat;834    cl_kernel kernel_mul_mv_q6_K_f32;835    cl_kernel kernel_mul_mv_q6_K_f32_flat;836    cl_kernel kernel_mul_mv_mxfp4_f32, kernel_mul_mv_mxfp4_f32_flat;837    cl_kernel kernel_mul_mv_q8_0_f32, kernel_mul_mv_q8_0_f32_flat;838    cl_kernel kernel_mul_mv_iq4_nl_f32;839    cl_kernel kernel_mul_mv_iq4_nl_f32_flat;840    cl_kernel kernel_solve_tri_f32;841    cl_kernel kernel_im2col_f32, kernel_im2col_f16;842    cl_kernel kernel_argsort_f32_i32;843    cl_kernel kernel_sum_rows_f32, kernel_sum_rows_f32_4;844    cl_kernel kernel_cumsum_blk, kernel_cumsum_add;845    cl_kernel kernel_repeat_f32;846    cl_kernel kernel_pad;847    cl_kernel kernel_tanh_f32, kernel_tanh_f32_4, kernel_tanh_f32_nc;848    cl_kernel kernel_tanh_f16, kernel_tanh_f16_4, kernel_tanh_f16_nc;849    cl_kernel kernel_neg_f32, kernel_neg_f32_4, kernel_neg_f32_nc;850    cl_kernel kernel_neg_f16, kernel_neg_f16_4, kernel_neg_f16_nc;851    cl_kernel kernel_exp_f32, kernel_exp_f32_4, kernel_exp_f32_nc;852    cl_kernel kernel_exp_f16, kernel_exp_f16_4, kernel_exp_f16_nc;853    cl_kernel kernel_expm1_f32, kernel_expm1_f32_4, kernel_expm1_f32_nc;854    cl_kernel kernel_expm1_f16, kernel_expm1_f16_4, kernel_expm1_f16_nc;855    cl_kernel kernel_abs_f32, kernel_abs_f32_4, kernel_abs_f32_nc;856    cl_kernel kernel_abs_f16, kernel_abs_f16_4, kernel_abs_f16_nc;857    cl_kernel kernel_softplus_f32, kernel_softplus_f32_4, kernel_softplus_f32_nc;858    cl_kernel kernel_softplus_f16, kernel_softplus_f16_4, kernel_softplus_f16_nc;859    cl_kernel kernel_upscale;860    cl_kernel kernel_upscale_bilinear;861    cl_kernel kernel_concat_f32, kernel_concat_f32_pack;862    cl_kernel kernel_conv_2d_f16;863    cl_kernel kernel_conv_2d_f32;864    cl_kernel kernel_conv_2d_f16_f32;865    cl_kernel kernel_ssm_conv_f32_f32, kernel_ssm_conv_f32_f32_4;866    // [size_idx][kda][tgpp] where size_idx: 0=S_V=16, 1=32, 2=64, 3=128; kda: 0 or 1.867    // tgpp 0 = TG variant (COLS_PER_LANE_GROUP=1), tgpp 1 = prefill variant (COLS_PER_LANE_GROUP=4).868    cl_kernel kernel_gated_delta_net_f32[4][2][2] = {};869    cl_kernel kernel_timestep_embedding;870    cl_kernel kernel_gemv_moe_q4_0_f32_ns, kernel_gemm_moe_q4_0_f32_ns, kernel_gemm_moe_q4_0_f32_ns_bin;871    cl_kernel kernel_gemm_moe_q8_0_f32_ns;872    cl_kernel kernel_gemv_moe_q4_1_f32_ns, kernel_gemm_moe_q4_1_f32_ns, kernel_gemm_moe_q4_1_f32_ns_bin;873    cl_kernel kernel_gemv_moe_q5_0_f32_ns, kernel_gemm_moe_q5_0_f32_ns;874    cl_kernel kernel_gemv_moe_q5_1_f32_ns, kernel_gemm_moe_q5_1_f32_ns;875    cl_kernel kernel_gemv_moe_q4_k_f32_ns, kernel_gemm_moe_q4_k_f32_ns, kernel_gemm_moe_q4_k_f32_ns_bin;876    cl_kernel kernel_gemv_moe_q4_k_f32_ns_wimg = nullptr;  // weight-as-texture MoE decode GEMV (opt-in)877    cl_kernel kernel_gemm_moe_q4_k_q8_1_dp4a = nullptr;    // dp4a (int8) prefill GEMM variant878    cl_kernel kernel_moe_reorder_quant_a_q8_1;   // fused reorder + q8_1 quant for the dp4a GEMM879    cl_kernel kernel_gemm_moe_q8_1_dp4a_q80 = nullptr;   // generic dp4a MoE GEMM (MOE_QT=80), opt-in880    cl_kernel kernel_moe_expand_scale_q8_0 = nullptr;    // q8_0 per-block d -> uniform scale[16]881    cl_kernel kernel_gemm_moe_q8_1_dp4a_q50 = nullptr;   // generic dp4a MoE GEMM (MOE_QT=50, q5_0), opt-in882    cl_kernel kernel_moe_expand_scale_q5_0 = nullptr;    // q5_0 d -> uniform scale[2]/min[1] per 32-block883    cl_kernel kernel_gemm_moe_q8_1_dp4a_q5k = nullptr;   // generic dp4a MoE GEMM (MOE_QT=5, q5_K), opt-in884    cl_kernel kernel_moe_expand_scale_q5_K = nullptr;    // q5_K 6-bit s[] -> uniform scale[16]/min[8]885    cl_kernel kernel_gemv_moe_q5_k_f32_ns, kernel_gemm_moe_q5_k_f32_ns;886    cl_kernel kernel_gemv_moe_q6_k_f32_ns, kernel_gemm_moe_q6_k_f32_ns, kernel_gemm_moe_q6_k_f32_ns_bin;887    cl_kernel kernel_gemm_moe_q6_k_q8_1_dp4a = nullptr;    // dp4a (int8) q6_K MoE prefill GEMM888    cl_kernel kernel_gemv_moe_mxfp4_f32, kernel_gemm_moe_mxfp4_f32;889    cl_kernel kernel_gemv_moe_mxfp4_f32_ns, kernel_gemm_moe_mxfp4_f32_ns, kernel_gemm_moe_mxfp4_f32_ns_bin;890    cl_kernel kernel_gemv_moe_mxfp4_f32_ns_wimg = nullptr;      // weight-as-texture MoE decode GEMV891    cl_kernel kernel_gemm_moe_mxfp4_q8_1_dp4a = nullptr;   // dp4a (int8) mxfp4 MoE prefill GEMM892    cl_kernel kernel_gemm_moe_q4_0_q8_1_dp4a = nullptr;    // dp4a (int8) q4_0 MoE prefill GEMM893    cl_kernel kernel_moe_reorder_b;894    cl_kernel kernel_moe_histogram, kernel_moe_scan, kernel_moe_fill, kernel_moe_scatter;895    cl_kernel kernel_moe_combine_f32 = nullptr;   // fused router-weight mul + cross-expert sum896    cl_kernel kernel_mul_mv_id_q4_0_f32_8x_flat;897    cl_kernel kernel_mul_mv_id_q8_0_f32, kernel_mul_mv_id_q8_0_f32_flat;898    cl_kernel kernel_mul_mv_id_mxfp4_f32;899    cl_kernel kernel_mul_mv_id_mxfp4_f32_flat;900    cl_kernel kernel_mul_mm_f32_f32_l4_lm;901    cl_kernel kernel_mul_mm_f16_f32_l4_lm;902    cl_kernel kernel_mul_mm_q1_0_f32_l4_lm;903    cl_kernel kernel_mul_mm_q4_0_f32_l4_lm;904    cl_kernel kernel_mul_mm_q4_1_f32_l4_lm;905    cl_kernel kernel_mul_mm_q5_0_f32_l4_lm;906    cl_kernel kernel_mul_mm_q5_1_f32_l4_lm;907    cl_kernel kernel_mul_mm_q8_0_f32_l4_lm;908    cl_kernel kernel_mul_mm_q4_k_f32_l4_lm;909    cl_kernel kernel_mul_mm_q5_k_f32_l4_lm;910    cl_kernel kernel_mul_mm_q6_k_f32_l4_lm;911    cl_kernel kernel_mul_mm_iq4_nl_f32_l4_lm;912 913    std::vector<ProfilingInfo> profiling_info;914    std::vector<ProfilingInfo> profiling_results;915 916    void flush_profiling_batch() {917        if (profiling_info.empty()) {918            return;919        }920 921        // Populate profiling info922        for (ProfilingInfo & info : profiling_info) {923            cl_ulong cmd_queued;924            cl_ulong cmd_submit;925            cl_ulong cmd_start;926            cl_ulong cmd_end;927            cl_ulong cmd_complete;928 929            CL_CHECK(clWaitForEvents(1, &info.evt));930            CL_CHECK(clGetEventProfilingInfo(931                info.evt, CL_PROFILING_COMMAND_QUEUED, sizeof(cl_ulong), &cmd_queued, NULL));932            CL_CHECK(clGetEventProfilingInfo(933                info.evt, CL_PROFILING_COMMAND_SUBMIT, sizeof(cl_ulong), &cmd_submit, NULL));934            CL_CHECK(clGetEventProfilingInfo(935                info.evt, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &cmd_start, NULL));936            CL_CHECK(clGetEventProfilingInfo(937                info.evt, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &cmd_end, NULL));938            CL_CHECK(clGetEventProfilingInfo(939                info.evt, CL_PROFILING_COMMAND_COMPLETE, sizeof(cl_ulong), &cmd_complete, NULL));940            CL_CHECK(clReleaseEvent(info.evt));941            info.evt = nullptr;942 943            char kernel_name[512];944            CL_CHECK(clGetKernelInfo(info.kernel, CL_KERNEL_FUNCTION_NAME,945                sizeof(kernel_name), kernel_name, NULL));946            info.kernel_name = kernel_name;947 948            info.cmd_queued = cmd_queued;949            info.cmd_submit = cmd_submit;950            info.cmd_start  = cmd_start;951            info.cmd_end    = cmd_end;952 953            info.cmd_queued_duration_ns     = cmd_submit    - cmd_queued;954            info.cmd_submit_duration_ns     = cmd_start     - cmd_submit;955            info.cmd_duration_ns            = cmd_end       - cmd_start;956            info.cmd_complete_duration_ns   = cmd_complete  - cmd_end;957            info.cmd_total_duration_ns      = cmd_complete  - cmd_queued;958        }959        profiling_results.insert(profiling_results.end(),960            std::make_move_iterator(profiling_info.begin()),961            std::make_move_iterator(profiling_info.end()));962        profiling_info.clear();963    }964 965    void write_profiling_info() {966        if (profiling_results.empty()) {967            return;968        }969 970        // Dump a csv971        FILE * fperf = fopen("cl_profiling.csv", "w");972        if (!fperf) {973            GGML_LOG_ERROR("Failed to open cl_profiling.csv\n");974            return;975        }976 977        fprintf(fperf, "op name, kernel name, exec duration (ms), global size, local size, output size\n");978        for (const ProfilingInfo & info : profiling_results) {979            fprintf(fperf, "%s,%s,%f,%zux%zux%zu,%zux%zux%zu,%zux%zux%zux%zu\n",980                info.op_name.c_str(), info.kernel_name.c_str(),981                info.cmd_duration_ns/1.e6f,982                info.global_size[0], info.global_size[1], info.global_size[2],983                info.local_size[0], info.local_size[1], info.local_size[2],984                info.output_size[0], info.output_size[1], info.output_size[2], info.output_size[3]);985        }986        fclose(fperf);987 988        // Dump a simple chrome trace989        FILE * ftrace = fopen("cl_trace.json", "w");990        if (!ftrace) {991            GGML_LOG_ERROR("Failed to open cl_trace.json\n");992            return;993        }994 995        fprintf(ftrace, "[\n");996        for (const ProfilingInfo & info : profiling_results) {997            fprintf(ftrace, "{\"name\": \"%s\", \"cat\": \"OpenCL\", \"ph\": \"B\", \"ts\": %" PRIu64 ", \"pid\": \"\", \"tid\": \"Host\"},\n",998                info.kernel_name.c_str(), info.cmd_queued/1000);999            fprintf(ftrace, "{\"name\": \"%s\", \"cat\": \"OpenCL\", \"ph\": \"E\", \"ts\": %" PRIu64 ", \"pid\": \"\", \"tid\": \"Host\"},\n",1000                info.kernel_name.c_str(), info.cmd_submit/1000);1001 1002            fprintf(ftrace, "{\"name\": \"%s\", \"cat\": \"OpenCL\", \"ph\": \"B\", \"ts\": %" PRIu64 ", \"pid\": \"\", \"tid\": \"Device\"},\n",1003                info.kernel_name.c_str(), info.cmd_start/1000);1004            fprintf(ftrace, "{\"name\": \"%s\", \"cat\": \"OpenCL\", \"ph\": \"E\", \"ts\": %" PRIu64 ", \"pid\": \"\", \"tid\": \"Device\"},\n",1005                info.kernel_name.c_str(), info.cmd_end/1000);1006        }1007        fprintf(ftrace, "]\n");1008        fclose(ftrace);1009    }1010 1011    size_t get_kernel_workgroup_size(cl_kernel kernel) const {1012        size_t workgroup_size = 0;1013        size_t ret_size = 0;1014        CL_CHECK(1015            clGetKernelWorkGroupInfo(kernel, device, CL_KERNEL_WORK_GROUP_SIZE,1016                sizeof(size_t), &workgroup_size, &ret_size));1017        GGML_ASSERT(sizeof(size_t) == ret_size);1018        return workgroup_size;1019    }1020 1021    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) {1022#ifdef GGML_OPENCL_PROFILING1023        cl_event evt;1024        CL_CHECK(clEnqueueNDRangeKernel(queue, kernel, work_dim, NULL, global_work_size, local_work_size, 0, NULL, &evt));1025 1026        profiling_info.emplace_back();1027        populateProfilingInfo(profiling_info.back(), evt, kernel, work_dim, global_work_size, local_work_size, tensor);1028        if (profiling_info.size() >= 2048) {1029            flush_profiling_batch();1030        }1031#else1032        GGML_UNUSED(tensor);1033        CL_CHECK(clEnqueueNDRangeKernel(queue, kernel, work_dim, NULL, global_work_size, local_work_size, 0, NULL, NULL));1034#endif1035    }1036 1037    const void * get_adreno_bin_kernel(const std::string &kernel_name, size_t *bin_size) const {1038        if (!get_adreno_bin_kernel_func) {1039            return nullptr;1040        }1041 1042        size_t sz;1043        const void * kernel_bin = get_adreno_bin_kernel_func(1044            kernel_name.c_str(), device_name.c_str(), driver_version.c_str(), &sz);1045        if (bin_size) {1046            *bin_size = sz;1047        }1048        return kernel_bin;1049    }1050 1051#ifdef GGML_OPENCL_USE_ADRENO_KERNELS1052    // Transpose kernels1053    cl_program program_transpose;1054 1055    cl_kernel kernel_transpose_32;1056    cl_kernel kernel_transpose_32_16;1057    cl_kernel kernel_transpose_16;1058    cl_kernel kernel_transpose_8_buf;1059    cl_kernel kernel_transpose_16_buf;1060    cl_kernel kernel_transpose_32_buf;1061    cl_kernel kernel_transpose_16_4x1;1062 1063    // Gemm and Gemv related programs, kernels, etc1064    cl_kernel kernel_gemm_noshuffle_q4_0_f32;1065    cl_kernel kernel_gemv_noshuffle_q4_0_f32;1066    cl_kernel kernel_gemv_noshuffle_q4_0_f32_4096_1_11008;1067    cl_kernel kernel_gemv_noshuffle_q4_0_f32_4096_1_4096;1068    cl_kernel kernel_gemv_noshuffle_q4_0_f32_11008_1_4096;1069    cl_kernel kernel_gemv_noshuffle_q4_0_f32_32000_1_4096;1070    cl_kernel kernel_gemv_noshuffle_q4_1_f32;1071    cl_kernel kernel_gemm_noshuffle_q4_1_f32;1072    cl_kernel kernel_gemm_noshuffle_q8_0_f32, kernel_gemm_noshuffle_q8_0_f32_bin;1073    cl_kernel kernel_gemm_noshuffle_q8_0_q8_1_dp4a = nullptr;  // dp4a (int8) dense q8_0 prefill GEMM (opt-in)1074    cl_kernel kernel_gemm_noshuffle_q8_0_q8_1_dp4a_wimg = nullptr;  // q8_0 dense dp4a, weights via texture (opt-in)1075    cl_kernel kernel_gemv_noshuffle_q8_0_f32;1076    cl_kernel kernel_gemm_noshuffle_q1_0_f32;1077    cl_kernel kernel_gemv_noshuffle_q1_0_f32;1078    cl_kernel kernel_gemv_noshuffle_q4_k_f32;1079    cl_kernel kernel_gemm_noshuffle_q4_k_f32;1080    cl_kernel kernel_gemm_noshuffle_q4_k_q8_1_dp4a = nullptr;  // dp4a (int8) dense prefill GEMM1081    cl_kernel kernel_gemm_noshuffle_q4_k_q8_1_dp4a_wimg = nullptr;  // dp4a dense prefill GEMM, weights via texture (X1 opt-in)1082    cl_kernel kernel_gemm_noshuffle_q5_k_q8_1_dp4a = nullptr;  // dp4a (int8) dense q5_K prefill GEMM1083    cl_kernel kernel_gemm_noshuffle_q6_k_q8_1_dp4a = nullptr;  // dp4a (int8) dense q6_K prefill GEMM1084    cl_kernel kernel_quant_a_q8_1;                    // plain activation q8_1 pre-pass1085    cl_kernel kernel_gemv_noshuffle_q6_K_f32;1086    cl_kernel kernel_gemm_noshuffle_q6_K_f32;1087    cl_kernel kernel_gemv_noshuffle_q5_k_f32;1088    cl_kernel kernel_gemm_noshuffle_q5_k_f32;1089    cl_kernel kernel_gemv_noshuffle_q5_0_f32;1090    cl_kernel kernel_gemm_noshuffle_q5_0_f32;1091    cl_kernel kernel_gemm_noshuffle_q5_0_q8_1_dp4a = nullptr;  // dp4a (int8) dense q5_0 prefill GEMM1092    cl_kernel kernel_gemm_noshuffle_q5_0_q8_1_dp4a_wimg = nullptr;  // q5_0 dense dp4a, qs plane via texture (opt-in)1093    cl_kernel kernel_gemv_noshuffle_q5_1_f32;1094    cl_kernel kernel_gemm_noshuffle_q5_1_f32;1095    cl_kernel kernel_gemv_noshuffle_iq4_nl_f32;1096    cl_kernel kernel_gemm_noshuffle_iq4_nl_f32;1097    cl_kernel kernel_gemm_noshuffle_iq4_nl_q8_1_dp4a = nullptr;  // dp4a (int8) dense IQ4_NL prefill GEMM1098    cl_kernel kernel_gemm_noshuffle_q4_0_q8_1_dp4a = nullptr;  // dp4a (int8) dense q4_0 prefill GEMM1099#endif // GGML_OPENCL_USE_ADRENO_KERNELS1100 1101    void free() {1102        clFinish(queue);1103 1104        ref_count--;1105        if (ref_count == 0) {1106#ifdef GGML_OPENCL_PROFILING1107            flush_profiling_batch();1108            write_profiling_info();1109            profiling_results.clear();1110#endif1111            // release pooled image1d_buffer views over KV cache layers.1112            for (auto & kv : kq_img_pool) {1113                if (kv.second.image)      { CL_CHECK(clReleaseMemObject(kv.second.image)); }1114                if (kv.second.sub_buffer) { CL_CHECK(clReleaseMemObject(kv.second.sub_buffer)); }1115            }1116            kq_img_pool.clear();1117            for (auto & kv : kqv_img_pool) {1118                if (kv.second.image)      { CL_CHECK(clReleaseMemObject(kv.second.image)); }1119                if (kv.second.sub_buffer) { CL_CHECK(clReleaseMemObject(kv.second.sub_buffer)); }1120            }1121            kqv_img_pool.clear();1122            for (auto & kv : dequant_f16_pool) {1123                if (kv.second.image) { CL_CHECK(clReleaseMemObject(kv.second.image)); }1124            }1125            dequant_f16_pool.clear();1126        }1127    }1128};1129 1130// All registered devices with a default device in the front.1131static std::vector<ggml_backend_device> g_ggml_backend_opencl_devices;1132// All device contexts associated with the devices above.1133// The devices live as long as the process, so do the contexts.1134static std::vector<std::unique_ptr<ggml_backend_opencl_device_context>> g_ggml_backend_opencl_dev_ctxs;1135 1136inline std::string read_file(const std::string &path) {1137    std::ifstream ifs(path);1138    if (!ifs) {1139        return "";1140    }1141    std::string text;1142    ifs.seekg(0, std::ios::end);1143    text.resize(ifs.tellg());1144    ifs.seekg(0, std::ios::beg);1145    ifs.read(&text[0], text.size());1146    return text;1147}1148 1149// fatal=false returns NULL on compile failure instead of aborting; used for1150// optional FA variants that may exhaust the Adreno compiler at large DK.1151// when the compiler returns CL_OUT_OF_HOST_MEMORY/CL_OUT_OF_RESOURCES (seen with DK>=256/512)1152// for FA programs, do clFinish the queue to free up resources, then rebuild (up to 3x)1153// if retry_queue is provided1154static cl_program build_program_from_source_ex(cl_context ctx, cl_device_id dev, const char* program_buffer, const std::string &compile_opts, bool fatal, const char *tag = nullptr, cl_command_queue retry_queue = nullptr) {1155    if (tag) { GGML_LOG_INFO("ggml_opencl: compiling %s\n", tag); }1156    cl_program p;1157    char *program_log;1158    size_t program_size;1159    size_t log_size;1160    int err;1161 1162    program_size = strlen(program_buffer);1163 1164    const int max_attempts = retry_queue ? 3 : 1;1165    for (int attempt = 0; attempt < max_attempts; ++attempt) {1166        p = clCreateProgramWithSource(ctx, 1, (const char**)&program_buffer, &program_size, &err);1167        if(err < 0) {1168            GGML_LOG_ERROR("OpenCL error creating program");1169            if (fatal) exit(1);1170            return NULL;1171        }1172 1173        err = clBuildProgram(p, 0, NULL, compile_opts.c_str(), NULL, NULL);1174        if (err == CL_SUCCESS) {1175            return p;1176        }1177 1178        const bool transient = (err == CL_OUT_OF_HOST_MEMORY || err == CL_OUT_OF_RESOURCES);1179        if (retry_queue && transient && attempt + 1 < max_attempts) {1180            clReleaseProgram(p);1181            GGML_LOG_WARN("ggml_opencl: transient compile failure (err=%d)%s%s — clFinish + retry (%d/%d)\n",1182                err, tag ? " building " : "", tag ? tag : "", attempt + 2, max_attempts);1183            clFinish(retry_queue);  // drain in-flight ops holding driver host-heap1184            continue;1185        }1186 1187        clGetProgramBuildInfo(p, dev, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);1188        program_log = (char*) malloc(log_size + 1);1189        program_log[log_size] = '\0';1190        clGetProgramBuildInfo(p, dev, CL_PROGRAM_BUILD_LOG, log_size + 1, program_log, NULL);1191        GGML_LOG_ERROR("ggml_opencl: kernel compile error (err=%d)%s%s:\n\n%s\n", err, tag ? " building " : "", tag ? tag : "", program_log);1192        free(program_log);1193        clReleaseProgram(p);1194        if (fatal) {1195            exit(1);1196        }1197        return nullptr;1198    }1199    return NULL;1200}

Showing the first 1,200 of 24840 lines. Download the file for the rest.