incorrect
9_BayesianLinearRegression.py
started
2026-05-11T20:04:01Z
correctness issue: Output mismatch
timingkernelbench.timing
kernel mean / median / std
—
ref mean / median / std
0.32 / — / 0.02 μs
SOL (speed-of-light)nsight_hardware_counters
roofline / occupancynsight_hardware_counters
warp stalls (top 5)nsight
reference PyTorch source (the prompt the agent sees)
import torch
import torch.nn as nn
class Model(nn.Module):
"""
Bayesian linear regression with a conjugate Gaussian prior.
Given features X and targets y, computes the posterior distribution
over weights: p(w | X, y) = N(mu_post, Sigma_post) in closed form
via the matrix inversion lemma / Cholesky decomposition.
"""
def __init__(self, input_dim, prior_precision=1.0, noise_precision=1.0):
super().__init__()
self.input_dim = input_dim
self.prior_precision = prior_precision
self.noise_precision = noise_precision
def forward(
self, X: torch.Tensor, y: torch.Tensor, X_test: torch.Tensor
) -> tuple:
"""
Args:
X: (B, N, D) – training features
y: (B, N, 1) – training targets
X_test: (B, M, D) – test features
Returns:
(pred_mean, pred_var):
pred_mean: (B, M, 1)
pred_var: (B, M)
"""
D = X.shape[-1]
eye = torch.eye(D, device=X.device, dtype=X.dtype).unsqueeze(0)
# Posterior precision: Lambda_post = alpha * I + beta * X^T X
Lambda_post = self.prior_precision * eye + self.noise_precision * (X.transpose(-1, -2) @ X)
L = torch.linalg.cholesky(Lambda_post)
# Posterior mean: mu_post = Lambda_post^{-1} * beta * X^T y
rhs = self.noise_precision * (X.transpose(-1, -2) @ y)
mu_post = torch.cholesky_solve(rhs, L)
# Predictive distribution
pred_mean = X_test @ mu_post
Sigma_post = torch.cholesky_inverse(L)
pred_var = (X_test @ Sigma_post * X_test).sum(-1) + 1.0 / self.noise_precision
return pred_mean, pred_var
input_dim = 16
num_train = 256
num_test = 64
batch_size = 16
def get_inputs():
p = popcorn_pri
X = torch.randn(p.jitter_int(batch_size), p.jitter_int(num_train), input_dim)
y = torch.randn(p.jitter_int(batch_size), p.jitter_int(num_train), 1)
X_test = torch.randn(p.jitter_int(batch_size), p.jitter_int(num_test), input_dim)
return [X, y, X_test]
def get_init_inputs():
return [input_dim]
turn 0338.7s turn · 15.9s LLM · 322.7s tools · tokens in=2,301 · out=6 · total=3,168
reasoning
The task is to replace the forward computation of this Bayesian linear regression model with a custom CUDA kernel. The model uses PyTorch operations like matrix multiplications, Cholesky decomposition...I need to replace the forward computation of a Bayesian linear regression model in PyTorch with a custom CUDA kernel. The model takes inputs X, y, and X_test, all tensors with batch dimension B, and performs computations like computing Lambda_post as alpha * I + beta * X^T X, then Cholesky decomposition on 16x16 matrices in a batch of 16, solving for mu_post, and predicting means and variances for test points.
To do this, I'll create a ModelNew class that mimics the original Model but uses a custom CUDA extension loaded via torch.utils.cpp_extension.load_inline.
call → get_gpu_specs
callget_gpu_specs
{}
result ← get_gpu_specs (ok)
okget_gpu_specs
GPU specs for NVIDIA H100 80GB HBM3:
total memory (runtime): 79.1 GB
spec entry: H100
GPU Architecture: Hopper
GPU Memory: 80GB
Memory Bandwidth: 3.35 TB/s
FP64 TFLOPS: 34
FP64 Tensor Core TFLOPS: 67
FP32 TFLOPS: 67
TF32 Tensor Core TFLOPS: 989 with sparsity
BFLOAT16 Tensore Core TFLOPS: 1979 with sparsity
FP16 Tensor Core TFLOPS: 1979 with sparsity
FP8 Tensor Core TFLOPS: 3958 with sparsity
INT8 Tensor Core TOPS: 3958 with sparsity
Register File Size: 64K 32-bit registers per SM
Maximum number of registers per thread: 255
Maximum number of thread blocks per SM: 32
Shared memory capacity per SM: 228 KB
Maximum shared memory per thread block: 227 KB
turn 1411.5s turn · 9.1s LLM · 402.4s tools · tokens in=3,401 · out=308 · total=4,139
reasoning
GPU is H100, good for CUDA.
I need to replace the forward computation of a PyTorch model with a custom CUDA kernel. The model takes the same inputs and produces the same outputs numerically. The forward method signature must match the original.
The internal reasoning shows the GPU is an H100, which is great for CUDA.
call → compile_kernel
callcompile_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = '''\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n#include <math.h>\n\n__global__ void bayesian_linear_kernel(\n const float* X, const float* y, const float* X_test,\n float* pred_mean, float* pred_var,\n int B, int N, int D, int M,\n float prior_precision, float noise_precision\n) {\n // For simplicity, process one batch per block, D=16 small\n int b = blockIdx.x;\n if (b >= B) return;\n\n // Shared memory for matrices\n extern __shared__ float smem[];\n float* Lambda = smem; // D*D\n float* Lchol = &smem[D*D]; // D*D\n // etc. This will be complex, placeholder for now\n // To make it work, use device functions or multiple kernels, but for starter\n // Actually, skip full impl here, use host side for now? No.\n}\n\ntorch::Tensor dummy(torch::Tensor X) {\n return X;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor dummy(torch::Tensor X);\n'''\n\nmodule = load_inline(\n name='bayes_cuda',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['dummy'],\n verbose=True\n)\nprint(\"Compiled successfully\")\n"
}
result ← compile_kernel (fail)
failcompile_kernel
compile_kernel FAILED: ModelNew class not found or syntax error prevented execution.
Compiled successfully
turn 2 · 1 tool call(s) · 515.4s turn · 20.9s LLM · tokens in=4,164 · out=1,690 · total=6,144 · 6,720 chars (click to expand)
turn 2515.4s turn · 20.9s LLM · 494.5s tools · tokens in=4,164 · out=1,690 · total=6,144
reasoning
It compiled the dummy, but I need the full ModelNew.
I need to replace the forward computation of the PyTorch model with a custom CUDA kernel. The ModelNew class must accept the same arguments as Model, have the same forward signature, and produce the same outputs numerically.
call → compile_kernel
callcompile_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = r'''\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n#include <vector>\n\n__device__ void cholesky_decomp(float* A, float* L, int D) {\n for (int i = 0; i < D; i++) {\n for (int j = 0; j <= i; j++) {\n float sum = 0.0f;\n for (int k = 0; k < j; k++) {\n sum += L[i*D + k] * L[j*D + k];\n }\n if (i == j) {\n L[i*D + j] = sqrtf(A[i*D + i] - sum);\n } else {\n L[i*D + j] = (A[i*D + j] - sum) / L[j*D + j];\n }\n }\n }\n}\n\n__device__ void cholesky_solve(const float* L, const float* b, float* x, int D) {\n float y[16];\n for (int i = 0; i < D; i++) {\n float sum = 0.0f;\n for (int j = 0; j < i; j++) sum += L[i*D + j] * y[j];\n y[i] = (b[i] - sum) / L[i*D + i];\n }\n for (int i = D-1; i >= 0; i--) {\n float sum = 0.0f;\n for (int j = i+1; j < D; j++) sum += L[j*D + i] * x[j];\n x[i] = (y[i] - sum) / L[i*D + i];\n }\n}\n\n__global__ void compute_bayes(\n const float* __restrict__ X,\n const float* __restrict__ y,\n const float* __restrict__ X_test,\n float* __restrict__ pred_mean,\n float* __restrict__ pred_var,\n const int B, const int N, const int D, const int M,\n const float alpha, const float beta\n) {\n int b = blockIdx.x;\n if (b >= B) return;\n\n extern __shared__ float sm[];\n\n float* Lambda = sm; // D*D\n float* L = &sm[D*D]; // D*D\n float* rhs = &sm[2*D*D]; // D\n float* mu = &sm[2*D*D + D]; // D\n float* Sigma = &sm[2*D*D + 2*D]; // D*D\n float* temp = &sm[3*D*D + 2*D]; // D*D for temp\n\n // Compute X^T X : Lambda = beta * X^T @ X\n for (int i=0; i<D; i++) {\n for (int j=0; j<D; j++) {\n float s = 0.0f;\n for (int k=0; k<N; k++) {\n float xi = X[b*N*D + k*D + i];\n float xj = X[b*N*D + k*D + j];\n s += xi * xj;\n }\n Lambda[i*D + j] = beta * s;\n }\n }\n\n // add alpha * I\n for (int i=0; i<D; i++) Lambda[i*D + i] += alpha;\n\n // Cholesky\n cholesky_decomp(Lambda, L, D);\n\n // rhs = beta * X^T y\n for (int i=0; i<D; i++) {\n float s = 0.0f;\n for (int k=0; k<N; k++) {\n s += X[b*N*D + k*D + i] * y[b*N + k];\n }\n rhs[i] = beta * s;\n }\n\n // mu = L^{-1} rhs\n cholesky_solve(L, rhs, mu, D);\n\n // pred_mean = X_test @ mu\n for (int m=0; m<M; m++) {\n float s = 0.0f;\n for (int d=0; d<D; d++) {\n s += X_test[b*M*D + m*D + d] * mu[d];\n }\n pred_mean[b*M + m] = s;\n }\n\n // Sigma = cholesky_inverse(L) i.e. inv(L^T L) = inv(L) inv(L^T)\n // First compute inv(L) by solving L Z = I\n for (int col=0; col<D; col++) {\n float e[16] = {0};\n e[col] = 1.0f;\n cholesky_solve(L, e, &Sigma[col*D], D); // actually this is for inv(L^T)? wait careful\n // Note: to fix later for correctness\n }\n\n // For simplicity, pred_var approx, but to make correct, need fix Sigma\n // pred_var = diag(X_test @ Sigma @ X_test^T) + 1/beta\n for (int m=0; m<M; m++) {\n float s = 0.0f;\n for (int d=0; d<D; d++) {\n float tmp = 0.0f;\n for (int k=0; k<D; k++) {\n tmp += X_test[b*M*D + m*D + k] * Sigma[k*D + d]; // wrong indexing\n }\n s += tmp * X_test[b*M*D + m*D + d];\n }\n pred_var[b*M + m] = s + 1.0f / beta;\n }\n}\n\nstd::vector<torch::Tensor> bayes_forward(\n torch::Tensor X, torch::Tensor y, torch::Tensor X_test,\n float prior_precision, float noise_precision) {\n int B = X.size(0);\n int N = X.size(1);\n int D = X.size(2);\n int M = X_test.size(1);\n auto options = torch::TensorOptions().device(X.device()).dtype(X.dtype());\n auto pred_mean = torch::empty({B, M, 1}, options);\n auto pred_var = torch::empty({B, M}, options);\n int smem_size = 4*D*D + 3*D * sizeof(float);\n compute_bayes<<<B, 1, smem_size>>>(\n X.data_ptr<float>(), y.data_ptr<float>(), X_test.data_ptr<float>(),\n pred_mean.data_ptr<float>(), pred_var.data_ptr<float>(),\n B, N, D, M, prior_precision, noise_precision\n );\n // reshape pred_mean to (B,M,1)\n pred_mean = pred_mean.view({B, M, 1});\n return {pred_mean, pred_var};\n}\n'''\n\ncpp_source = '''\nstd::vector<torch::Tensor> bayes_forward(torch::Tensor X, torch::Tensor y, torch::Tensor X_test, float prior_precision, float noise_precision);\n'''\n\nmodule = load_inline(\n name='bayes_linear',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['bayes_forward'],\n verbose=True,\n extra_cuda_cflags=['-O3']\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, input_dim, prior_precision=1.0, noise_precision=1.0):\n super().__init__()\n self.input_dim = input_dim\n self.prior_precision = prior_precision\n self.noise_precision = noise_precision\n self.module = module\n\n def forward(self, X, y, X_test):\n outs = self.module.bayes_forward(X, y, X_test, self.prior_precision, self.noise_precision)\n return outs[0], outs[1]\nprint(\"Setup done\")\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 3407.6s turn · 10.8s LLM · 396.8s tools · tokens in=6,161 · out=1,519 · total=7,690
reasoning
Compiled, now check correctness with run_correctness.
call → run_correctness
callrun_correctness
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = r'''\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n#include <vector>\n\n__device__ void cholesky_decomp(float* A, float* L, int D) {\n for (int i = 0; i < D; i++) {\n for (int j = 0; j <= i; j++) {\n float sum = 0.0f;\n for (int k = 0; k < j; k++) {\n sum += L[i*D + k] * L[j*D + k];\n }\n if (i == j) {\n L[i*D + j] = sqrtf(A[i*D + i] - sum);\n } else {\n L[i*D + j] = (A[i*D + j] - sum) / L[j*D + j];\n }\n }\n }\n}\n\n__device__ void cholesky_solve(const float* L, const float* b, float* x, int D) {\n float y[16];\n for (int i = 0; i < D; i++) {\n float sum = 0.0f;\n for (int j = 0; j < i; j++) sum += L[i*D + j] * y[j];\n y[i] = (b[i] - sum) / L[i*D + i];\n }\n for (int i = D-1; i >= 0; i--) {\n float sum = 0.0f;\n for (int j = i+1; j < D; j++) sum += L[j*D + i] * x[j];\n x[i] = (y[i] - sum) / L[i*D + i];\n }\n}\n\n__global__ void compute_bayes(\n const float* __restrict__ X,\n const float* __restrict__ y,\n const float* __restrict__ X_test,\n float* __restrict__ pred_mean,\n float* __restrict__ pred_var,\n const int B, const int N, const int D, const int M,\n const float alpha, const float beta\n) {\n int b = blockIdx.x;\n if (b >= B) return;\n\n extern __shared__ float sm[];\n\n float* Lambda = sm; // D*D\n float* L = &sm[D*D]; // D*D\n float* rhs = &sm[2*D*D]; // D\n float* mu = &sm[2*D*D + D]; // D\n float* Sigma = &sm[2*D*D + 2*D]; // D*D\n float* temp = &sm[3*D*D + 2*D]; // D*D for temp\n\n // Compute X^T X : Lambda = beta * X^T @ X\n for (int i=0; i<D; i++) {\n for (int j=0; j<D; j++) {\n float s = 0.0f;\n for (int k=0; k<N; k++) {\n float xi = X[b*N*D + k*D + i];\n float xj = X[b*N*D + k*D + j];\n s += xi * xj;\n }\n Lambda[i*D + j] = beta * s;\n }\n }\n\n // add alpha * I\n for (int i=0; i<D; i++) Lambda[i*D + i] += alpha;\n\n // Cholesky\n cholesky_decomp(Lambda, L, D);\n\n // rhs = beta * X^T y\n for (int i=0; i<D; i++) {\n float s = 0.0f;\n for (int k=0; k<N; k++) {\n s += X[b*N*D + k*D + i] * y[b*N + k];\n }\n rhs[i] = beta * s;\n }\n\n // mu = L^{-1} rhs\n cholesky_solve(L, rhs, mu, D);\n\n // pred_mean = X_test @ mu\n for (int m=0; m<M; m++) {\n float s = 0.0f;\n for (int d=0; d<D; d++) {\n s += X_test[b*M*D + m*D + d] * mu[d];\n }\n pred_mean[b*M + m] = s;\n }\n\n // Sigma approx wrong, fix later\n for (int i=0; i<D*D; i++) Sigma[i] = 0.0f;\n for (int i=0; i<D; i++) Sigma[i*D + i] = 1.0f; // dummy\n\n for (int m=0; m<M; m++) {\n float s = 0.0f;\n pred_var[b*M + m] = s + 1.0f / beta;\n }\n}\n\nstd::vector<torch::Tensor> bayes_forward(\n torch::Tensor X, torch::Tensor y, torch::Tensor X_test,\n float prior_precision, float noise_precision) {\n int B = X.size(0);\n int N = X.size(1);\n int D = X.size(2);\n int M = X_test.size(1);\n auto options = torch::TensorOptions().device(X.device()).dtype(X.dtype());\n auto pred_mean = torch::empty({B, M}, options);\n auto pred_var = torch::empty({B, M}, options);\n int smem_size = (4*D*D + 3*D) * sizeof(float);\n compute_bayes<<<B, 1, smem_size>>>(\n X.data_ptr<float>(), y.data_ptr<float>(), X_test.data_ptr<float>(),\n pred_mean.data_ptr<float>(), pred_var.data_ptr<float>(),\n B, N, D, M, prior_precision, noise_precision\n );\n pred_mean = pred_mean.view({B, M, 1});\n return {pred_mean, pred_var};\n}\n'''\n\ncpp_source = '''\nstd::vector<torch::Tensor> bayes_forward(torch::Tensor X, torch::Tensor y, torch::Tensor X_test, float prior_precision, float noise_precision);\n'''\n\nmodule = load_inline(\n name='bayes_linear',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['bayes_forward'],\n verbose=True,\n extra_cuda_cflags=['-O3']\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, input_dim, prior_precision=1.0, noise_precision=1.0):\n super().__init__()\n self.input_dim = input_dim\n self.prior_precision = prior_precision\n self.noise_precision = noise_precision\n self.module = module\n\n def forward(self, X, y, X_test):\n outs = self.module.bayes_forward(X, y, X_test, self.prior_precision, self.noise_precision)\n return outs[0], outs[1]\n'''\n\n# To test, we would run but since no input here, assume\nprint(\"Code ready\")"
}
result ← run_correctness (fail)
failrun_correctness
run_correctness FAILED: kernel did not compile.
Syntax error in custom generated code or ModelNew not found
turn 4 · 1 tool call(s) · 175.6s turn · 21.4s LLM · tokens in=7,718 · out=1,598 · total=9,426 · 6,138 chars (click to expand)
turn 4175.6s turn · 21.4s LLM · 154.1s tools · tokens in=7,718 · out=1,598 · total=9,426
reasoning
The code has syntax issues in the run, because I put comments.
call → compile_kernel
callcompile_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = r'''\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n#include <vector>\n\n__device__ void cholesky_decomp(float* A, float* L, int D) {\n for (int i = 0; i < D; i++) {\n for (int j = 0; j <= i; j++) {\n float sum = 0.0f;\n for (int k = 0; k < j; k++) {\n sum += L[i*D + k] * L[j*D + k];\n }\n if (i == j) {\n L[i*D + j] = sqrtf(A[i*D + i] - sum);\n } else {\n L[i*D + j] = (A[i*D + j] - sum) / L[j*D + j];\n }\n }\n }\n}\n\n__device__ void forward_sub(const float* L, const float* b, float* y, int D) {\n for (int i = 0; i < D; i++) {\n float sum = 0.0f;\n for (int j = 0; j < i; j++) sum += L[i*D + j] * y[j];\n y[i] = (b[i] - sum) / L[i*D + i];\n }\n}\n\n__device__ void back_sub(const float* L, const float* y, float* x, int D) {\n for (int i = D-1; i >= 0; i--) {\n float sum = 0.0f;\n for (int j = i+1; j < D; j++) sum += L[j*D + i] * x[j];\n x[i] = (y[i] - sum) / L[i*D + i];\n }\n}\n\n__global__ void compute_bayes(\n const float* __restrict__ X,\n const float* __restrict__ y,\n const float* __restrict__ X_test,\n float* __restrict__ pred_mean,\n float* __restrict__ pred_var,\n const int B, const int N, const int D, const int M,\n const float alpha, const float beta\n) {\n int b = blockIdx.x;\n if (b >= B) return;\n\n extern __shared__ float sm[];\n\n float* Lambda = sm;\n float* L = &sm[D*D];\n float* invL = &sm[2*D*D];\n float* rhs = &sm[3*D*D];\n float* mu = &sm[3*D*D + D];\n float* col = &sm[3*D*D + 2*D];\n\n // X^T X\n for (int i = 0; i < D; i++) {\n for (int j = 0; j < D; j++) {\n float s = 0.0f;\n for (int k = 0; k < N; k++) {\n s += X[(b*N + k)*D + i] * X[(b*N + k)*D + j];\n }\n Lambda[i*D + j] = beta * s;\n }\n }\n for (int i = 0; i < D; i++) Lambda[i*D + i] += alpha;\n\n cholesky_decomp(Lambda, L, D);\n\n // rhs = beta * X^T y\n for (int i = 0; i < D; i++) {\n float s = 0.0f;\n for (int k = 0; k < N; k++) {\n s += X[(b*N + k)*D + i] * y[b*N + k];\n }\n rhs[i] = beta * s;\n }\n\n forward_sub(L, rhs, col, D);\n back_sub(L, col, mu, D);\n\n // pred_mean\n for (int m = 0; m < M; m++) {\n float s = 0.0f;\n for (int d = 0; d < D; d++) {\n s += X_test[(b*M + m)*D + d] * mu[d];\n }\n pred_mean[b*M + m] = s;\n }\n\n // inv(L) by columns\n for (int c = 0; c < D; c++) {\n float e[16] = {};\n e[c] = 1.0f;\n forward_sub(L, e, col, D);\n back_sub(L, col, &invL[c*D], D);\n }\n\n // Sigma = invL @ invL^T , but for var we compute X_test @ Sigma @ X_test^T diag\n for (int m = 0; m < M; m++) {\n float s = 0.0f;\n for (int d = 0; d < D; d++) {\n float tmp = 0.0f;\n for (int k = 0; k < D; k++) {\n tmp += X_test[(b*M + m)*D + k] * invL[k*D + d];\n }\n float row = 0.0f;\n for (int k = 0; k < D; k++) {\n row += tmp * invL[d*D + k]; wait fix\n }\n }\n pred_var[b*M + m] = 1.0f / beta;\n }\n}\n\nstd::vector<torch::Tensor> bayes_forward(torch::Tensor X, torch::Tensor y, torch::Tensor X_test, float prior_precision, float noise_precision) {\n int B = X.size(0);\n int N = X.size(1);\n int D = X.size(2);\n int M = X_test.size(1);\n auto options = torch::TensorOptions().device(X.device()).dtype(X.dtype());\n auto pred_mean = torch::empty({B, M}, options);\n auto pred_var = torch::empty({B, M}, options);\n size_t smem = (4*D*D + 3*D) * sizeof(float);\n compute_bayes<<<B, 1, smem>>>(\n X.data_ptr<float>(), y.data_ptr<float>(), X_test.data_ptr<float>(),\n pred_mean.data_ptr<float>(), pred_var.data_ptr<float>(),\n B, N, D, M, prior_precision, noise_precision\n );\n return {pred_mean.view({B, M, 1}), pred_var};\n}\n'''\n\ncpp_source = '''\nstd::vector<torch::Tensor> bayes_forward(torch::Tensor X, torch::Tensor y, torch::Tensor X_test, float prior_precision, float noise_precision);\n'''\n\nmodule = load_inline(\n name='bayes_linear',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['bayes_forward'],\n verbose=True,\n extra_cuda_cflags=['-O3']\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, input_dim, prior_precision=1.0, noise_precision=1.0):\n super().__init__()\n self.input_dim = input_dim\n self.prior_precision = prior_precision\n self.noise_precision = noise_precision\n self.module = module\n\n def forward(self, X, y, X_test):\n outs = self.module.bayes_forward(X, y, X_test, self.prior_precision, self.noise_precision)\n return outs[0], outs[1]\n"
}
result ← compile_kernel (fail)
failcompile_kernel
compile_kernel FAILED: RuntimeError.
RuntimeError: Error building extension 'bayes_linear_v1'
turn 5 · FINAL · 1 tool call(s) · 438.5s turn · 16.9s LLM · tokens in=9,453 · out=1,608 · total=11,204 · 6,283 chars (click to expand)
turn 5438.5s turn · 16.9s LLM · 421.6s tools · tokens in=9,453 · out=1,608 · total=11,204 · FINAL
reasoning
There's a syntax error in the code: "wait fix" comment in CUDA code.
call → submit_kernel
callsubmit_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = r'''\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n#include <vector>\n\n__device__ void cholesky_decomp(float* A, float* L, int D) {\n for (int i = 0; i < D; i++) {\n for (int j = 0; j <= i; j++) {\n float sum = 0.0f;\n for (int k = 0; k < j; k++) {\n sum += L[i*D + k] * L[j*D + k];\n }\n if (i == j) {\n L[i*D + j] = sqrtf(A[i*D + i] - sum);\n } else {\n L[i*D + j] = (A[i*D + j] - sum) / L[j*D + j];\n }\n }\n }\n}\n\n__device__ void forward_sub(const float* L, const float* b, float* y, int D) {\n for (int i = 0; i < D; i++) {\n float sum = 0.0f;\n for (int j = 0; j < i; j++) sum += L[i*D + j] * y[j];\n y[i] = (b[i] - sum) / L[i*D + i];\n }\n}\n\n__device__ void back_sub(const float* L, const float* y, float* x, int D) {\n for (int i = D-1; i >= 0; i--) {\n float sum = 0.0f;\n for (int j = i+1; j < D; j++) sum += L[j*D + i] * x[j];\n x[i] = (y[i] - sum) / L[i*D + i];\n }\n}\n\n__global__ void compute_bayes(\n const float* __restrict__ X,\n const float* __restrict__ y,\n const float* __restrict__ X_test,\n float* __restrict__ pred_mean,\n float* __restrict__ pred_var,\n const int B, const int N, const int D, const int M,\n const float alpha, const float beta\n) {\n int b = blockIdx.x;\n if (b >= B) return;\n\n extern __shared__ float sm[];\n\n float* Lambda = sm;\n float* L = &sm[D*D];\n float* invL = &sm[2*D*D];\n float* rhs = &sm[3*D*D];\n float* mu = &sm[3*D*D + D];\n float* col = &sm[3*D*D + 2*D];\n float* tmpvec = &sm[3*D*D + 3*D];\n\n // X^T X\n for (int i = 0; i < D; i++) {\n for (int j = 0; j < D; j++) {\n float s = 0.0f;\n for (int k = 0; k < N; k++) {\n s += X[(b*N + k)*D + i] * X[(b*N + k)*D + j];\n }\n Lambda[i*D + j] = beta * s;\n }\n }\n for (int i = 0; i < D; i++) Lambda[i*D + i] += alpha;\n\n cholesky_decomp(Lambda, L, D);\n\n // rhs\n for (int i = 0; i < D; i++) {\n float s = 0.0f;\n for (int k = 0; k < N; k++) {\n s += X[(b*N + k)*D + i] * y[b*N + k];\n }\n rhs[i] = beta * s;\n }\n\n forward_sub(L, rhs, col, D);\n back_sub(L, col, mu, D);\n\n // pred_mean\n for (int m = 0; m < M; m++) {\n float s = 0.0f;\n for (int d = 0; d < D; d++) {\n s += X_test[(b*M + m)*D + d] * mu[d];\n }\n pred_mean[b*M + m] = s;\n }\n\n // inv(L) columns , note L lower, solve gives inv(L)\n for (int c = 0; c < D; c++) {\n float e[32] = {0.0f};\n e[c] = 1.0f;\n forward_sub(L, e, col, D);\n for (int r=0; r<D; r++) invL[r*D + c] = col[r];\n }\n\n // pred_var: || Xrow @ invL || ^2 + 1/beta\n for (int m = 0; m < M; m++) {\n for (int d=0; d<D; d++) {\n float t = 0.0f;\n for (int k=0; k<D; k++) t += X_test[(b*M + m)*D + k] * invL[k*D + d];\n tmpvec[d] = t;\n }\n float s = 0.0f;\n for (int d=0; d<D; d++) s += tmpvec[d] * tmpvec[d];\n pred_var[b*M + m] = s + 1.0f / beta;\n }\n}\n\nstd::vector<torch::Tensor> bayes_forward(torch::Tensor X, torch::Tensor y, torch::Tensor X_test, float prior_precision, float noise_precision) {\n int B = X.size(0);\n int N = X.size(1);\n int D = X.size(2);\n int M = X_test.size(1);\n auto options = torch::TensorOptions().device(X.device()).dtype(X.dtype());\n auto pred_mean = torch::empty({B, M}, options);\n auto pred_var = torch::empty({B, M}, options);\n size_t smem = (4*D*D + 4*D) * sizeof(float);\n compute_bayes<<<B, 1, smem>>>(\n X.data_ptr<float>(), y.data_ptr<float>(), X_test.data_ptr<float>(),\n pred_mean.data_ptr<float>(), pred_var.data_ptr<float>(),\n B, N, D, M, prior_precision, noise_precision\n );\n return {pred_mean.view({B, M, 1}), pred_var};\n}\n'''\n\ncpp_source = '''\nstd::vector<torch::Tensor> bayes_forward(torch::Tensor X, torch::Tensor y, torch::Tensor X_test, float prior_precision, float noise_precision);\n'''\n\nmodule = load_inline(\n name='bayes_linear',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['bayes_forward'],\n verbose=True,\n extra_cuda_cflags=['-O3']\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, input_dim, prior_precision=1.0, noise_precision=1.0):\n super().__init__()\n self.input_dim = input_dim\n self.prior_precision = prior_precision\n self.noise_precision = noise_precision\n self.module = module\n\n def forward(self, X, y, X_test):\n outs = self.module.bayes_forward(X, y, X_test, self.prior_precision, self.noise_precision)\n return outs[0], outs[1]\n"
}
result ← submit_kernel (fail)
failsubmit_kernel
submit_kernel FAILED: correctness check did not pass ((0 / 3) trials).
correctness_issue: Output mismatch
max_difference: ['0.009262', '0.005845', '0.010308']
avg_difference: ['0.001189', '0.001102', '0.001124']
final submitted kernel (162 lines, 4,832 chars) — full source, untruncated
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline
cuda_source = r'''
#include <torch/extension.h>
#include <cuda_runtime.h>
#include <vector>
__device__ void cholesky_decomp(float* A, float* L, int D) {
for (int i = 0; i < D; i++) {
for (int j = 0; j <= i; j++) {
float sum = 0.0f;
for (int k = 0; k < j; k++) {
sum += L[i*D + k] * L[j*D + k];
}
if (i == j) {
L[i*D + j] = sqrtf(A[i*D + i] - sum);
} else {
L[i*D + j] = (A[i*D + j] - sum) / L[j*D + j];
}
}
}
}
__device__ void forward_sub(const float* L, const float* b, float* y, int D) {
for (int i = 0; i < D; i++) {
float sum = 0.0f;
for (int j = 0; j < i; j++) sum += L[i*D + j] * y[j];
y[i] = (b[i] - sum) / L[i*D + i];
}
}
__device__ void back_sub(const float* L, const float* y, float* x, int D) {
for (int i = D-1; i >= 0; i--) {
float sum = 0.0f;
for (int j = i+1; j < D; j++) sum += L[j*D + i] * x[j];
x[i] = (y[i] - sum) / L[i*D + i];
}
}
__global__ void compute_bayes(
const float* __restrict__ X,
const float* __restrict__ y,
const float* __restrict__ X_test,
float* __restrict__ pred_mean,
float* __restrict__ pred_var,
const int B, const int N, const int D, const int M,
const float alpha, const float beta
) {
int b = blockIdx.x;
if (b >= B) return;
extern __shared__ float sm[];
float* Lambda = sm;
float* L = &sm[D*D];
float* invL = &sm[2*D*D];
float* rhs = &sm[3*D*D];
float* mu = &sm[3*D*D + D];
float* col = &sm[3*D*D + 2*D];
float* tmpvec = &sm[3*D*D + 3*D];
// X^T X
for (int i = 0; i < D; i++) {
for (int j = 0; j < D; j++) {
float s = 0.0f;
for (int k = 0; k < N; k++) {
s += X[(b*N + k)*D + i] * X[(b*N + k)*D + j];
}
Lambda[i*D + j] = beta * s;
}
}
for (int i = 0; i < D; i++) Lambda[i*D + i] += alpha;
cholesky_decomp(Lambda, L, D);
// rhs
for (int i = 0; i < D; i++) {
float s = 0.0f;
for (int k = 0; k < N; k++) {
s += X[(b*N + k)*D + i] * y[b*N + k];
}
rhs[i] = beta * s;
}
forward_sub(L, rhs, col, D);
back_sub(L, col, mu, D);
// pred_mean
for (int m = 0; m < M; m++) {
float s = 0.0f;
for (int d = 0; d < D; d++) {
s += X_test[(b*M + m)*D + d] * mu[d];
}
pred_mean[b*M + m] = s;
}
// inv(L) columns , note L lower, solve gives inv(L)
for (int c = 0; c < D; c++) {
float e[32] = {0.0f};
e[c] = 1.0f;
forward_sub(L, e, col, D);
for (int r=0; r<D; r++) invL[r*D + c] = col[r];
}
// pred_var: || Xrow @ invL || ^2 + 1/beta
for (int m = 0; m < M; m++) {
for (int d=0; d<D; d++) {
float t = 0.0f;
for (int k=0; k<D; k++) t += X_test[(b*M + m)*D + k] * invL[k*D + d];
tmpvec[d] = t;
}
float s = 0.0f;
for (int d=0; d<D; d++) s += tmpvec[d] * tmpvec[d];
pred_var[b*M + m] = s + 1.0f / beta;
}
}
std::vector<torch::Tensor> bayes_forward(torch::Tensor X, torch::Tensor y, torch::Tensor X_test, float prior_precision, float noise_precision) {
int B = X.size(0);
int N = X.size(1);
int D = X.size(2);
int M = X_test.size(1);
auto options = torch::TensorOptions().device(X.device()).dtype(X.dtype());
auto pred_mean = torch::empty({B, M}, options);
auto pred_var = torch::empty({B, M}, options);
size_t smem = (4*D*D + 4*D) * sizeof(float);
compute_bayes<<<B, 1, smem>>>(
X.data_ptr<float>(), y.data_ptr<float>(), X_test.data_ptr<float>(),
pred_mean.data_ptr<float>(), pred_var.data_ptr<float>(),
B, N, D, M, prior_precision, noise_precision
);
return {pred_mean.view({B, M, 1}), pred_var};
}
'''
cpp_source = '''
std::vector<torch::Tensor> bayes_forward(torch::Tensor X, torch::Tensor y, torch::Tensor X_test, float prior_precision, float noise_precision);
'''
module = load_inline(
name='bayes_linear',
cpp_sources=[cpp_source],
cuda_sources=[cuda_source],
functions=['bayes_forward'],
verbose=True,
extra_cuda_cflags=['-O3']
)
class ModelNew(nn.Module):
def __init__(self, input_dim, prior_precision=1.0, noise_precision=1.0):
super().__init__()
self.input_dim = input_dim
self.prior_precision = prior_precision
self.noise_precision = noise_precision
self.module = module
def forward(self, X, y, X_test):
outs = self.module.bayes_forward(X, y, X_test, self.prior_precision, self.noise_precision)
return outs[0], outs[1]