static int load_symbol(void *lib, const char *name, void **out) {
    *out = dlsym(lib, name);
    return *out != NULL;
}

static int load_cuda(struct cuda_api *api) {
    memset(api, 0, sizeof(*api));
    api->lib = dlopen("libcuda.so.1", RTLD_NOW | RTLD_LOCAL);
    if (!api->lib) {
        return 0;
    }
    if (!(
        load_symbol(api->lib, "cuInit", (void **)&api->cuInit) &&
        load_symbol(api->lib, "cuDeviceGetCount", (void **)&api->cuDeviceGetCount) &&
        load_symbol(api->lib, "cuDeviceGet", (void **)&api->cuDeviceGet) &&
        load_symbol(api->lib, "cuDeviceGetName", (void **)&api->cuDeviceGetName) &&
        load_symbol(api->lib, "cuDeviceGetAttribute", (void **)&api->cuDeviceGetAttribute) &&
        load_symbol(api->lib, "cuCtxCreate_v2", (void **)&api->cuCtxCreate) &&
        load_symbol(api->lib, "cuCtxDestroy_v2", (void **)&api->cuCtxDestroy) &&
        load_symbol(api->lib, "cuCtxSynchronize", (void **)&api->cuCtxSynchronize) &&
        load_symbol(api->lib, "cuMemAlloc_v2", (void **)&api->cuMemAlloc) &&
        load_symbol(api->lib, "cuMemFree_v2", (void **)&api->cuMemFree) &&
        load_symbol(api->lib, "cuMemsetD8_v2", (void **)&api->cuMemsetD8) &&
        load_symbol(api->lib, "cuMemcpyHtoD_v2", (void **)&api->cuMemcpyHtoD) &&
        load_symbol(api->lib, "cuMemcpyDtoH_v2", (void **)&api->cuMemcpyDtoH) &&
        load_symbol(api->lib, "cuModuleLoadDataEx", (void **)&api->cuModuleLoadDataEx) &&
        load_symbol(api->lib, "cuModuleGetFunction", (void **)&api->cuModuleGetFunction) &&
        load_symbol(api->lib, "cuLaunchKernel", (void **)&api->cuLaunchKernel))) {
        dlclose(api->lib);
        memset(api, 0, sizeof(*api));
        return 0;
    }
    load_symbol(api->lib, "cuMemGetInfo_v2", (void **)&api->cuMemGetInfo);
    load_symbol(api->lib, "cuStreamCreate", (void **)&api->cuStreamCreate);
    if (!load_symbol(api->lib, "cuStreamDestroy_v2", (void **)&api->cuStreamDestroy)) {
        load_symbol(api->lib, "cuStreamDestroy", (void **)&api->cuStreamDestroy);
    }
    return 1;
}

static const char *cu_error_name(struct cuda_api *api, CUresult rc) {
    const char *value = NULL;
    if (api->cuGetErrorName && api->cuGetErrorName(rc, &value) == CU_SUCCESS && value) {
        return value;
    }
    return "CUDA_ERROR";
}

static const char *cu_error_string(struct cuda_api *api, CUresult rc) {
    const char *value = NULL;
    if (api->cuGetErrorString && api->cuGetErrorString(rc, &value) == CU_SUCCESS && value) {
        return value;
    }
    return "unknown";
}

static int check_rc(struct cuda_api *api, const char *step, CUresult rc) {
    if (rc == CU_SUCCESS) {
        return 1;
    }
    fprintf(stderr, "%s failed: %s (%s)\n", step, cu_error_name(api, rc), cu_error_string(api, rc));
    return 0;
}

static double now_seconds(void) {
    struct timespec ts;
    clock_gettime(CLOCK_MONOTONIC, &ts);
    return (double)ts.tv_sec + ((double)ts.tv_nsec / 1000000000.0);
}

static size_t round_down_size(size_t value, size_t multiple) {
    if (multiple == 0 || value < multiple) {
        return value;
    }
    return value - (value % multiple);
}

static int query_compute_capability(struct cuda_api *api, CUdevice dev, int *major, int *minor) {
    int cc_major = 0;
    int cc_minor = 0;
    if (!check_rc(api,
                  "cuDeviceGetAttribute(major)",
                  api->cuDeviceGetAttribute(&cc_major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, dev))) {
        return 0;
    }
    if (!check_rc(api,
                  "cuDeviceGetAttribute(minor)",
                  api->cuDeviceGetAttribute(&cc_minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, dev))) {
        return 0;
    }
    *major = cc_major;
    *minor = cc_minor;
    return 1;
}

static int query_multiprocessor_count(struct cuda_api *api, CUdevice dev, int *count) {
    int mp_count = 0;
    if (!check_rc(api,
                  "cuDeviceGetAttribute(multiprocessors)",
                  api->cuDeviceGetAttribute(&mp_count, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT, dev))) {
        return 0;
    }
    *count = mp_count;
    return 1;
}

static size_t clamp_budget_to_free_memory(struct cuda_api *api, size_t requested_bytes) {
    size_t free_bytes = 0;
    size_t total_bytes = 0;
    size_t max_bytes = requested_bytes;

    if (!api->cuMemGetInfo) {
        return requested_bytes;
    }
    if (api->cuMemGetInfo(&free_bytes, &total_bytes) != CU_SUCCESS || free_bytes == 0) {
        return requested_bytes;
    }

    max_bytes = (free_bytes * 9u) / 10u;
    if (max_bytes < (size_t)4u * 1024u * 1024u) {
        max_bytes = (size_t)4u * 1024u * 1024u;
    }
    if (requested_bytes > max_bytes) {
        return max_bytes;
    }
    return requested_bytes;
}

static int choose_stream_count(int mp_count, int planned_profiles, size_t total_budget, int have_streams) {
    int stream_count = 1;
    if (!have_streams || mp_count <= 0 || planned_profiles <= 0) {
        return 1;
    }

    stream_count = mp_count / 8;
    if (stream_count < 2) {
        stream_count = 2;
    }
    if (stream_count > MAX_STRESS_STREAMS) {
        stream_count = MAX_STRESS_STREAMS;
    }

    while (stream_count > 1) {
        size_t per_stream_budget = total_budget / ((size_t)planned_profiles * (size_t)stream_count);
        if (per_stream_budget >= MIN_STREAM_BUDGET_BYTES) {
            break;
        }
        stream_count--;
    }
    return stream_count;
}

#if HAVE_CUBLASLT_HEADERS
static size_t clamp_single_precision_profile_budget(size_t profile_budget_bytes) {
    if (profile_budget_bytes > MAX_SINGLE_PRECISION_PROFILE_BUDGET_BYTES) {
        return MAX_SINGLE_PRECISION_PROFILE_BUDGET_BYTES;
    }
    return profile_budget_bytes;
}
#endif

static void destroy_streams(struct cuda_api *api, CUstream *streams, int count) {
    if (!api->cuStreamDestroy) {
        return;
    }
    for (int i = 0; i < count; i++) {
        if (streams[i]) {
            api->cuStreamDestroy(streams[i]);
            streams[i] = NULL;
        }
    }
}

#if HAVE_CUBLASLT_HEADERS
static void append_detail(char *buf, size_t cap, const char *fmt, ...) {
    size_t len = strlen(buf);
    if (len >= cap) {
        return;
    }
    va_list ap;
    va_start(ap, fmt);
    vsnprintf(buf + len, cap - len, fmt, ap);
    va_end(ap);
}
#endif

static int run_ptx_fallback(struct cuda_api *api,
                            CUdevice dev,
                            const char *device_name,
                            int cc_major,
                            int cc_minor,
                            int seconds,
                            int size_mb,
                            struct stress_report *report) {
    CUcontext ctx = NULL;
    CUmodule module = NULL;
    CUfunction kernel = NULL;
    uint32_t sample[256];
    CUdeviceptr device_mem[MAX_STRESS_STREAMS] = {0};
    CUstream streams[MAX_STRESS_STREAMS] = {0};
    uint32_t words[MAX_STRESS_STREAMS] = {0};
    uint32_t rounds[MAX_STRESS_STREAMS] = {0};
    void *params[MAX_STRESS_STREAMS][3];
    size_t bytes_per_stream[MAX_STRESS_STREAMS] = {0};
    unsigned long iterations = 0;
    int mp_count = 0;
    int stream_count = 1;

    memset(report, 0, sizeof(*report));
    snprintf(report->backend, sizeof(report->backend), "driver-ptx");
    snprintf(report->device, sizeof(report->device), "%s", device_name);
    report->cc_major = cc_major;
    report->cc_minor = cc_minor;
    report->buffer_mb = size_mb;

    if (!check_rc(api, "cuCtxCreate", api->cuCtxCreate(&ctx, 0, dev))) {
        return 0;
    }

    size_t requested_bytes = (size_t)size_mb * 1024u * 1024u;
    if (requested_bytes < MIN_PROFILE_BUDGET_BYTES) {
        requested_bytes = MIN_PROFILE_BUDGET_BYTES;
    }
    size_t total_bytes = clamp_budget_to_free_memory(api, requested_bytes);
    if (total_bytes < MIN_PROFILE_BUDGET_BYTES) {
        total_bytes = MIN_PROFILE_BUDGET_BYTES;
    }
    report->buffer_mb = (int)(total_bytes / (1024u * 1024u));

    if (query_multiprocessor_count(api, dev, &mp_count) &&
        api->cuStreamCreate &&
        api->cuStreamDestroy) {
        stream_count = choose_stream_count(mp_count, 1, total_bytes, 1);
    }
    if (stream_count > 1) {
        int created = 0;
        for (; created < stream_count; created++) {
            if (!check_rc(api, "cuStreamCreate", api->cuStreamCreate(&streams[created], 0))) {
                destroy_streams(api, streams, created);
                stream_count = 1;
                break;
            }
        }
    }
    report->stream_count = stream_count;

    for (int lane = 0; lane < stream_count; lane++) {
        size_t slice = total_bytes / (size_t)stream_count;
        if (lane == stream_count - 1) {
            slice = total_bytes - ((size_t)lane * (total_bytes / (size_t)stream_count));
        }
        slice = round_down_size(slice, sizeof(uint32_t));
        if (slice < MIN_PROFILE_BUDGET_BYTES) {
            slice = MIN_PROFILE_BUDGET_BYTES;
        }
        bytes_per_stream[lane] = slice;
        words[lane] = (uint32_t)(slice / sizeof(uint32_t));

        if (!check_rc(api, "cuMemAlloc", api->cuMemAlloc(&device_mem[lane], slice))) {
            goto fail;
        }
        if (!check_rc(api, "cuMemsetD8", api->cuMemsetD8(device_mem[lane], 0, slice))) {
            goto fail;
        }
        rounds[lane] = 2048;
        params[lane][0] = &device_mem[lane];
        params[lane][1] = &words[lane];
        params[lane][2] = &rounds[lane];
    }

    if (!check_rc(api,
                  "cuModuleLoadDataEx",
                  api->cuModuleLoadDataEx(&module, ptx_source, 0, NULL, NULL))) {
        goto fail;
    }
    if (!check_rc(api, "cuModuleGetFunction", api->cuModuleGetFunction(&kernel, module, "burn"))) {
        goto fail;
    }

    unsigned int threads = 256;

    double deadline = now_seconds() + (double)seconds;
    double next_sync = now_seconds() + 1.0;
    while (now_seconds() < deadline) {
        int launched = 0;
        for (int lane = 0; lane < stream_count; lane++) {
            unsigned int blocks = (unsigned int)((words[lane] + threads - 1) / threads);
            if (!check_rc(api,
                          "cuLaunchKernel",
                          api->cuLaunchKernel(kernel,
                                              blocks,
                                              1,
                                              1,
                                              threads,
                                              1,
                                              1,
                                              0,
                                              streams[lane],
                                              params[lane],
                                              NULL))) {
                goto fail;
            }
            launched++;
            iterations++;
        }
        if (launched <= 0) {
            goto fail;
        }
        double now = now_seconds();
        if (now >= next_sync || now >= deadline) {
            if (!check_rc(api, "cuCtxSynchronize", api->cuCtxSynchronize())) {
                goto fail;
            }
            next_sync = now + 1.0;
        }
    }
    api->cuCtxSynchronize();

    if (!check_rc(api, "cuMemcpyDtoH", api->cuMemcpyDtoH(sample, device_mem[0], sizeof(sample)))) {
        goto fail;
    }

    for (size_t i = 0; i < sizeof(sample) / sizeof(sample[0]); i++) {
        report->checksum += sample[i];
    }
    report->iterations = iterations;
    snprintf(report->details,
             sizeof(report->details),
             "fallback_int32=OK requested_mb=%d actual_mb=%d streams=%d per_stream_mb=%zu iterations=%lu\n",
             size_mb,
             report->buffer_mb,
             report->stream_count,
             bytes_per_stream[0] / (1024u * 1024u),
             iterations);

    for (int lane = 0; lane < stream_count; lane++) {
        if (device_mem[lane]) {
            api->cuMemFree(device_mem[lane]);
        }
    }
    destroy_streams(api, streams, stream_count);
    api->cuCtxDestroy(ctx);
    return 1;

fail:
    for (int lane = 0; lane < MAX_STRESS_STREAMS; lane++) {
        if (device_mem[lane]) {
            api->cuMemFree(device_mem[lane]);
        }
    }
    destroy_streams(api, streams, MAX_STRESS_STREAMS);
    if (ctx) {
        api->cuCtxDestroy(ctx);
    }
    return 0;
}