correct
23_GibbsSamplingStep.py
started
2026-05-11T22:04:21Z
timingkernelbench.timing
kernel mean / median / std
0.62 / — / 0.01 μs
ref mean / median / std
3.51 / — / 0.36 μs
numerical precisioneval_kernel_against_ref
gpu memoryextended_metrics.measure_memory
kernel launches / fusiontorch.profiler
fusion ratio (ref / kernel)
4.68x
SOL (speed-of-light)nsight_hardware_counters
roofline / occupancynsight_hardware_counters
dram bandwidth
334.13 GB/s
warp stalls (top 5)nsight
reference PyTorch source (the prompt the agent sees)
import torch
import torch.nn as nn
class Model(nn.Module):
"""
One full sweep of Gibbs sampling for a multivariate Gaussian with
known precision matrix. Each coordinate is sampled from its full
conditional distribution while holding all other coordinates fixed.
Benchmarks the sequential-scan pattern common in Bayesian models
with tractable conditionals (e.g., topic models, Boltzmann machines).
"""
def __init__(self, dim):
super().__init__()
self.dim = dim
# Precision matrix: diagonally dominant for well-conditioning
L = torch.randn(dim, dim) * 0.1
precision = L @ L.t() + torch.eye(dim) * 2.0
self.register_buffer("precision", precision)
self.register_buffer("mu", torch.zeros(dim))
def forward(self, x: torch.Tensor) -> torch.Tensor:
"""
Args:
x: (B, dim) – current state
Returns:
x_new: (B, dim) – state after one full Gibbs sweep
"""
x = x.clone()
for i in range(self.dim):
prec_ii = self.precision[i, i]
residual = x @ self.precision[i] - x[:, i] * prec_ii
cond_mean = self.mu[i] - residual / prec_ii
cond_var = 1.0 / prec_ii
x[:, i] = cond_mean + torch.sqrt(cond_var) * torch.randn(x.shape[0], device=x.device)
return x
dim = 32
batch_size = 4096
def get_inputs():
return [torch.randn(popcorn_pri.jitter_int(batch_size), dim)]
def get_init_inputs():
return [dim]
turn 033.5s turn · 33.5s LLM · 0.0s tools · tokens in=2,125 · out=726 · total=4,618
reasoning
The task is to replace the forward computation of this PyTorch model with a custom CUDA kernel. The model does one full sweep of Gibbs sampling for a multivariate Gaussian.
The task involves replacing the forward pass of a PyTorch model with a custom CUDA kernel. The original model performs a full sweep of Gibbs sampling for a multivariate Gaussian, looping over each dimension i, computing a conditional mean and variance based on the current state of x, and then sampling a new value for x[:, i] from a normal distribution conditioned on the rest.
Since the updates are sequential and each step depends on the previous updates to x, the CUDA kernel must replicate this order. With a batch size of 4096 and dimension 32, the computation is manageable, but I need to handle the randomness properly because the original uses torch .randn, which generates random numbers each time.
The challenge is achieving numerically equivalent outputs, as the kernel's RNG (like curand) won't produce the same numbers as PyTorch's RNG without careful seeding or passing the random values explicitly. This means the outputs won't match exactly due to randomness, but the task requires matching the 32 full sweeps of Gibbs sampling for a multivariate Gaussian.
To ensure numerical equivalence, one approach is to generate all the needed random numbers on the host using PyTorch's RNG, then pass them to the CUDA kernel as an additional input. However, the forward signature must remain the same (forward(self, x)), so I can't add arguments. Instead, I can generate the random numbers inside the forward method using the same RNG state as the reference, then use them in the conditional calculations and pass them to the CUDA kernel for the actual computation.
For the kernel, since CUDA kernels can't easily access host RNG, passing pre-generated tensors of random numbers 32 times batch_size numbers per sweep, for 32 sweeps, but since it's one full sweep, it's dim * batch_size numbers.
The reference generates dim times batch_size numbers, one batch per dimension.
assistant```python
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline
cuda_source = '''
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
__global__ void gibbs_kernel(
const float* __restrict__ x_in,
float* __restrict__ x_out,
const float* __restrict__ precision,
const float* __restrict__ rands,
const float* __restrict__ mu,
int batch_size,
int dim
) {
int b = blockIdx.x * blockDim.x + threadIdx.x;
if (b >= batch_size) return;
// Each thread processes one batch element, sequential over dim
extern __shared__ float s_prec[]; // not used for simplicity
for (int i = 0; i < dim; ++i) {
float prec_ii = precision[i * dim + i];
float residual = 0.0f;
for (int j = 0; j < dim; ++j) {
float xj = x_out[b * dim + j]; // use current values
residual += xj * precision[i * dim + j];
}
residual -= x_out[b * dim + i] * prec_ii;
float cond_mean = mu[i] - residual / prec_ii;
float cond_var = 1.0f / prec_ii;
float std = sqrtf(cond_var);
float rand_val = rands[b * dim + i];
x_out[b * dim + i] = cond_mean + std * rand_val;
}
}
torch::Tensor gibbs_cuda(
torch::Tensor x,
torch::Tensor precision,
torch::Tensor mu,
torch::Tensor rands
) {
auto batch_size = x.size(0);
auto dim = x.size(1);
auto x_out = x.clone();
const int threads = 256;
const int blocks = (batch_size + threads - 1) / threads;
gibbs_kernel<<<blocks, threads>>>(
x.data_ptr<float>(),
x_out.data_ptr<float>(),
precision.data_ptr<float>(),
rands.data_ptr<float>(),
mu.data_ptr<float>(),
batch_size,
dim
);
return x_out;
}
'''
cpp_source = '''
torch::Tensor gibbs_cuda(
torch::Tensor x,
torch::Tensor precision,
torch::Tensor mu,
torch::Tensor rands
);
'''
module = load_inline(
name='gibbs_sampler',
cpp_sources=cpp_source,
cuda_sources=cuda_source,
functions=['gibbs_cuda'],
verbose=False
)
class ModelNew(nn.Module):
def __init__(self, dim):
super().__init__()
self.dim = dim
L = torch.randn(dim, dim) * 0.1
precision = L @ L.t() + torch.eye(dim) * 2.0
self.register_buffer("precision", precision)
self.register_buffer("mu", torch.zeros(dim))
def forward(self, x: torch.Tensor) -> torch.Tensor:
batch_size = x.size(0)
flat_rand = torch.randn(self.dim * batch_size, device=x.device)
rands = flat_rand.view(self.dim, batch_size).T.contiguous()
return module.gibbs_cuda(x, self.precision, self.mu, rands)
```
turn 1320.9s turn · 5.4s LLM · 315.6s tools · tokens in=4,686 · out=698 · total=5,459
reasoning
The instruction says I must call at least one tool this turn. I need to use compile_kernel or run_correctness to iterate, or submit_kernel.
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.h>\n#include <cuda_runtime.h>\n\n__global__ void gibbs_kernel(\n const float* __restrict__ x_in,\n float* __restrict__ x_out,\n const float* __restrict__ precision,\n const float* __restrict__ rands,\n const float* __restrict__ mu,\n int batch_size,\n int dim\n) {\n int b = blockIdx.x * blockDim.x + threadIdx.x;\n if (b >= batch_size) return;\n\n for (int i = 0; i < dim; ++i) {\n float prec_ii = precision[i * dim + i];\n \n float residual = 0.0f;\n for (int j = 0; j < dim; ++j) {\n float xj = x_out[b * dim + j];\n residual += xj * precision[i * dim + j];\n }\n residual -= x_out[b * dim + i] * prec_ii;\n \n float cond_mean = mu[i] - residual / prec_ii;\n float cond_var = 1.0f / prec_ii;\n float std = sqrtf(cond_var);\n \n float rand_val = rands[b * dim + i];\n x_out[b * dim + i] = cond_mean + std * rand_val;\n }\n}\n\ntorch::Tensor gibbs_cuda(\n torch::Tensor x,\n torch::Tensor precision,\n torch::Tensor mu,\n torch::Tensor rands\n) {\n auto batch_size = x.size(0);\n auto dim = x.size(1);\n auto x_out = x.clone();\n \n const int threads = 256;\n const int blocks = (batch_size + threads - 1) / threads;\n \n gibbs_kernel<<<blocks, threads>>>(\n x.data_ptr<float>(),\n x_out.data_ptr<float>(),\n precision.data_ptr<float>(),\n rands.data_ptr<float>(),\n mu.data_ptr<float>(),\n batch_size,\n dim\n );\n \n return x_out;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor gibbs_cuda(\n torch::Tensor x,\n torch::Tensor precision,\n torch::Tensor mu,\n torch::Tensor rands\n);\n'''\n\nmodule = load_inline(\n name='gibbs_sampler',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['gibbs_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, dim):\n super().__init__()\n self.dim = dim\n L = torch.randn(dim, dim) * 0.1\n precision = L @ L.t() + torch.eye(dim) * 2.0\n self.register_buffer(\"precision\", precision)\n self.register_buffer(\"mu\", torch.zeros(dim))\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n batch_size = x.size(0)\n flat_rand = torch.randn(self.dim * batch_size, device=x.device)\n rands = flat_rand.view(self.dim, batch_size).T.contiguous()\n return module.gibbs_cuda(x, self.precision, self.mu, rands)\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 285.8s turn · 5.9s LLM · 79.9s tools · tokens in=5,476 · out=699 · total=6,190
reasoning
The kernel compiled successfully. Now I should test 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 = '''\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\n__global__ void gibbs_kernel(\n const float* __restrict__ x_in,\n float* __restrict__ x_out,\n const float* __restrict__ precision,\n const float* __restrict__ rands,\n const float* __restrict__ mu,\n int batch_size,\n int dim\n) {\n int b = blockIdx.x * blockDim.x + threadIdx.x;\n if (b >= batch_size) return;\n\n for (int i = 0; i < dim; ++i) {\n float prec_ii = precision[i * dim + i];\n \n float residual = 0.0f;\n for (int j = 0; j < dim; ++j) {\n float xj = x_out[b * dim + j];\n residual += xj * precision[i * dim + j];\n }\n residual -= x_out[b * dim + i] * prec_ii;\n \n float cond_mean = mu[i] - residual / prec_ii;\n float cond_var = 1.0f / prec_ii;\n float std = sqrtf(cond_var);\n \n float rand_val = rands[b * dim + i];\n x_out[b * dim + i] = cond_mean + std * rand_val;\n }\n}\n\ntorch::Tensor gibbs_cuda(\n torch::Tensor x,\n torch::Tensor precision,\n torch::Tensor mu,\n torch::Tensor rands\n) {\n auto batch_size = x.size(0);\n auto dim = x.size(1);\n auto x_out = x.clone();\n \n const int threads = 256;\n const int blocks = (batch_size + threads - 1) / threads;\n \n gibbs_kernel<<<blocks, threads>>>(\n x.data_ptr<float>(),\n x_out.data_ptr<float>(),\n precision.data_ptr<float>(),\n rands.data_ptr<float>(),\n mu.data_ptr<float>(),\n batch_size,\n dim\n );\n \n return x_out;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor gibbs_cuda(\n torch::Tensor x,\n torch::Tensor precision,\n torch::Tensor mu,\n torch::Tensor rands\n);\n'''\n\nmodule = load_inline(\n name='gibbs_sampler',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['gibbs_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, dim):\n super().__init__()\n self.dim = dim\n L = torch.randn(dim, dim) * 0.1\n precision = L @ L.t() + torch.eye(dim) * 2.0\n self.register_buffer(\"precision\", precision)\n self.register_buffer(\"mu\", torch.zeros(dim))\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n batch_size = x.size(0)\n flat_rand = torch.randn(self.dim * batch_size, device=x.device)\n rands = flat_rand.view(self.dim, batch_size).T.contiguous()\n return module.gibbs_cuda(x, self.precision, self.mu, rands)\n"
}
result ← run_correctness (fail)
failrun_correctness
run_correctness FAILED: (0 / 3) trials did not all match.
correctness_issue: Output mismatch
max_difference: ['4.077594', '4.665485', '4.354724']
avg_difference: ['0.721141', '0.721029', '0.722997']
turn 3185.6s turn · 36.1s LLM · 149.5s tools · tokens in=6,263 · out=705 · total=8,647
reasoning
There's an output mismatch. The rands ordering is wrong, leading to different values.
The issue is an output mismatch caused by the wrong ordering of random numbers. In the reference implementation, the random numbers are generated per dimension, with each dimension i getting a sequence of batch_size numbers in order. But in my code, I flattened the randoms and reshaped them incorrectly, leading to a different ordering.
To fix it, generate the flat_rand as torch.randn(dim * batch_size, ...), then reshape it to (batch_size, dim) instead of (dim, batch_size). This ensures that for each dimension i, the numbers are consecutive in blocks of batch_size. Transposing afterward isn't necessary if I directly use the ~~~truncated~~~
Truncated due to length, but the key is reshaping to (batch_size, dim) to match the reference's sequential generation.
I think the problem was in the reshape: using .view(self.dim, batch_size).T gave the wrong layout. The correct way is to use .view(batch_size, ~~~
Truncated due to length, but focus on the core: the bug is in the random ordering, and the fix is to adjust the reshape to match the reference order.
There seems to be no obvious bug in the residual calculation or other logic, so the mismatch is due to the rands ordering.
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.h>\n#include <cuda_runtime.h>\n\n__global__ void gibbs_kernel(\n const float* __restrict__ x_in,\n float* __restrict__ x_out,\n const float* __restrict__ precision,\n const float* __restrict__ rands,\n const float* __restrict__ mu,\n int batch_size,\n int dim\n) {\n int b = blockIdx.x * blockDim.x + threadIdx.x;\n if (b >= batch_size) return;\n\n for (int i = 0; i < dim; ++i) {\n float prec_ii = precision[i * dim + i];\n \n float residual = 0.0f;\n for (int j = 0; j < dim; ++j) {\n float xj = x_out[b * dim + j];\n residual += xj * precision[i * dim + j];\n }\n residual -= x_out[b * dim + i] * prec_ii;\n \n float cond_mean = mu[i] - residual / prec_ii;\n float cond_var = 1.0f / prec_ii;\n float std = sqrtf(cond_var);\n \n float rand_val = rands[b * dim + i];\n x_out[b * dim + i] = cond_mean + std * rand_val;\n }\n}\n\ntorch::Tensor gibbs_cuda(\n torch::Tensor x,\n torch::Tensor precision,\n torch::Tensor mu,\n torch::Tensor rands\n) {\n auto batch_size = x.size(0);\n auto dim = x.size(1);\n auto x_out = x.clone();\n \n const int threads = 256;\n const int blocks = (batch_size + threads - 1) / threads;\n \n gibbs_kernel<<<blocks, threads>>>(\n x.data_ptr<float>(),\n x_out.data_ptr<float>(),\n precision.data_ptr<float>(),\n rands.data_ptr<float>(),\n mu.data_ptr<float>(),\n batch_size,\n dim\n );\n \n return x_out;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor gibbs_cuda(\n torch::Tensor x,\n torch::Tensor precision,\n torch::Tensor mu,\n torch::Tensor rands\n);\n'''\n\nmodule = load_inline(\n name='gibbs_sampler',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['gibbs_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, dim):\n super().__init__()\n self.dim = dim\n L = torch.randn(dim, dim) * 0.1\n precision = L @ L.t() + torch.eye(dim) * 2.0\n self.register_buffer(\"precision\", precision)\n self.register_buffer(\"mu\", torch.zeros(dim))\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n batch_size = x.size(0)\n rands = torch.zeros(batch_size, self.dim, device=x.device)\n for i in range(self.dim):\n rands[:, i] = torch.randn(batch_size, device=x.device)\n return module.gibbs_cuda(x, self.precision, self.mu, rands)\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 4148.0s turn · 4.6s LLM · 143.4s tools · tokens in=8,664 · out=706 · total=9,376
reasoning
Compiled. Now 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 = '''\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\n__global__ void gibbs_kernel(\n const float* __restrict__ x_in,\n float* __restrict__ x_out,\n const float* __restrict__ precision,\n const float* __restrict__ rands,\n const float* __restrict__ mu,\n int batch_size,\n int dim\n) {\n int b = blockIdx.x * blockDim.x + threadIdx.x;\n if (b >= batch_size) return;\n\n for (int i = 0; i < dim; ++i) {\n float prec_ii = precision[i * dim + i];\n \n float residual = 0.0f;\n for (int j = 0; j < dim; ++j) {\n float xj = x_out[b * dim + j];\n residual += xj * precision[i * dim + j];\n }\n residual -= x_out[b * dim + i] * prec_ii;\n \n float cond_mean = mu[i] - residual / prec_ii;\n float cond_var = 1.0f / prec_ii;\n float std = sqrtf(cond_var);\n \n float rand_val = rands[b * dim + i];\n x_out[b * dim + i] = cond_mean + std * rand_val;\n }\n}\n\ntorch::Tensor gibbs_cuda(\n torch::Tensor x,\n torch::Tensor precision,\n torch::Tensor mu,\n torch::Tensor rands\n) {\n auto batch_size = x.size(0);\n auto dim = x.size(1);\n auto x_out = x.clone();\n \n const int threads = 256;\n const int blocks = (batch_size + threads - 1) / threads;\n \n gibbs_kernel<<<blocks, threads>>>(\n x.data_ptr<float>(),\n x_out.data_ptr<float>(),\n precision.data_ptr<float>(),\n rands.data_ptr<float>(),\n mu.data_ptr<float>(),\n batch_size,\n dim\n );\n \n return x_out;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor gibbs_cuda(\n torch::Tensor x,\n torch::Tensor precision,\n torch::Tensor mu,\n torch::Tensor rands\n);\n'''\n\nmodule = load_inline(\n name='gibbs_sampler',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['gibbs_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, dim):\n super().__init__()\n self.dim = dim\n L = torch.randn(dim, dim) * 0.1\n precision = L @ L.t() + torch.eye(dim) * 2.0\n self.register_buffer(\"precision\", precision)\n self.register_buffer(\"mu\", torch.zeros(dim))\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n batch_size = x.size(0)\n rands = torch.zeros(batch_size, self.dim, device=x.device)\n for i in range(self.dim):\n rands[:, i] = torch.randn(batch_size, device=x.device)\n return module.gibbs_cuda(x, self.precision, self.mu, rands)\n"
}
result ← run_correctness (ok)
okrun_correctness
run_correctness PASSED: (3 / 3) trials all matched the reference.
Numerical precision: max_abs_err=1.79e-06 mean_abs_err=5.69e-08 max_rel_err=2.41e-02
turn 5208.9s turn · 6.0s LLM · 202.9s tools · tokens in=9,436 · out=705 · total=10,276 · FINAL
reasoning
Correctness passed! Great, the loop generation of rands matches exactly.
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 = '''\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\n__global__ void gibbs_kernel(\n const float* __restrict__ x_in,\n float* __restrict__ x_out,\n const float* __restrict__ precision,\n const float* __restrict__ rands,\n const float* __restrict__ mu,\n int batch_size,\n int dim\n) {\n int b = blockIdx.x * blockDim.x + threadIdx.x;\n if (b >= batch_size) return;\n\n for (int i = 0; i < dim; ++i) {\n float prec_ii = precision[i * dim + i];\n \n float residual = 0.0f;\n for (int j = 0; j < dim; ++j) {\n float xj = x_out[b * dim + j];\n residual += xj * precision[i * dim + j];\n }\n residual -= x_out[b * dim + i] * prec_ii;\n \n float cond_mean = mu[i] - residual / prec_ii;\n float cond_var = 1.0f / prec_ii;\n float std = sqrtf(cond_var);\n \n float rand_val = rands[b * dim + i];\n x_out[b * dim + i] = cond_mean + std * rand_val;\n }\n}\n\ntorch::Tensor gibbs_cuda(\n torch::Tensor x,\n torch::Tensor precision,\n torch::Tensor mu,\n torch::Tensor rands\n) {\n auto batch_size = x.size(0);\n auto dim = x.size(1);\n auto x_out = x.clone();\n \n const int threads = 256;\n const int blocks = (batch_size + threads - 1) / threads;\n \n gibbs_kernel<<<blocks, threads>>>(\n x.data_ptr<float>(),\n x_out.data_ptr<float>(),\n precision.data_ptr<float>(),\n rands.data_ptr<float>(),\n mu.data_ptr<float>(),\n batch_size,\n dim\n );\n \n return x_out;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor gibbs_cuda(\n torch::Tensor x,\n torch::Tensor precision,\n torch::Tensor mu,\n torch::Tensor rands\n);\n'''\n\nmodule = load_inline(\n name='gibbs_sampler',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['gibbs_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, dim):\n super().__init__()\n self.dim = dim\n L = torch.randn(dim, dim) * 0.1\n precision = L @ L.t() + torch.eye(dim) * 2.0\n self.register_buffer(\"precision\", precision)\n self.register_buffer(\"mu\", torch.zeros(dim))\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n batch_size = x.size(0)\n rands = torch.zeros(batch_size, self.dim, device=x.device)\n for i in range(self.dim):\n rands[:, i] = torch.randn(batch_size, device=x.device)\n return module.gibbs_cuda(x, self.precision, self.mu, rands)\n"
}
result ← submit_kernel (ok)
oksubmit_kernel
submit_kernel PASSED: (3 / 3) correctness trials all passed.
Kernel runtime: 0.62 μs
Runtime stats: mean=0.62μs median=?μs std=0.01μs
Numerical precision: max_abs_err=1.79e-06 mean_abs_err=5.69e-08 max_rel_err=2.41e-02
Memory: 34.51 MB peak (ref: 34.09 MB, ratio: 1.0124x)
Kernel launches: 199 (ref: 931, fusion ratio: 4.6784)
SOL score: 0.1888
final submitted kernel (100 lines, 2,605 chars) — full source, untruncated
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline
cuda_source = '''
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
__global__ void gibbs_kernel(
const float* __restrict__ x_in,
float* __restrict__ x_out,
const float* __restrict__ precision,
const float* __restrict__ rands,
const float* __restrict__ mu,
int batch_size,
int dim
) {
int b = blockIdx.x * blockDim.x + threadIdx.x;
if (b >= batch_size) return;
for (int i = 0; i < dim; ++i) {
float prec_ii = precision[i * dim + i];
float residual = 0.0f;
for (int j = 0; j < dim; ++j) {
float xj = x_out[b * dim + j];
residual += xj * precision[i * dim + j];
}
residual -= x_out[b * dim + i] * prec_ii;
float cond_mean = mu[i] - residual / prec_ii;
float cond_var = 1.0f / prec_ii;
float std = sqrtf(cond_var);
float rand_val = rands[b * dim + i];
x_out[b * dim + i] = cond_mean + std * rand_val;
}
}
torch::Tensor gibbs_cuda(
torch::Tensor x,
torch::Tensor precision,
torch::Tensor mu,
torch::Tensor rands
) {
auto batch_size = x.size(0);
auto dim = x.size(1);
auto x_out = x.clone();
const int threads = 256;
const int blocks = (batch_size + threads - 1) / threads;
gibbs_kernel<<<blocks, threads>>>(
x.data_ptr<float>(),
x_out.data_ptr<float>(),
precision.data_ptr<float>(),
rands.data_ptr<float>(),
mu.data_ptr<float>(),
batch_size,
dim
);
return x_out;
}
'''
cpp_source = '''
torch::Tensor gibbs_cuda(
torch::Tensor x,
torch::Tensor precision,
torch::Tensor mu,
torch::Tensor rands
);
'''
module = load_inline(
name='gibbs_sampler',
cpp_sources=cpp_source,
cuda_sources=cuda_source,
functions=['gibbs_cuda'],
verbose=False
)
class ModelNew(nn.Module):
def __init__(self, dim):
super().__init__()
self.dim = dim
L = torch.randn(dim, dim) * 0.1
precision = L @ L.t() + torch.eye(dim) * 2.0
self.register_buffer("precision", precision)
self.register_buffer("mu", torch.zeros(dim))
def forward(self, x: torch.Tensor) -> torch.Tensor:
batch_size = x.size(0)
rands = torch.zeros(batch_size, self.dim, device=x.device)
for i in range(self.dim):
rands[:, i] = torch.randn(batch_size, device=x.device)
return module.gibbs_cuda(x, self.precision, self.mu, rands)