#include <cuda.h>

#include <cstdio>
#include <cstdlib>

#define CU_CHECK(call)                                                        \
    do {                                                                      \
        const CUresult status = (call);                                       \
        if (status != CUDA_SUCCESS) {                                         \
            const char* message = nullptr;                                    \
            cuGetErrorString(status, &message);                               \
            std::fprintf(stderr, "%s:%d Driver error: %s\n", __FILE__,       \
                         __LINE__, message ? message : "unknown");            \
            std::exit(EXIT_FAILURE);                                          \
        }                                                                     \
    } while (0)

int main(int argc, char** argv) {
    if (argc != 2) {
        std::fprintf(stderr, "usage: %s kernel.cubin\n", argv[0]);
        return EXIT_FAILURE;
    }

    CU_CHECK(cuInit(0));
    CUdevice device{};
    CUcontext context{};
    CUmodule module{};
    CUfunction function{};
    CUdeviceptr output{};
    CU_CHECK(cuDeviceGet(&device, 0));
    CU_CHECK(cuCtxCreate(&context, 0, device));
    CU_CHECK(cuModuleLoad(&module, argv[1]));
    CU_CHECK(cuModuleGetFunction(&function, module, "fill_kernel"));

    int n = 1024;
    float value = 7.0F;
    CU_CHECK(cuMemAlloc(&output, n * sizeof(float)));
    void* arguments[] = {&output, &n, &value};
    CU_CHECK(cuLaunchKernel(function, (n + 255) / 256, 1, 1, 256, 1, 1, 0,
                            nullptr, arguments, nullptr));
    CU_CHECK(cuCtxSynchronize());

    float first = 0.0F;
    CU_CHECK(cuMemcpyDtoH(&first, output, sizeof(first)));
    std::printf("module=%s first=%.1f\n", argv[1], first);

    CU_CHECK(cuMemFree(output));
    CU_CHECK(cuModuleUnload(module));
    CU_CHECK(cuCtxDestroy(context));
    return first == value ? EXIT_SUCCESS : EXIT_FAILURE;
}