auto_tile.h

#pragma once
// Runtime tile-width picker.
//
// Each version (V2..V6) sizes its working set as
//     bytes(T) = (slabs + plan_ws_slabs) * streams * N * sizeof(cufftComplex)
// where slabs and streams depend on the version. We probe free GPU memory
// with cudaMemGetInfo and return the largest T in {1024, 512, 256, 128}
// that fits inside an 80% budget.
//
// plan_ws_slabs is set to 2 because cuFFT Bluestein workspaces for
// prime-ish N can reach ~2x the slab footprint (measured: 410 MB at
// N=52329 T=512, vs 214 MB per slab).

#include <cstdio>
#include <cstring>
#include <cuda_runtime.h>
#include <cufft.h>

namespace sbfft {

inline int pick_tile_T(int N, const char* version,
                       int n_streams = 1, bool verbose = true) {
    size_t free_b = 0, total_b = 0;
    if (cudaMemGetInfo(&free_b, &total_b) != cudaSuccess) {
        if (verbose) std::fprintf(stderr,
            "[tile] cudaMemGetInfo failed; defaulting to T=512\n");
        return 512;
    }
    const size_t budget = static_cast<size_t>(static_cast<double>(free_b) * 0.80);

    int slabs = 1;
    int streams = 1;
    if (std::strcmp(version, "v4") == 0 || std::strcmp(version, "v6") == 0) {
        slabs = 2;
    } else if (std::strcmp(version, "v5") == 0) {
        slabs = 2;
        streams = n_streams > 0 ? n_streams : 2;
    }
    const int plan_ws_slabs = 2;
    const size_t bytes_per_T =
        static_cast<size_t>(slabs + plan_ws_slabs) *
        static_cast<size_t>(streams) *
        static_cast<size_t>(N) *
        sizeof(cufftComplex);

    int chosen = 128;
    const int candidates[] = {1024, 512, 256, 128};
    for (int T : candidates) {
        if (static_cast<size_t>(T) * bytes_per_T <= budget) { chosen = T; break; }
    }
    if (verbose) {
        std::printf("[tile] auto: N=%d %s streams=%d -> T=%d "
                    "(free=%.2f GB, est=%.2f GB)\n",
                    N, version, streams, chosen,
                    static_cast<double>(free_b) / (1ull << 30),
                    static_cast<double>(static_cast<size_t>(chosen) * bytes_per_T)
                        / (1ull << 30));
    }
    return chosen;
}

} // namespace sbfft