/**
 * ruRE CUDA launch plan, part 2 (ADR-501): what `analyzeCuda` found, turned into a plan for measuring. For each kernel of each architecture: the
 * theoretical occupancy at a set of block sizes, what limits it, and a few plain flags; and, on request, a small CUDA driver-API program (as text)
 * that a person can build and run to check those numbers on their own GPU. Pure: data in, data and text out. Nothing here compiles, loads or runs
 * GPU code, and the harness text holds nothing taken from the analyzed file (the kernel name and the file are arguments to the program).
 */
import type { CudaInfo } from './cuda-parse.js';
import { type Limiter, type Provenance } from './cuda-occupancy.js';
export declare const DEFAULT_BLOCK_SIZES: readonly number[];
export declare const MAX_BLOCK_SIZES = 24;
export declare const MAX_DYNAMIC_SHARED_BYTES: number;
export declare const MAX_PLANNED_KERNELS = 200;
export interface PlanOptions {
    /** Only this architecture (`sm_86`). */
    arch?: string;
    blockSizes?: readonly number[];
    /** Dynamic shared memory per block the launch will ask for; it is not in the file, so it is a parameter. */
    dynamicSharedBytes?: number;
    includeHarness?: boolean;
}
export type PlanFlag = 'local-memory' | 'register-limited' | 'shared-memory-limited';
export interface PlannedBlockSize {
    blockSize: number;
    activeBlocksPerSM: number;
    activeWarpsPerSM: number;
    /** 0 to 1, four places. */
    occupancy: number;
    limitedBy: Limiter[];
}
export interface PlannedKernel {
    name: string;
    identifier: string | null;
    /** `staticSharedBytes` is the figure the driver and occupancy calculation use; `sharedSectionBytes` is what the file's section says (cuobjdump's SHARED), which differs on sm_90 and later. */
    resources: {
        registers: number | null;
        staticSharedBytes: number;
        sharedSectionBytes: number;
        localBytes: number;
        constantBytes: number | null;
        codeBytes: number | null;
    };
    /** Absent when the file gave no register count: occupancy cannot be worked out without it. */
    blockSizes?: PlannedBlockSize[];
    best?: {
        blockSizes: number[];
        occupancy: number;
        limitedBy: Limiter[];
    };
    flags: PlanFlag[];
    note?: string;
}
export interface PlannedArch {
    arch: string;
    /** False when there is no limits table for this architecture; then there are no kernels below. */
    supported: boolean;
    provenance: Provenance | null;
    limits?: {
        maxThreadsPerSM: number;
        maxBlocksPerSM: number;
        regsPerSM: number;
        sharedPerSM: number;
        sharedReservedPerBlock: number;
    };
    kernels: PlannedKernel[];
    note?: string;
}
export interface LaunchPlan {
    archs: PlannedArch[];
    blockSizes: number[];
    dynamicSharedBytes: number;
    assumptions: string[];
    harness?: string;
    truncated: boolean;
}
export declare const ASSUMPTIONS: readonly string[];
export declare function planCuda(info: CudaInfo, options?: PlanOptions): LaunchPlan;
/**
 * A program a person can build and run. It is text, fixed in this file: it contains nothing from the analyzed file. Run as
 *   ./occ_check <file.cubin|file.fatbin> <kernel name as listed> [dynamicSharedBytes]
 * it asks the CUDA driver how many blocks fit per multiprocessor at each block size (compare with the plan), and, if built with -DRURE_BENCH and a
 * rure_bench.h of your own that sets up the kernel's arguments, times launches. It loads the file's code onto the GPU; the benchmark mode runs it.
 */
export declare const HARNESS_SOURCE = "// occ_check.cpp: written by ruflo's rure_cuda_launch_plan. ruRE did not build or run this; building and running it is your decision.\n// Build:  nvcc -O2 -o occ_check occ_check.cpp -lcuda            (add -DRURE_BENCH with your own rure_bench.h to time launches)\n// Run:    ./occ_check <file.cubin|file.fatbin> <kernel name> [dynamicSharedBytes]\n// Without RURE_BENCH it loads the code onto device 0 and asks the driver for occupancy; it launches nothing.\n// With RURE_BENCH it LAUNCHES THE KERNEL with the arguments your rure_bench.h gives it. Only do that for code you trust, on a GPU you can spare.\n#include <cuda.h>\n#include <algorithm>\n#include <cstdio>\n#include <cstdlib>\n#include <fstream>\n#include <iterator>\n#include <vector>\n#ifdef RURE_BENCH\n// rure_bench.h must define:\n//   bool rure_prepare(void*** kernelParams);   allocate buffers, fill an array of pointers to the kernel's arguments\n//   unsigned rure_grid(int blockSize);         how many blocks to launch at this block size\n//   void rure_release();\n#include \"rure_bench.h\"\n#endif\n\n#define CK(call) do { CUresult r_ = (call); if (r_ != CUDA_SUCCESS) { const char* m_ = nullptr; cuGetErrorString(r_, &m_); fprintf(stderr, \"line %d: %s\\n\", __LINE__, m_ ? m_ : \"error\"); return 1; } } while (0)\n\nint main(int argc, char** argv) {\n  if (argc < 3) { fprintf(stderr, \"usage: %s <file.cubin|file.fatbin> <kernel name> [dynamicSharedBytes]\\n\", argv[0]); return 2; }\n  size_t dyn = argc > 3 ? strtoull(argv[3], nullptr, 10) : 0;\n  std::ifstream in(argv[1], std::ios::binary);\n  std::vector<char> image((std::istreambuf_iterator<char>(in)), std::istreambuf_iterator<char>());\n  if (image.empty()) { fprintf(stderr, \"cannot read %s\\n\", argv[1]); return 2; }\n\n  CK(cuInit(0));\n  CUdevice dev; CK(cuDeviceGet(&dev, 0));\n  CUcontext ctx; CK(cuDevicePrimaryCtxRetain(&ctx, dev)); CK(cuCtxSetCurrent(ctx));\n  CUmodule mod; CK(cuModuleLoadData(&mod, image.data()));\n  CUfunction fn; CK(cuModuleGetFunction(&fn, mod, argv[2]));\n\n  int regs, shared, local, maxThreads, maxThreadsSM, major, minor, blocksSM, regsSM, sharedSM, reserved, optin;\n  CK(cuFuncGetAttribute(&regs, CU_FUNC_ATTRIBUTE_NUM_REGS, fn));\n  CK(cuFuncGetAttribute(&shared, CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES, fn));\n  CK(cuFuncGetAttribute(&local, CU_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES, fn));\n  CK(cuFuncGetAttribute(&maxThreads, CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, fn));\n  CK(cuDeviceGetAttribute(&maxThreadsSM, CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_MULTIPROCESSOR, dev));\n  CK(cuDeviceGetAttribute(&major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, dev));\n  CK(cuDeviceGetAttribute(&minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, dev));\n  CK(cuDeviceGetAttribute(&blocksSM, CU_DEVICE_ATTRIBUTE_MAX_BLOCKS_PER_MULTIPROCESSOR, dev));\n  CK(cuDeviceGetAttribute(&regsSM, CU_DEVICE_ATTRIBUTE_MAX_REGISTERS_PER_MULTIPROCESSOR, dev));\n  CK(cuDeviceGetAttribute(&sharedSM, CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_MULTIPROCESSOR, dev));\n  CK(cuDeviceGetAttribute(&reserved, CU_DEVICE_ATTRIBUTE_RESERVED_SHARED_MEMORY_PER_BLOCK, dev));\n  CK(cuDeviceGetAttribute(&optin, CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK_OPTIN, dev));\n  printf(\"device sm_%d%d: maxThreadsPerSM %d, maxBlocksPerSM %d, regsPerSM %d, sharedPerSM %d, sharedReservedPerBlock %d, sharedPerBlockOptin %d\\n\", major, minor, maxThreadsSM, blocksSM, regsSM, sharedSM, reserved, optin);\n  printf(\"kernel: registers %d, static shared %d, local %d, max threads per block %d\\n\", regs, shared, local, maxThreads);\n  printf(\"blockSize activeBlocksPerSM activeWarpsPerSM occupancy%s\\n\",\n#ifdef RURE_BENCH\n         \" medianMs\");\n#else\n         \"\");\n#endif\n\n  const int sizes[] = {32, 64, 96, 128, 192, 256, 384, 512, 768, 1024};\n  for (int b : sizes) {\n    int blocks = 0;\n    if (cuOccupancyMaxActiveBlocksPerMultiprocessor(&blocks, fn, b, dyn) != CUDA_SUCCESS) blocks = 0;\n    int warps = blocks * ((b + 31) / 32);\n    printf(\"%d %d %d %.4f\", b, blocks, warps, (double)warps / (maxThreadsSM / 32));\n#ifdef RURE_BENCH\n    double ms = -1;\n    if (blocks > 0) {\n      void** params = nullptr;\n      if (!rure_prepare(&params)) { fprintf(stderr, \"rure_prepare failed\\n\"); return 1; }\n      CUevent t0, t1; CK(cuEventCreate(&t0, CU_EVENT_DEFAULT)); CK(cuEventCreate(&t1, CU_EVENT_DEFAULT));\n      std::vector<float> times;\n      for (int i = 0; i < 23; i++) {\n        CK(cuEventRecord(t0, 0));\n        CK(cuLaunchKernel(fn, rure_grid(b), 1, 1, b, 1, 1, (unsigned)dyn, 0, params, 0));\n        CK(cuEventRecord(t1, 0)); CK(cuEventSynchronize(t1));\n        float t; CK(cuEventElapsedTime(&t, t0, t1));\n        if (i >= 3) times.push_back(t);\n      }\n      std::sort(times.begin(), times.end());\n      ms = times[times.size() / 2];\n      rure_release();\n    }\n    printf(\" %.4f\", ms);\n#endif\n    printf(\"\\n\");\n  }\n  return 0;\n}\n";
//# sourceMappingURL=cuda-plan.d.ts.map