v1 init

.gitattributes

@@ -33,3 +33,4 @@ saved_model/**/* filter=lfs diff=lfs merge=lfs -text
 *.zip filter=lfs diff=lfs merge=lfs -text
 *.zst filter=lfs diff=lfs merge=lfs -text
 *tfevents* filter=lfs diff=lfs merge=lfs -text
+*.so filter=lfs diff=lfs merge=lfs -text

Dockerfile

@@ -0,0 +1,37 @@
+FROM nvidia/cuda:12.4.0-devel-ubuntu22.04
+
+ENV DEBIAN_FRONTEND=noninteractive
+RUN apt-get update && \
+    apt-get install -y --no-install-recommends wget ca-certificates && \
+    wget -q https://developer.nvidia.com/downloads/assets/tools/secure/nsight-systems/2025_3/NsightSystems-linux-cli-public-2025.3.1.90-3582212.deb -O /tmp/nsys.deb && \
+    apt-get install -y /tmp/nsys.deb && \
+    rm -f /tmp/nsys.deb && \
+    apt-get install -y python3.11 python3.11-distutils python3.11-dev python3-pip && \
+    update-alternatives --install /usr/bin/python python /usr/bin/python3.11 1 && \
+    python -m pip install --upgrade pip && \
+    rm -rf /var/lib/apt/lists/*
+
+RUN apt-get update && apt-get install -y gcc g++ build-essential && \
+    python -m pip install --no-cache-dir nuitka
+
+# 可选：创建非 root 用户
+RUN useradd -m -u 1000 user
+USER user
+ENV PATH="/home/user/.local/bin:$PATH"
+
+WORKDIR /app
+
+# 安装 Python 依赖
+COPY --chown=user requirements.txt .
+RUN pip install --no-cache-dir -r requirements.txt
+
+# 复制源码
+COPY --chown=user app.py .
+COPY --chown=user tool /app/tool
+RUN chown -R user:user /app
+
+# Gradio 默认监听 7860
+EXPOSE 7860
+
+# 直接启动 Python 脚本
+CMD ["python", "app.py"]

app.py

@@ -0,0 +1,151 @@
+import os, tempfile, time
+import gradio as gr
+from tool.test import run_autotune_pipeline, DATA_DIR
+
+# ---------- Core callback ----------
+
+def get_test_text(test_file, test_data_input):
+    if test_file is not None:
+        if hasattr(test_file, "read"):
+            return test_file.read().decode("utf-8")
+        elif hasattr(test_file, "data"):
+            return test_file.data if isinstance(test_file.data, str) else test_file.data.decode("utf-8")
+        elif hasattr(test_file, "name") and os.path.exists(test_file.name):
+            with open(test_file.name, "r", encoding="utf-8") as f:
+                return f.read()
+    # fallback to textbox
+    return test_data_input or ""
+
+def generate_kernel(text_input, test_data_input, test_file, n_iters, progress=gr.Progress()):
+    """
+    text_input : string from textbox (NL description or base CUDA code)
+    test_data_input: test data (variable name, data)
+    file_input : gr.File upload object (or None)
+    Returns   : (kernel_code_str, downloadable_file_path)
+    """
+    progress((0, n_iters), desc="Initializing...")
+    # 1) Select input source
+
+    if not text_input.strip():
+        return "⚠️ Please paste a description or baseline CUDA code."
+        
+    # td = tempfile.mkdtemp(prefix="auto_")
+
+    # # ------- select test data source -------
+    # if test_file is not None and test_file.size > 0:
+    #     test_text = test_file.read().decode("utf-8")
+    # elif test_data_input.strip():
+    #     test_text = test_data_input
+    # else:
+    #     return "Test data required: either fill Test Data Input or upload a .txt file.", "", None
+    
+    # src_path = os.path.join(td, f"input_{int(time.time())}.txt")
+    # test_path = os.path.join(td, f"test_data_{int(time.time())}.txt")
+    
+    # with open(src_path, "w") as f:
+    #     f.write(text_input)
+
+    # with open(test_path, "w") as f:
+    #     f.write(test_data_input or "")
+
+    # if test_file is not None:
+    #     test_text = test_file.read().decode("utf-8")
+    # else:
+    #     test_text = test_data_input
+
+    test_text = get_test_text(test_file, test_data_input)
+
+    if not test_text.strip():
+        return "⚠️ Test data required."
+
+
+    best_code = ""
+    for info in run_autotune_pipeline(
+        input_code=text_input,
+        test_data_input=test_text,
+        test_file=None,
+        bin_dir=DATA_DIR,
+        max_iterations=int(n_iters)
+    ):
+        # 1) update progress bar (if iteration known)
+        if info["iteration"] is not None:
+            # print(f"Iteration {info['iteration']} / {n_iters}: {info['message']}")
+            progress((info["iteration"], n_iters), desc=info["message"])
+
+        # 3) kernel output only when we get new code
+        if info["code"]:
+            best_code = info["code"]
+
+    # TBD: download button
+    return best_code
+
+
+# ---------- Gradio UI ----------
+with gr.Blocks(
+    title="KernelPilot", 
+    theme=gr.themes.Soft(
+        text_size="lg", 
+        font=[
+            "system-ui",
+            "-apple-system",
+            "BlinkMacSystemFont",
+            "Segoe UI",
+            "Roboto",
+            "Helvetica Neue",
+            "Arial",
+            "Noto Sans",
+            "sans-serif"
+    ])) as demo:
+    gr.Markdown(
+        """# 🚀 KernelPilot Optimizer  
+Enter a code, test data, then click **Generate** to obtain the optimized kernel function."""
+    )
+
+    with gr.Row():
+        txt_input = gr.Textbox(
+            label="📝 Input",
+            lines=10,
+            placeholder="Enter the code",
+            scale=3
+        )
+        level = gr.Number(
+            label="Optimazation Level",
+            minimum=1,
+            maximum=5,
+            value=5,
+            step=1,
+            scale=1
+        )
+
+    with gr.Row():
+            test_data_input = gr.Textbox(
+                label="Test Data Input",
+                lines=10,
+                placeholder="<number_of_test_cases>\n<number_of_variables>\n\n<variable_1_name>\n<variable_1_testcase_1_data>\n<variable_1_testcase_2_data>\n...\n<variable_1_testcase_N_data>\n\n<variable_2_name>\n<variable_2_testcase_1_data>\n...\n<variable_2_testcase_N_data>\n\n...",
+                scale=2
+            )
+            test_file = gr.File(
+                label="Upload Test Data (.txt)",
+                file_types=["text"],
+                scale=1
+            )
+
+    gen_btn = gr.Button("⚡ Generate")
+
+    kernel_output = gr.Code(
+        label="🎯 Tuned CUDA Kernel",
+        language="cpp"
+    )
+
+    gen_btn.click(
+        fn=generate_kernel,
+        inputs=[txt_input, test_data_input, test_file, level],
+        outputs=[kernel_output],
+        queue=True,               # keeps requests queued
+        show_progress=True,  # show progress bar
+        show_progress_on=kernel_output  # update log box with progress
+    )
+
+if __name__ == "__main__":
+    demo.queue(default_concurrency_limit=1, max_size=50)
+    demo.launch(server_name="0.0.0.0", server_port=7860)

requirements.txt

@@ -0,0 +1,4 @@
+gradio>=4.30
+openai>=1.30.5
+httpx>=0.27.0
+numpy>=1.26.0 

tool/formatted_file_input.txt

@@ -0,0 +1,14 @@
+2
+3
+
+mse_preds
+[1.0, 2.0, 3.0, 4.0]
+[10.0, 20.0, 30.0]
+
+mse_targets
+[1.5, 2.5, 3.5, 4.5]
+[12.0, 18.0, 33.0]
+
+mse_ref
+0.25
+5.67

tool/input.txt

@@ -0,0 +1,84 @@
+#include <iostream>
+#include <cuda_runtime.h>
+#include <fstream>
+#include <cmath>
+#include <vector>
+#include <string>
+
+__global__ void mseKernel(const float* predictions, const float* targets, size_t numElements, float* sum) {
+    size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
+    if (idx < numElements) {
+        float diff = predictions[idx] - targets[idx];
+        float sq_diff = diff * diff;
+        atomicAdd(sum, sq_diff);
+    }
+}
+
+void read_binary(const std::string& filename, float* data, size_t size) {
+    std::ifstream in(filename, std::ios::binary);
+    if (!in) {
+        std::cerr << "Cannot open file: " << filename << std::endl;
+        exit(1);
+    }
+    in.read(reinterpret_cast<char*>(data), size * sizeof(float));
+    in.close();
+}
+
+// test
+bool compare_scalar(float a, float b, float tol = 1e-1f) {
+    return std::fabs(a - b) < tol;
+}
+
+int main() {
+    std::vector<size_t> sizes = {1 << 10, 1 << 12, 1 << 14, 1 << 16, 1 << 18};
+    bool all_passed = true;
+
+    for (int t = 0; t < sizes.size(); ++t) {
+        size_t N = sizes[t];
+        size_t input_size = N * sizeof(float);
+        
+        // test
+        std::string pred_file = "data/mse_preds_" + std::to_string(t + 1) + ".bin";
+        std::string target_file = "data/mse_targets_" + std::to_string(t + 1) + ".bin";
+        std::string ref_file = "data/mse_ref_" + std::to_string(t + 1) + ".bin";
+
+        float* h_preds = (float*)malloc(input_size);
+        float* h_targets = (float*)malloc(input_size);
+        float h_mse_ref;
+
+        read_binary(pred_file, h_preds, N);
+        read_binary(target_file, h_targets, N);
+        read_binary(ref_file, &h_mse_ref, 1);
+
+        float *d_preds, *d_targets, *d_sum;
+        cudaMalloc(&d_preds, input_size);
+        cudaMalloc(&d_targets, input_size);
+        cudaMalloc(&d_sum, sizeof(float));
+        cudaMemcpy(d_preds, h_preds, input_size, cudaMemcpyHostToDevice);
+        cudaMemcpy(d_targets, h_targets, input_size, cudaMemcpyHostToDevice);
+        cudaMemset(d_sum, 0, sizeof(float));
+
+        int threads = 256;
+        int blocks = (N + threads - 1) / threads;
+        mseKernel<<<blocks, threads>>>(d_preds, d_targets, N, d_sum);
+
+        float h_sum = 0.0f;
+        cudaMemcpy(&h_sum, d_sum, sizeof(float), cudaMemcpyDeviceToHost);
+        float mse = h_sum / N;
+
+        if (!compare_scalar(mse, h_mse_ref)) {
+            std::cout << "F" << std::endl;
+            all_passed = false;
+
+            cudaFree(d_preds); cudaFree(d_targets); cudaFree(d_sum);
+            free(h_preds); free(h_targets);
+            break;
+        }
+
+        cudaFree(d_preds); cudaFree(d_targets); cudaFree(d_sum);
+        free(h_preds); free(h_targets);
+    }
+
+    if (all_passed) std::cout << "T" << std::endl;
+    return 0;
+}

tool/prompt_a.txt

@@ -0,0 +1,22 @@
+The task is to write a CUDA kernel function on GPU, and we have the benchmark code for this task: 
+[code]
+
+Optimize the kernel function for less execution time on GPU.
+The output should be the content of whole .cu file containing ONE kernel function.
+Do not modify the test part. Note the test data contains exactly four input sets. The generated .cu file must ensure that for each input set, the kernel function is called exactly once, resulting in a total of four kernel invocations. Do not include any extra timing logic, profiling wrappers, or repeat kernel calls that could cause each input to trigger multiple kernel launches.
+
+When generating CUDA code, you must produce a complete, standalone, compilable program, not just a kernel or code fragment.
+The program should include headers, data structures, the kernel definition, a main() function that allocates memory, launches the kernel, and prints "T" or "F" based on correctness.
+Follow these strict rules:
+1. Always make the code fully self-contained and directly compilable with nvcc file.cu -o file. No missing functions, dependencies, or external headers.
+2. Do not use std::max, std::min, or std::abs in device code. Always use fmaxf, fminf, and fabsf instead.
+3. Do not use INFINITY; use CUDART_INF_F or a large constant (e.g., 1e30f) instead.
+4. Include all required headers: <vector>, <cmath>, <cstdint>, and <cuda_runtime.h>.
+5. Avoid duplicate includes and never mark __global__ functions as inline.
+
+Please focus your modifications only in the CUDA kernel function part. Avoid changing any other parts of the program, including:
+- data loading, I/O, or test logic
+- host-side function definitions
+- main() function or CUDA memory allocation logic
+
+Do not rename variables or structs defined outside the kernel region.

tool/prompt_b.txt

@@ -0,0 +1,5 @@
+Modify the code with the execution error result.
+The output should be the content of whole .cu file containing ONE kernel function.
+Do not modify the test part. Note the test data contains exactly four input sets. The generated .cu file must ensure that for each input set, the kernel function is called exactly once, resulting in a total of four kernel invocations. Do not include any extra timing logic, profiling wrappers, or repeat kernel calls that could cause each input to trigger multiple kernel launches.
+
+The execution output is:

tool/prompt_c.txt

@@ -0,0 +1,3 @@
+Optimize the kernel function for less execution time on GPU.
+The output should be the content of whole .cu file containing ONE kernel function.
+Do not modify the test part. Note the test data contains exactly four input sets. The generated .cu file must ensure that for each input set, the kernel function is called exactly once, resulting in a total of four kernel invocations. Do not include any extra timing logic, profiling wrappers, or repeat kernel calls that could cause each input to trigger multiple kernel launches.

tool/prompt_e.txt

@@ -0,0 +1,3 @@
+The code failed to launch the kernel. Modify the code.
+The output should be the content of whole .cu file containing ONE kernel function.
+Do not modify the test part. Note the test data contains exactly four input sets. The generated .cu file must ensure that for each input set, the kernel function is called exactly once, resulting in a total of four kernel invocations. Do not include any extra timing logic, profiling wrappers, or repeat kernel calls that could cause each input to trigger multiple kernel launches.

tool/prompt_f.txt

@@ -0,0 +1,3 @@
+The result is not the same with the reference output. Modify the code.
+The output should be the content of whole .cu file containing ONE kernel function.
+Do not modify the test part. Note the test data contains exactly four input sets. The generated .cu file must ensure that for each input set, the kernel function is called exactly once, resulting in a total of four kernel invocations. Do not include any extra timing logic, profiling wrappers, or repeat kernel calls that could cause each input to trigger multiple kernel launches.

tool/prompt_norm.txt

@@ -0,0 +1,134 @@
+You are an expert CUDA engineer.
+
+Your task is:
+1. Completely ignore all input-loading code, file paths, main() logic, correctness checks,
+   device allocations, free(), test loops, or anything unrelated to the CUDA kernel.
+2. Extract ONLY the CUDA kernel from the provided code:
+    - The kernel must keep the original function name.
+    - The kernel must keep the original parameter list.
+    - The kernel body may be reformatted but not logically changed.
+3. Rebuild a complete, compilable CUDA program using the fixed template provided below.
+4. Insert the extracted kernel EXACTLY into the placeholder {{KERNEL_CODE}}.
+5. Do NOT modify the test harness structure.
+   The test harness will load .bin files automatically and compare results.
+6. Use the following information:
+    - Binary directory: {{BIN_DIR}}
+    - Variables to load: {{VARNAMES}}
+    - Number of test cases: {{NUM_TESTS}}
+    - The test data files follow the format: data/{variable_name}_{case_idx}.bin, where {variable_name} is one of: {{VARNAMES}}
+
+The output **must be a COMPLETE, COMPILABLE .cu file**.
+It must:
+- include exactly one main() from the template,
+- print 'T' if all tests pass,
+- print 'F' otherwise,
+- compile with nvcc.
+
+Below is the CUDA test harness template you must strictly follow.
+You MUST follow the template's structure exactly. All sections in the template (read_bin_f32_dyn, CUDA kernel, correctness check, main(), variable loading, allocation, launch, correctness, free) must appear exactly once.
+Insert your kernel where {{KERNEL_CODE}} is located.
+
+```cpp
+#include <cuda_runtime.h>
+#include <iostream>
+#include <fstream>
+#include <vector>
+#include <string>
+#include <cmath>
+
+using std::vector;
+using std::string;
+
+#define TOLERANCE 1e-3f  // The TOLERANCE can be modified according to different tasks
+#define NUM_TESTS {{NUM_TESTS}}
+
+// --------------------------
+// read_bin_f32_dyn: read .bin function
+// --------------------------
+vector<float> read_bin_f32_dyn(const string& filename) {
+    FILE* fp = fopen(filename.c_str(), "rb");
+    if (!fp) {
+        std::cerr << "Cannot open bin file: " << filename << std::endl;
+        exit(1);
+    }
+
+    fseek(fp, 0, SEEK_END);
+    long fsize = ftell(fp);
+    rewind(fp);
+
+    size_t n_elem = fsize / sizeof(float);
+    vector<float> data(n_elem);
+    fread(data.data(), sizeof(float), n_elem, fp);
+    fclose(fp);
+    return data;
+}
+
+// --------------------------
+// CUDA kernel (inserted here)
+// --------------------------
+{{KERNEL_CODE}}
+
+// --------------------------
+// correctness check
+// --------------------------
+bool check_equal(const vector<float>& a, const vector<float>& b) {
+    if (a.size() != b.size()) return false;
+    for (size_t i = 0; i < a.size(); ++i) {
+        if (fabs(a[i] - b[i]) > TOLERANCE) return false;
+    }
+    return true;
+}
+
+// --------------------------
+// main test harness
+// --------------------------
+int main() {
+    bool all_pass = true;
+
+    for (int case_id = 1; case_id <= NUM_TESTS; ++case_id) {
+
+        // --------------------------
+        // load all variables
+        // --------------------------
+        {{VAR_LOADING_BLOCK}}
+
+        // --------------------------
+        // allocate GPU buffers
+        // --------------------------
+        {{ALLOCATE_BLOCK}}
+
+        // --------------------------
+        // kernel launch
+        // --------------------------
+        {{KERNEL_LAUNCH_BLOCK}}
+
+        // --------------------------
+        // correctness check
+        // --------------------------
+        {{CORRECTNESS_BLOCK}}
+
+        // --------------------------
+        // free GPU memory
+        // --------------------------
+        {{FREE_BLOCK}}
+    }
+
+    std::cout << (all_pass ? "T" : "F") << std::endl;
+    return 0;
+}
+
+Below is the original CUDA source code.
+Extract ONLY the kernel and rebuild the program using the above template.
+
+======== ORIGINAL CUDA CODE START ========
+{{CODE}}
+======== ORIGINAL CUDA CODE END ==========
+
+IMPORTANT RULES (strict):
+1. Remove or replace any undefined types (e.g., BlockMessage). If the structure definition is missing, expand it using equivalent primitive variables, e.g., int region1_index = blockMessage_region1_index[blockID];
+2. Do not generate any non-ASCII characters or Chinese comments. All output must be strictly ASCII-only.
+3. If the kernel uses external macros (such as BlockSizeInEval, MaxIntervalCount), detect them from the input code and insert default definitions at the top of the normalized code.
+4. Remove the inline qualifier from any __global__ kernel function. (inline is meaningless and should not be applied to __global__ functions.)
+5. Ensure that all parameters referenced inside the __global__ kernel are actually defined within the function scope. No undefined external dependencies are allowed.
+6. If many numerical mismatches occur in fabs(a[i] - b[i]), allow increasing the tolerance in the template (e.g., from 1e-3f to 1e-1f).
+7. Ignore warnings about unused local variables. These do not affect correctness or compilation and should not be treated as errors.