Server2

.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"]

README.md

@@ -1,5 +1,5 @@
 ---
-title: KernelPilot V1 Server2
+title: KernelPliot-V1 Server2
 emoji: 🔥
 colorFrom: indigo
 colorTo: pink

app.py

@@ -0,0 +1,92 @@
+import os, tempfile, time
+import gradio as gr
+from tool.testv3 import run_autotune_pipeline
+
+# ---------- Core callback ----------
+def generate_kernel(text_input, n_iters, progress=gr.Progress()):
+    """
+    text_input : string from textbox (NL description or base CUDA code)
+    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.", "", None
+        
+    td = tempfile.mkdtemp(prefix="auto_")
+    src_path = os.path.join(td, f"input_{int(time.time())}.txt")
+    with open(src_path, "w") as f:
+        f.write(text_input)
+
+    best_code = ""
+    for info in run_autotune_pipeline(src_path, 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"]
+
+
+    # last yield enables the 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  
+Enter a natural‑language description,  
+then click **Generate** to obtain the kernel function."""
+    )
+
+    with gr.Row():
+        txt_input = gr.Textbox(
+            label="📝 Input",
+            lines=10,
+            placeholder="Describe the kernel",
+            scale=3
+        )
+        level = gr.Number(
+            label="Optimization Level",
+            minimum=1,
+            maximum=5,
+            value=5,
+            step=1,
+            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, 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/input.txt

@@ -0,0 +1 @@
+Write a program that transposes a matrix of 32-bit floating point numbers on a GPU. The transpose of a matrix switches its rows and columns. Given a matrix $A$ of dimensions rows $\times$ cols, the transpose $A^T$ will have dimensions cols $\times$ rows. All matrices are stored in row-major format.

tool/prompt_a.txt

@@ -0,0 +1,7 @@
+The task is to write a CUDA kernel function on GPU, we have the input described as below:
+[input.txt]
+And we have also generated the benchmark code for this task: 
+[benchmark 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 five 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 five 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_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 five 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 five 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 five 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 five 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 five 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 five 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 five 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 five 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/sample.txt

@@ -0,0 +1,140 @@
+Given the following task description:
+[input.txt]
+Write a complete CUDA program (.cu file) that solves this task using a basic and correct algorithm. The implementation should include a kernel function and the main function that tests it.
+There are 5 sets of binary input data. The main function should test all five datasets. If all of them pass, the program should print "T", otherwise print "F". The output must strictly be either "T" or "F". Do not write any extra output.
+Here is a reference style for the structure of the .cu file:
+```cu
+#include <iostream>
+#include <cuda_runtime.h>
+#include <device_launch_parameters.h>
+#include <cmath>
+#include <fstream>
+#include <vector>
+
+#define C 10    
+
+__global__ void cross_entropy_kernel(const float* logits, const int* labels, float* loss_sum, int N) {
+    int j = blockIdx.x * blockDim.x + threadIdx.x; 
+    if (j >= N) return;
+
+    float max_logit = -1e20f;
+    for (int k = 0; k < C; ++k) {
+        float z = logits[j * C + k];
+        if (z > max_logit) max_logit = z;
+    }
+
+    float exp_sum = 0.0f;
+    for (int k = 0; k < C; ++k) {
+        exp_sum += expf(logits[j * C + k] - max_logit);
+    }
+
+    float log_softmax_sum = logf(exp_sum);
+    float loss_j = log_softmax_sum + max_logit - logits[j * C + labels[j]];
+    atomicAdd(loss_sum, loss_j / N); 
+}
+
+void read_binary_float(const std::string& filename, float* data, size_t size) {
+    std::ifstream in(filename, std::ios::binary);
+    if (!in) {
+        std::cerr << "Cannot open: " << filename << std::endl;
+        exit(1);
+    }
+    in.read(reinterpret_cast<char*>(data), size * sizeof(float));
+    in.close();
+}
+
+void read_binary_int(const std::string& filename, int* data, size_t size) {
+    std::ifstream in(filename, std::ios::binary);
+    if (!in) {
+        std::cerr << "Cannot open: " << filename << std::endl;
+        exit(1);
+    }
+    in.read(reinterpret_cast<char*>(data), size * sizeof(int));
+    in.close();
+}
+
+// test
+bool compare_scalar(float a, float b, float tol = 1e-2f) {
+    return fabs(a - b) < tol;
+}
+
+int main() {
+    std::vector<size_t> Ns = {1<<14, 1<<16, 1<<18, 1<<20, 1<<22};
+    bool all_match = true;
+
+    for (int idx = 0; idx < Ns.size(); ++idx) {
+        size_t N = Ns[idx];
+        size_t logits_size = N * C;
+        size_t logits_bytes = logits_size * sizeof(float);
+        size_t labels_bytes = N * sizeof(int);
+
+        // test
+        std::string logits_file = "data/ce_logits_" + std::to_string(idx + 1) + ".bin";
+        std::string labels_file = "data/ce_labels_" + std::to_string(idx + 1) + ".bin";
+        std::string ref_file    = "data/ce_ref_"    + std::to_string(idx + 1) + ".bin";
+
+        float* h_logits = (float*)malloc(logits_bytes);
+        int* h_labels   = (int*)malloc(labels_bytes);
+        float h_ref;
+
+        read_binary_float(logits_file, h_logits, logits_size);
+        read_binary_int(labels_file, h_labels, N);
+        read_binary_float(ref_file, &h_ref, 1);
+
+        float *d_logits, *d_loss;
+        int* d_labels;
+        cudaMalloc(&d_logits, logits_bytes);
+        cudaMalloc(&d_labels, labels_bytes);
+        cudaMalloc(&d_loss, sizeof(float));
+        cudaMemcpy(d_logits, h_logits, logits_bytes, cudaMemcpyHostToDevice);
+        cudaMemcpy(d_labels, h_labels, labels_bytes, cudaMemcpyHostToDevice);
+        cudaMemset(d_loss, 0, sizeof(float));
+
+        int threads = 256;
+        int blocks = (N + threads - 1) / threads;
+        cross_entropy_kernel<<<blocks, threads>>>(d_logits, d_labels, d_loss, N);
+
+        float h_loss;
+        cudaMemcpy(&h_loss, d_loss, sizeof(float), cudaMemcpyDeviceToHost);
+
+        if (!compare_scalar(h_loss, h_ref)) {
+            std::cout << "F" << std::endl;
+            all_match = false;
+            cudaFree(d_logits); cudaFree(d_labels); cudaFree(d_loss);
+            free(h_logits); free(h_labels);
+            break;
+        }
+
+        cudaFree(d_logits); cudaFree(d_labels); cudaFree(d_loss);
+        free(h_logits); free(h_labels);
+    }
+
+    if (all_match) std::cout << "T" << std::endl;
+    return 0;
+}
+```
+You also need to write the Python script gen_test_data.py to generate the test binary files. Use the following as a style reference. Make sure all binary files are stored under a folder named data/. Avoid generating extremely large data. Suggested sizes are shown below:
+```py
+import numpy as np
+import os
+
+np.random.seed(30)
+
+C = 10
+sizes = [2**14, 2**16, 2**18, 2**20, 2**22]
+
+for idx, N in enumerate(sizes):
+    logits = (np.random.randn(N, C) * 3).astype(np.float32)
+    labels = np.random.randint(0, C, size=N, dtype=np.int32)
+
+    logits_max = logits.max(axis=1, keepdims=True)
+    exp_logits = np.exp(logits - logits_max)
+    log_sum_exp = np.log(exp_logits.sum(axis=1))
+    losses = log_sum_exp + logits_max[:, 0] - logits[np.arange(N), labels]
+    avg_loss = losses.mean().astype(np.float32)
+
+    logits.tofile(f"ce_logits_{idx+1}.bin")
+    labels.tofile(f"ce_labels_{idx+1}.bin")
+    np.array([avg_loss], dtype=np.float32).tofile(f"ce_ref_{idx+1}.bin")
+```
+You should generate these two code blocks at once.

tool/testv3 copy.so

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:9911ef353f67e6f1bd1ea5c7fea5e9bea863983dfbe8945069079343a25a4be3
+size 487120

tool/testv3.so

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:9911ef353f67e6f1bd1ea5c7fea5e9bea863983dfbe8945069079343a25a4be3
+size 487120