// ============================================================
// bench_dram_scalar_float.cu
//
// Usage:
//   ./bench_dram_scalar_float <cs|cg|cv> <launch|kernel> [working_set_mib]
//
// Examples:
//   ./bench_dram_scalar_float cs launch
//   ./bench_dram_scalar_float cg launch
//   ./bench_dram_scalar_float cv launch
//   ./bench_dram_scalar_float cs kernel
//   ./bench_dram_scalar_float cg kernel
//   ./bench_dram_scalar_float cv kernel
//   ./bench_dram_scalar_float cs launch 4096
//   ./bench_dram_scalar_float cv launch 4096
//
// 参数：
//   cache policy:
//     cs : 使用 ld.global.cs.f32 / st.global.cs.f32
//     cg : 使用 ld.global.cg.f32 / st.global.cg.f32
//     cv : 使用 ld.global.cv.f32 / st.global.cs.f32
//
//   loop mode:
//     launch : host 端多次 launch，每个 kernel 只做一次 pass
//     kernel : kernel 内部循环 iters 次，每个 kernel launch 做多次 pass
//
// 说明：
//   这是 scalar float 版本，每条 load/store 指令访问 4 bytes。
//   为了同时满足“地址可合并”和“尽量减少缓存干扰”，这里采用：
//   1) 顺序 grid-stride 访问；
//   2) 不做 warmup；
//   3) 默认每个测试只做 10 次 pass；
//   4) working set 远大于 L2。
//
// 注意：
//   PTX 中 .cv 是 load cache operator，合法形式如：
//     ld.global.cv.f32
//   但 store 没有 st.global.cv 这个合法 cache operator。
//   因此 cv 模式下 store 使用：
//     st.global.cs.f32
// ============================================================

#include <cuda_runtime.h>

#include <cstdio>
#include <cstdlib>
#include <vector>
#include <algorithm>
#include <cstdint>
#include <cstring>
#include <cmath>
#include <string>

#define CHECK_CUDA(call)                                                   \
    do {                                                                   \
        cudaError_t err = call;                                            \
        if (err != cudaSuccess) {                                          \
            fprintf(stderr, "CUDA error %s:%d: %s\n",                      \
                    __FILE__, __LINE__, cudaGetErrorString(err));          \
            std::exit(EXIT_FAILURE);                                       \
        }                                                                  \
    } while (0)

enum CachePolicy {
    CACHE_CS = 0,
    CACHE_CG = 1,
    CACHE_CV = 2
};

enum LoopMode {
    LOOP_LAUNCH = 0,
    LOOP_KERNEL = 1
};

// ------------------------------------------------------------
// .cs scalar float load
// ------------------------------------------------------------
__device__ __forceinline__ float load_cs_float(const float* ptr) {
    float v;

    asm volatile(
        "ld.global.cs.f32 %0, [%1];"
        : "=f"(v)
        : "l"(ptr)
        : "memory"
    );

    return v;
}

// ------------------------------------------------------------
// .cg scalar float load
// ------------------------------------------------------------
__device__ __forceinline__ float load_cg_float(const float* ptr) {
    float v;

    asm volatile(
        "ld.global.cg.f32 %0, [%1];"
        : "=f"(v)
        : "l"(ptr)
        : "memory"
    );

    return v;
}

// ------------------------------------------------------------
// .cv scalar float load
// ------------------------------------------------------------
__device__ __forceinline__ float load_cv_float(const float* ptr) {
    float v;

    asm volatile(
        "ld.global.cv.f32 %0, [%1];"
        : "=f"(v)
        : "l"(ptr)
        : "memory"
    );

    return v;
}

// ------------------------------------------------------------
// 按策略选择 load
// ------------------------------------------------------------
__device__ __forceinline__ float load_policy_float(
    const float* ptr,
    int policy
) {
    if (policy == CACHE_CG) {
        return load_cg_float(ptr);
    } else if (policy == CACHE_CV) {
        return load_cv_float(ptr);
    } else {
        return load_cs_float(ptr);
    }
}

// ------------------------------------------------------------
// .cs scalar float store
// ------------------------------------------------------------
__device__ __forceinline__ void store_cs_float(
    float* ptr,
    float v
) {
    asm volatile(
        "st.global.cs.f32 [%0], %1;"
        :
        : "l"(ptr), "f"(v)
        : "memory"
    );
}

// ------------------------------------------------------------
// .cg scalar float store
// ------------------------------------------------------------
__device__ __forceinline__ void store_cg_float(
    float* ptr,
    float v
) {
    asm volatile(
        "st.global.cg.f32 [%0], %1;"
        :
        : "l"(ptr), "f"(v)
        : "memory"
    );
}

// ------------------------------------------------------------
// 按策略选择 store
// ------------------------------------------------------------
__device__ __forceinline__ void store_policy_float(
    float* ptr,
    float v,
    int policy
) {
    if (policy == CACHE_CG) {
        store_cg_float(ptr, v);
    } else if (policy == CACHE_CV) {
        store_cs_float(ptr, v);
    } else {
        store_cs_float(ptr, v);
    }
}

// ------------------------------------------------------------
// 生成普通的确定性 scalar float 值
// ------------------------------------------------------------
__device__ __forceinline__ float make_normal_value_float(
    size_t i,
    int it
) {
    return 1.0f + static_cast<float>((i + static_cast<size_t>(it)) & 1023) * 0.001f;
}

// ------------------------------------------------------------
// Read kernel
//
// 顺序 grid-stride 访问，保证 warp 内地址可合并。
// ------------------------------------------------------------
__global__ void dram_read_kernel(
    const float* __restrict__ src,
    float* __restrict__ sink,
    size_t elems,
    int iters,
    int policy
) {
    size_t tid =
        static_cast<size_t>(blockIdx.x) * blockDim.x + threadIdx.x;

    size_t stride =
        static_cast<size_t>(gridDim.x) * blockDim.x;

    float acc = 0.0f;

    for (int it = 0; it < iters; ++it) {
        for (size_t i = tid; i < elems; i += stride) {
            float v = load_policy_float(src + i, policy);
            acc += v;
        }
    }

    sink[tid] = acc;
}

// ------------------------------------------------------------
// Write kernel
//
// 写入普通规律值，地址顺序可合并。
// ------------------------------------------------------------
__global__ void dram_write_kernel(
    float* __restrict__ dst,
    size_t elems,
    int iters,
    int policy
) {
    size_t tid =
        static_cast<size_t>(blockIdx.x) * blockDim.x + threadIdx.x;

    size_t stride =
        static_cast<size_t>(gridDim.x) * blockDim.x;

    for (int it = 0; it < iters; ++it) {
        for (size_t i = tid; i < elems; i += stride) {
            float v = make_normal_value_float(i, it);
            store_policy_float(dst + i, v, policy);
        }
    }
}

// ------------------------------------------------------------
// Copy kernel
//
// 顺序读取 + 顺序写入。
// ------------------------------------------------------------
__global__ void dram_copy_kernel(
    const float* __restrict__ src,
    float* __restrict__ dst,
    size_t elems,
    int iters,
    int policy
) {
    size_t tid =
        static_cast<size_t>(blockIdx.x) * blockDim.x + threadIdx.x;

    size_t stride =
        static_cast<size_t>(gridDim.x) * blockDim.x;

    for (int it = 0; it < iters; ++it) {
        for (size_t i = tid; i < elems; i += stride) {
            float v = load_policy_float(src + i, policy);
            store_policy_float(dst + i, v, policy);
        }
    }
}

// ------------------------------------------------------------
// Mixed 5-read-1-write kernel
//
// 5 路顺序读，1 路顺序写。
// ------------------------------------------------------------
__global__ void dram_mixed_5r1w_kernel(
    const float* __restrict__ a0,
    const float* __restrict__ a1,
    const float* __restrict__ a2,
    const float* __restrict__ a3,
    const float* __restrict__ a4,
    float* __restrict__ dst,
    size_t elems,
    int iters,
    int policy
) {
    size_t tid =
        static_cast<size_t>(blockIdx.x) * blockDim.x + threadIdx.x;

    size_t stride =
        static_cast<size_t>(gridDim.x) * blockDim.x;

    for (int it = 0; it < iters; ++it) {
        for (size_t i = tid; i < elems; i += stride) {
            float v0 = load_policy_float(a0 + i, policy);
            float v1 = load_policy_float(a1 + i, policy);
            float v2 = load_policy_float(a2 + i, policy);
            float v3 = load_policy_float(a3 + i, policy);
            float v4 = load_policy_float(a4 + i, policy);

            float out = v0 + v1 + v2 + v3 + v4;

            store_policy_float(dst + i, out, policy);
        }
    }
}

// ------------------------------------------------------------
// 用普通规律值初始化 float 数组
// ------------------------------------------------------------
static void init_float_normal(float* d, size_t elems, float base) {
    const size_t chunk = 1 << 20;
    std::vector<float> h(chunk);

    size_t off = 0;

    while (off < elems) {
        size_t now = std::min(chunk, elems - off);

        for (size_t i = 0; i < now; ++i) {
            size_t idx = off + i;
            h[i] = base + static_cast<float>(idx & 1023) * 0.001f;
        }

        CHECK_CUDA(cudaMemcpy(
            d + off,
            h.data(),
            now * sizeof(float),
            cudaMemcpyHostToDevice
        ));

        off += now;
    }
}

// ------------------------------------------------------------
// 统一计时函数
// ------------------------------------------------------------
template <typename Launcher>
static float time_benchmark(
    Launcher launcher,
    int warmup_iters,
    int measured_iters,
    LoopMode mode
) {
    if (warmup_iters > 0) {
        if (mode == LOOP_LAUNCH) {
            for (int r = 0; r < warmup_iters; ++r) {
                launcher(1);
            }
        } else {
            launcher(warmup_iters);
        }

        CHECK_CUDA(cudaDeviceSynchronize());
        CHECK_CUDA(cudaGetLastError());
    }

    cudaEvent_t start, stop;

    CHECK_CUDA(cudaEventCreate(&start));
    CHECK_CUDA(cudaEventCreate(&stop));

    CHECK_CUDA(cudaEventRecord(start));

    if (mode == LOOP_LAUNCH) {
        for (int r = 0; r < measured_iters; ++r) {
            launcher(1);
        }
    } else {
        launcher(measured_iters);
    }

    CHECK_CUDA(cudaEventRecord(stop));
    CHECK_CUDA(cudaEventSynchronize(stop));

    CHECK_CUDA(cudaGetLastError());

    float ms = 0.0f;
    CHECK_CUDA(cudaEventElapsedTime(&ms, start, stop));

    CHECK_CUDA(cudaEventDestroy(start));
    CHECK_CUDA(cudaEventDestroy(stop));

    return ms;
}

// ------------------------------------------------------------
// bytes / ms -> GB/s
// ------------------------------------------------------------
static double gbps(double bytes, double ms) {
    return bytes / (ms / 1000.0) / 1e9;
}

static const char* policy_name(CachePolicy p) {
    if (p == CACHE_CG) {
        return "cg";
    }

    if (p == CACHE_CV) {
        return "cv";
    }

    return "cs";
}

static const char* mode_name(LoopMode m) {
    return m == LOOP_KERNEL ? "kernel" : "launch";
}

static bool parse_policy(const char* s, CachePolicy& p) {
    if (std::strcmp(s, "cs") == 0 || std::strcmp(s, ".cs") == 0) {
        p = CACHE_CS;
        return true;
    }

    if (std::strcmp(s, "cg") == 0 || std::strcmp(s, ".cg") == 0) {
        p = CACHE_CG;
        return true;
    }

    if (std::strcmp(s, "cv") == 0 || std::strcmp(s, ".cv") == 0) {
        p = CACHE_CV;
        return true;
    }

    return false;
}

static bool parse_mode(const char* s, LoopMode& m) {
    if (std::strcmp(s, "launch") == 0) {
        m = LOOP_LAUNCH;
        return true;
    }

    if (std::strcmp(s, "kernel") == 0) {
        m = LOOP_KERNEL;
        return true;
    }

    return false;
}

static void print_usage(const char* prog) {
    fprintf(stderr,
            "Usage:\n"
            "  %s <cs|cg|cv> <launch|kernel> [working_set_mib]\n\n"
            "Examples:\n"
            "  %s cs launch\n"
            "  %s cg launch\n"
            "  %s cv launch\n"
            "  %s cs kernel\n"
            "  %s cg kernel\n"
            "  %s cv kernel 4096\n",
            prog, prog, prog, prog, prog, prog, prog);
}

int main(int argc, char** argv) {
    if (argc < 3) {
        print_usage(argv[0]);
        return 1;
    }

    CachePolicy cache_policy = CACHE_CS;
    LoopMode loop_mode = LOOP_LAUNCH;

    if (!parse_policy(argv[1], cache_policy)) {
        fprintf(stderr, "Invalid cache policy: %s\n", argv[1]);
        print_usage(argv[0]);
        return 1;
    }

    if (!parse_mode(argv[2], loop_mode)) {
        fprintf(stderr, "Invalid loop mode: %s\n", argv[2]);
        print_usage(argv[0]);
        return 1;
    }

    int dev = 0;
    CHECK_CUDA(cudaSetDevice(dev));

    cudaDeviceProp prop{};
    CHECK_CUDA(cudaGetDeviceProperties(&prop, dev));

    int sms = prop.multiProcessorCount;
    size_t l2_bytes = prop.l2CacheSize;
    size_t global_mem = prop.totalGlobalMem;

    // ------------------------------------------------------------
    // 默认 working set：
    // 至少为 64 倍 L2，或者 2 GiB，
    // 然后限制在显存的 1/8 以内（防止爆显存）。
    // ------------------------------------------------------------
    size_t default_bytes = std::max<size_t>(
        l2_bytes * 64ull,
        2ull * 1024ull * 1024ull * 1024ull
    );

    default_bytes = std::min(default_bytes, global_mem / 8);

    size_t bytes = default_bytes;

    if (argc >= 4) {
        double mib = std::atof(argv[3]);

        if (mib <= 0.0) {
            fprintf(stderr, "Invalid working_set_mib: %s\n", argv[3]);
            return 1;
        }

        bytes = static_cast<size_t>(mib * 1024.0 * 1024.0);
    }

    size_t elems = bytes / sizeof(float);
    bytes = elems * sizeof(float);

    if (elems == 0) {
        fprintf(stderr, "Working set too small.\n");
        return 1;
    }

    if (bytes * 6 > global_mem * 9 / 10) {
        fprintf(stderr,
                "Requested size too large for mixed 5R1W test.\n"
                "Need about 6x working set.\n"
                "Requested working set %.2f MiB, total needed %.2f GiB.\n",
                bytes / 1024.0 / 1024.0,
                bytes * 6.0 / 1024.0 / 1024.0 / 1024.0);
        return 1;
    }

    printf("GPU: %s\n", prop.name);
    printf("SM count: %d\n", sms);
    printf("Max threads per block: %d\n", prop.maxThreadsPerBlock);
    printf("L2 cache size: %.2f MiB\n", l2_bytes / 1024.0 / 1024.0);
    printf("Global memory: %.2f GiB\n", global_mem / 1024.0 / 1024.0 / 1024.0);
    printf("Working set per array: %.2f MiB\n", bytes / 1024.0 / 1024.0);
    printf("Float elems per array: %zu\n", elems);
    printf("Cache policy: .%s\n", policy_name(cache_policy));
    printf("Loop mode: %s\n", mode_name(loop_mode));
    printf("Address mode: coalesced grid-stride\n");
    printf("Element type: float\n");
    printf("Bytes per load/store instruction: 4\n");

    if (cache_policy == CACHE_CS) {
        printf("Load policy:  ld.global.cs.f32\n");
        printf("Store policy: st.global.cs.f32\n");
    } else if (cache_policy == CACHE_CG) {
        printf("Load policy:  ld.global.cg.f32\n");
        printf("Store policy: st.global.cg.f32\n");
    } else {
        printf("Load policy:  ld.global.cv.f32\n");
        printf("Store policy: st.global.cs.f32\n");
        printf("Note: .cv is a load cache operator; store uses .cs because st.global.cv is not valid PTX.\n");
    }

    printf("\n");

    // ------------------------------------------------------------
    // 扫描配置
    // ------------------------------------------------------------
    std::vector<int> thread_candidates;
    for (int i = 32; i <= 1536; i += 32) {
        thread_candidates.push_back(i);
    }

    std::vector<int> block_per_sm_candidates;
    for (int i = 1; i <= 24; ++i) {
        block_per_sm_candidates.push_back(i);
    }
    // std::vector<int> thread_candidates = {
    //     // 512,
    //     // 768,
    //     // 1024,
    //     // 1024
    // };

    // std::vector<int> block_per_sm_candidates = {
    //     // 12,
    //     // 1,
    //     // 24,
    //     // 24
    // };

    std::vector<int> thread_list;

    for (int t : thread_candidates) {
        if (t <= prop.maxThreadsPerBlock) {
            thread_list.push_back(t);
        }
    }

    if (thread_list.empty()) {
        fprintf(stderr, "No valid thread block size found.\n");
        return 1;
    }

    int max_threads =
        *std::max_element(thread_list.begin(), thread_list.end());

    int max_blocks_per_sm =
        *std::max_element(
            block_per_sm_candidates.begin(),
            block_per_sm_candidates.end()
        );

    int max_blocks = sms * max_blocks_per_sm;

    // ------------------------------------------------------------
    // 分配 device memory
    // ------------------------------------------------------------
    float* d_a0 = nullptr;
    float* d_a1 = nullptr;
    float* d_a2 = nullptr;
    float* d_a3 = nullptr;
    float* d_a4 = nullptr;
    float* d_dst = nullptr;

    float* d_sink = nullptr;

    CHECK_CUDA(cudaMalloc(&d_a0, bytes));
    CHECK_CUDA(cudaMalloc(&d_a1, bytes));
    CHECK_CUDA(cudaMalloc(&d_a2, bytes));
    CHECK_CUDA(cudaMalloc(&d_a3, bytes));
    CHECK_CUDA(cudaMalloc(&d_a4, bytes));
    CHECK_CUDA(cudaMalloc(&d_dst, bytes));

    CHECK_CUDA(cudaMalloc(
        &d_sink,
        static_cast<size_t>(max_blocks) *
        static_cast<size_t>(max_threads) *
        sizeof(float)
    ));

    printf("Initializing arrays with normal deterministic values...\n");

    init_float_normal(d_a0, elems, 1.0f);
    init_float_normal(d_a1, elems, 2.0f);
    init_float_normal(d_a2, elems, 3.0f);
    init_float_normal(d_a3, elems, 4.0f);
    init_float_normal(d_a4, elems, 5.0f);

    CHECK_CUDA(cudaMemset(d_dst, 0, bytes));

    CHECK_CUDA(cudaMemset(
        d_sink,
        0,
        static_cast<size_t>(max_blocks) *
        static_cast<size_t>(max_threads) *
        sizeof(float)
    ));

    printf("Initialization done.\n\n");

    // ------------------------------------------------------------
    // 为减少缓存干扰：
    // 1) 不做 warmup
    // 2) 默认每个测试只做 10 次 pass
    // ------------------------------------------------------------
    int read_iters  = 10;
    int write_iters = 10;
    int copy_iters  = 10;
    int mixed_iters = 10;

    int warmup_read  = 0;
    int warmup_write = 0;
    int warmup_copy  = 0;
    int warmup_mixed = 0;

    double read_bytes =
        static_cast<double>(bytes) * read_iters;

    double write_bytes =
        static_cast<double>(bytes) * write_iters;

    double copy_bytes =
        static_cast<double>(bytes) * copy_iters * 2.0;

    double mixed_bytes =
        static_cast<double>(bytes) * mixed_iters * 6.0;

    printf("Scan settings:\n");

    printf("  thread candidates: ");
    for (int t : thread_list) {
        printf("%d ", t);
    }
    printf("\n");

    printf("  GridBlk/SM candidates: ");
    for (int bpsm : block_per_sm_candidates) {
        printf("%d ", bpsm);
    }
    printf("\n");

    printf("  read_iters  = %d, warmup = %d\n", read_iters, warmup_read);
    printf("  write_iters = %d, warmup = %d\n", write_iters, warmup_write);
    printf("  copy_iters  = %d, warmup = %d\n", copy_iters, warmup_copy);
    printf("  mixed_iters = %d, warmup = %d\n", mixed_iters, warmup_mixed);
    printf("\n");

    printf("%10s %8s %10s %14s %16s %12s %12s %16s\n",
           "GridBlk/SM",
           "Threads",
           "Blocks",
           "GridWarp/SM",
           "Test",
           "Iters",
           "Time_ms",
           "GB/s");

    printf("----------------------------------------------------------------------------------------------------------\n");

    double best_read_gbps  = 0.0;
    double best_write_gbps = 0.0;
    double best_copy_gbps  = 0.0;
    double best_mixed_gbps = 0.0;

    int best_read_threads = 0;
    int best_read_bpsm = 0;

    int best_write_threads = 0;
    int best_write_bpsm = 0;

    int best_copy_threads = 0;
    int best_copy_bpsm = 0;

    int best_mixed_threads = 0;
    int best_mixed_bpsm = 0;

    int policy_int = static_cast<int>(cache_policy);

    // ------------------------------------------------------------
    // 正式扫描
    // ------------------------------------------------------------
    for (int threads : thread_list) {
        for (int blocks_per_sm : block_per_sm_candidates) {
            int blocks = sms * blocks_per_sm;

            double grid_warps_per_sm =
                static_cast<double>(blocks_per_sm) *
                static_cast<double>(threads) / 32.0;

            size_t sink_elems =
                static_cast<size_t>(blocks) *
                static_cast<size_t>(threads);

            CHECK_CUDA(cudaMemset(
                d_sink,
                0,
                sink_elems * sizeof(float)
            ));

            // ---------------------------
            // Read-only
            // ---------------------------
            float ms_read = time_benchmark(
                [&](int kernel_iters) {
                    dram_read_kernel<<<blocks, threads>>>(
                        d_a0,
                        d_sink,
                        elems,
                        kernel_iters,
                        policy_int
                    );
                },
                warmup_read,
                read_iters,
                loop_mode
            );

            double read_gbps = gbps(read_bytes, ms_read);

            printf("%10d %8d %10d %14.1f %16s %12d %12.3f %16.2f\n",
                   blocks_per_sm,
                   threads,
                   blocks,
                   grid_warps_per_sm,
                   "Read",
                   read_iters,
                   ms_read,
                   read_gbps);

            if (read_gbps > best_read_gbps) {
                best_read_gbps = read_gbps;
                best_read_threads = threads;
                best_read_bpsm = blocks_per_sm;
            }

            // ---------------------------
            // Write-only
            // ---------------------------
            float ms_write = time_benchmark(
                [&](int kernel_iters) {
                    dram_write_kernel<<<blocks, threads>>>(
                        d_dst,
                        elems,
                        kernel_iters,
                        policy_int
                    );
                },
                warmup_write,
                write_iters,
                loop_mode
            );

            double write_gbps = gbps(write_bytes, ms_write);

            printf("%10d %8d %10d %14.1f %16s %12d %12.3f %16.2f\n",
                   blocks_per_sm,
                   threads,
                   blocks,
                   grid_warps_per_sm,
                   "Write",
                   write_iters,
                   ms_write,
                   write_gbps);

            if (write_gbps > best_write_gbps) {
                best_write_gbps = write_gbps;
                best_write_threads = threads;
                best_write_bpsm = blocks_per_sm;
            }

            // ---------------------------
            // Copy 1R1W
            // ---------------------------
            float ms_copy = time_benchmark(
                [&](int kernel_iters) {
                    dram_copy_kernel<<<blocks, threads>>>(
                        d_a0,
                        d_dst,
                        elems,
                        kernel_iters,
                        policy_int
                    );
                },
                warmup_copy,
                copy_iters,
                loop_mode
            );

            double copy_gbps = gbps(copy_bytes, ms_copy);

            printf("%10d %8d %10d %14.1f %16s %12d %12.3f %16.2f\n",
                   blocks_per_sm,
                   threads,
                   blocks,
                   grid_warps_per_sm,
                   "Copy_1R1W",
                   copy_iters,
                   ms_copy,
                   copy_gbps);

            if (copy_gbps > best_copy_gbps) {
                best_copy_gbps = copy_gbps;
                best_copy_threads = threads;
                best_copy_bpsm = blocks_per_sm;
            }

            // ---------------------------
            // Mixed 5R1W
            // ---------------------------
            float ms_mixed = time_benchmark(
                [&](int kernel_iters) {
                    dram_mixed_5r1w_kernel<<<blocks, threads>>>(
                        d_a0,
                        d_a1,
                        d_a2,
                        d_a3,
                        d_a4,
                        d_dst,
                        elems,
                        kernel_iters,
                        policy_int
                    );
                },
                warmup_mixed,
                mixed_iters,
                loop_mode
            );

            double mixed_gbps = gbps(mixed_bytes, ms_mixed);

            printf("%10d %8d %10d %14.1f %16s %12d %12.3f %16.2f\n",
                   blocks_per_sm,
                   threads,
                   blocks,
                   grid_warps_per_sm,
                   "Mixed_5R1W",
                   mixed_iters,
                   ms_mixed,
                   mixed_gbps);

            if (mixed_gbps > best_mixed_gbps) {
                best_mixed_gbps = mixed_gbps;
                best_mixed_threads = threads;
                best_mixed_bpsm = blocks_per_sm;
            }

            printf("----------------------------------------------------------------------------------------------------------\n");
        }
    }

    printf("\nBest results:\n");

    printf("%16s %14s %12s %16s\n",
           "Test",
           "GridBlk/SM",
           "Threads",
           "Best GB/s");

    printf("%16s %14d %12d %16.2f\n",
           "Read",
           best_read_bpsm,
           best_read_threads,
           best_read_gbps);

    printf("%16s %14d %12d %16.2f\n",
           "Write",
           best_write_bpsm,
           best_write_threads,
           best_write_gbps);

    printf("%16s %14d %12d %16.2f\n",
           "Copy_1R1W",
           best_copy_bpsm,
           best_copy_threads,
           best_copy_gbps);

    printf("%16s %14d %12d %16.2f\n",
           "Mixed_5R1W",
           best_mixed_bpsm,
           best_mixed_threads,
           best_mixed_gbps);

    printf("\nConfiguration summary:\n");
    printf("  Cache policy: .%s\n", policy_name(cache_policy));
    printf("  Loop mode: %s\n", mode_name(loop_mode));
    printf("  Address mode: coalesced grid-stride\n");
    printf("  Element type: float\n");

    if (cache_policy == CACHE_CS) {
        printf("  Load instruction:  ld.global.cs.f32\n");
        printf("  Store instruction: st.global.cs.f32\n");
    } else if (cache_policy == CACHE_CG) {
        printf("  Load instruction:  ld.global.cg.f32\n");
        printf("  Store instruction: st.global.cg.f32\n");
    } else {
        printf("  Load instruction:  ld.global.cv.f32\n");
        printf("  Store instruction: st.global.cs.f32\n");
        printf("  Note: .cv applies to load only; store uses .cs because st.global.cv is not valid PTX.\n");
    }

    printf("  Bytes per load/store instruction: 4\n");

    if (loop_mode == LOOP_LAUNCH) {
        printf("  Meaning of Iters: number of host-side kernel launches.\n");
        printf("  Each kernel launch internally uses iters = 1.\n");
    } else {
        printf("  Meaning of Iters: number of inner loops inside one kernel launch.\n");
        printf("  Host launches one timed kernel per test/configuration.\n");
    }

    CHECK_CUDA(cudaFree(d_a0));
    CHECK_CUDA(cudaFree(d_a1));
    CHECK_CUDA(cudaFree(d_a2));
    CHECK_CUDA(cudaFree(d_a3));
    CHECK_CUDA(cudaFree(d_a4));
    CHECK_CUDA(cudaFree(d_dst));
    CHECK_CUDA(cudaFree(d_sink));

    return 0;
}