// repro_binary_constant_arg.cpp
//
// Minimal reproducer for an IMG DDK 24.2 / PowerVR A-Series AXE-1-16M
// regression on the TI SK-AM62B-P1 (libPVROCL.so.24.2.6643903):
//
//   Kernels created via `clCreateProgramWithBinary` lose the per-arg
//   metadata for arguments in the `__constant` address space, *when the
//   bound buffer is larger than ~4 KB*. `clSetKernelArg` then SIGSEGVs
//   inside libPVROCL.so the moment it binds such a buffer. The same
//   buffer binds cleanly to the source-built sibling of the same kernel.
//
// We narrowed the threshold:
//   - `__constant` buffer <= 4035 bytes: bind succeeds on the binary
//     kernel (the driver appears to take a small-constant inline path
//     that does not consult the missing metadata).
//   - `__constant` buffer >= 4036 bytes: bind segfaults on the binary
//     kernel. 4036 is 60 bytes below the 4 KB page; we suspect the
//     driver reserves a small metadata footer per page in the path that
//     reads the lost address-space attributes.
//
// The kernel below is the minimum we found: one __global float* output,
// one __constant float* LUT, body reads the LUT and writes the buffer.
// No __local memory, no barriers, no `max_constant_size` attribute, no
// `reqd_work_group_size`. Removing the body so the kernel makes no use
// of `lut` also reproduces the crash — only the parameter declaration
// matters.
//
// Build:  g++ -std=c++17 repro_binary_constant_arg.cpp -lOpenCL -o repro
// Run:    ./repro

#define CL_TARGET_OPENCL_VERSION 300
#include <CL/cl.h>

#include <cstdio>
#include <cstdlib>
#include <vector>

static const char* kSource = R"CL(
__kernel void use_constant(__global float* out, __constant float* lut) {
    uint gid = get_global_id(0);
    out[gid] = lut[gid & 15u];
}
)CL";

#define CHECK(e) do { cl_int _e = (e); \
    if (_e != CL_SUCCESS) { std::fprintf(stderr, \
        "%s:%d: OpenCL error %d\n", __FILE__, __LINE__, _e); std::exit(1); } \
} while (0)

int main() {
    cl_platform_id plat = nullptr;
    CHECK(clGetPlatformIDs(1, &plat, nullptr));
    cl_device_id dev = nullptr;
    CHECK(clGetDeviceIDs(plat, CL_DEVICE_TYPE_GPU, 1, &dev, nullptr));
    char name[256] = {};
    clGetDeviceInfo(dev, CL_DEVICE_NAME, sizeof(name), name, nullptr);
    std::printf("device: %s\n\n", name);

    cl_int err = CL_SUCCESS;
    cl_context ctx = clCreateContext(nullptr, 1, &dev, nullptr, nullptr, &err);
    CHECK(err);

    // 1) Source build, capture binary, reload it via clCreateProgramWithBinary.
    cl_program src_prog = clCreateProgramWithSource(
        ctx, 1, &kSource, nullptr, &err); CHECK(err);
    CHECK(clBuildProgram(src_prog, 1, &dev, nullptr, nullptr, nullptr));
    size_t bin_size = 0;
    CHECK(clGetProgramInfo(src_prog, CL_PROGRAM_BINARY_SIZES,
                           sizeof(bin_size), &bin_size, nullptr));
    std::vector<unsigned char> binary(bin_size);
    unsigned char* binary_ptr = binary.data();
    CHECK(clGetProgramInfo(src_prog, CL_PROGRAM_BINARIES,
                           sizeof(unsigned char*), &binary_ptr, nullptr));
    cl_int bin_status = 0, bin_err = 0;
    const unsigned char* bin_const = binary.data();
    cl_program bin_prog = clCreateProgramWithBinary(
        ctx, 1, &dev, &bin_size, &bin_const, &bin_status, &bin_err);
    CHECK(bin_err);
    CHECK(clBuildProgram(bin_prog, 1, &dev, nullptr, nullptr, nullptr));
    std::printf("Step 1: source build + binary round-trip OK\n"
                "        captured %zu-byte binary, link returned CL_SUCCESS\n\n",
                bin_size);

    // 2) Create the kernel from both programs and a single __global buffer.
    cl_kernel k_src = clCreateKernel(src_prog, "use_constant", &err); CHECK(err);
    cl_kernel k_bin = clCreateKernel(bin_prog, "use_constant", &err); CHECK(err);
    const size_t N = 1024;
    cl_mem out_buf = clCreateBuffer(ctx, CL_MEM_READ_WRITE,
                                    sizeof(float) * N, nullptr, &err); CHECK(err);
    CHECK(clSetKernelArg(k_src, 0, sizeof(cl_mem), &out_buf));
    CHECK(clSetKernelArg(k_bin, 0, sizeof(cl_mem), &out_buf));

    // 3) Bind a SMALL (1 KB) __constant buffer to both kernels — both succeed.
    {
        cl_mem lut_small = clCreateBuffer(ctx, CL_MEM_READ_ONLY,
                                          1024, nullptr, &err); CHECK(err);
        std::printf("Step 2: bind 1024-byte __constant buffer\n");
        cl_int e1 = clSetKernelArg(k_src, 1, sizeof(cl_mem), &lut_small);
        std::printf("        SOURCE-built kernel  arg 1 err=%d\n", (int)e1);
        cl_int e2 = clSetKernelArg(k_bin, 1, sizeof(cl_mem), &lut_small);
        std::printf("        BINARY-loaded kernel arg 1 err=%d\n\n", (int)e2);
        clReleaseMemObject(lut_small);
    }

    // 4) Bind a LARGER (8 KB) __constant buffer:
    //      - SOURCE-built kernel: succeeds
    //      - BINARY-loaded kernel: SIGSEGV inside libPVROCL.so
    //    The threshold is around 4 KB. We use 8 KB here to be unambiguous;
    //    the bug also fires at anything >= 4036 bytes on this driver.
    {
        cl_mem lut_big = clCreateBuffer(ctx, CL_MEM_READ_ONLY,
                                        8192, nullptr, &err); CHECK(err);
        std::printf("Step 3: bind 8192-byte __constant buffer\n");
        cl_int e1 = clSetKernelArg(k_src, 1, sizeof(cl_mem), &lut_big);
        std::printf("        SOURCE-built kernel  arg 1 err=%d\n", (int)e1);
        std::printf("        BINARY-loaded kernel arg 1 ... "); std::fflush(stdout);
        cl_int e2 = clSetKernelArg(k_bin, 1, sizeof(cl_mem), &lut_big);
        // On the affected driver the previous call SIGSEGVs and this
        // line never prints.
        std::printf("err=%d\n", (int)e2);
        clReleaseMemObject(lut_big);
    }

    clReleaseMemObject(out_buf);
    clReleaseKernel(k_src);
    clReleaseKernel(k_bin);
    clReleaseProgram(src_prog);
    clReleaseProgram(bin_prog);
    clReleaseContext(ctx);
    return 0;
}