correct
26_ImportanceSampling.py
started
2026-05-11T17:36:19Z
timingkernelbench.timing
kernel mean / median / std
0.18 / — / 0.02 μs
ref mean / median / std
0.43 / — / 0.06 μs
numerical precisioneval_kernel_against_ref
gpu memoryextended_metrics.measure_memory
kernel launches / fusiontorch.profiler
fusion ratio (ref / kernel)
2.38x
SOL (speed-of-light)nsight_hardware_counters
roofline / occupancynsight_hardware_counters
dram bandwidth
210.96 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):
"""
Self-normalized importance sampling estimate.
Given samples from a proposal distribution q, computes importance
weights w_i = p(x_i)/q(x_i), normalizes them, and returns both
the weighted estimate of E_p[f(x)] and the effective sample size.
Target p is a mixture of Gaussians; proposal q is a single Gaussian.
"""
def __init__(self, dim, num_components=3):
super().__init__()
self.dim = dim
self.num_components = num_components
# Target mixture components
self.register_buffer("mix_weights", torch.ones(num_components) / num_components)
mus = torch.randn(num_components, dim) * 2.0
self.register_buffer("mix_mus", mus)
self.register_buffer("mix_stds", torch.ones(num_components, dim))
# Proposal
self.register_buffer("q_mu", torch.zeros(dim))
self.register_buffer("q_std", torch.ones(dim) * 3.0)
def _log_p(self, x):
"""Log density of target mixture."""
diff = x.unsqueeze(1) - self.mix_mus.unsqueeze(0) # (B, K, D)
log_comp = -0.5 * (diff / self.mix_stds.unsqueeze(0)) ** 2 - self.mix_stds.log().unsqueeze(0)
log_comp = log_comp.sum(-1) # (B, K)
return torch.logsumexp(log_comp + self.mix_weights.log().unsqueeze(0), dim=-1)
def _log_q(self, x):
"""Log density of proposal."""
return (-0.5 * ((x - self.q_mu) / self.q_std) ** 2 - self.q_std.log()).sum(-1)
def forward(self, samples: torch.Tensor) -> tuple:
"""
Args:
samples: (N, dim) – samples drawn from q
Returns:
(weighted_mean, ess):
weighted_mean: (dim,) – importance-weighted mean estimate of E_p[x]
ess: scalar – effective sample size
"""
log_w = self._log_p(samples) - self._log_q(samples)
log_w = log_w - torch.logsumexp(log_w, dim=0)
w = torch.exp(log_w)
weighted_mean = (w.unsqueeze(-1) * samples).sum(dim=0)
ess = 1.0 / (w ** 2).sum()
return weighted_mean, ess
dim = 16
num_samples = 4096
def get_inputs():
q_std = 3.0
n = popcorn_pri.jitter_int(num_samples)
return [torch.randn(n, dim) * q_std]
def get_init_inputs():
return [dim]
turn 07.0s turn · 7.0s LLM · 0.0s tools · tokens in=2,364 · out=6 · total=2,850
reasoning
The task is to replace the forward computation of this PyTorch model with a custom CUDA kernel. The model computes self-normalized importance sampling estimates.
I need to replace the forward computation of a PyTorch model with a custom CUDA kernel. The model performs self-normalized importance sampling estimates. It uses buffers for mixture weights, mus, stds for the target distribution, and q_mu, q_std for the proposal.
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 187.9s turn · 22.0s LLM · 65.9s tools · tokens in=3,083 · out=560 · total=3,937
reasoning
H100, high bandwidth.
Now, I need to implement the 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#include <math.h>\n\n__global__ void importance_sampling_kernel(\n const float* samples, // (N, D)\n const float* mix_mus, // (K, D)\n const float* mix_stds, // (K, D)\n const float* mix_weights_log, // (K,)\n const float* q_mu, // (D,)\n const float* q_std, // (D,)\n float* log_ws, // (N,)\n int N, int D, int K\n) {\n int i = blockIdx.x * blockDim.x + threadIdx.x;\n if (i >= N) return;\n\n float log_p = -INFINITY;\n const float* x = samples + i * D;\n for (int k = 0; k < K; k++) {\n float log_comp = 0.0f;\n for (int d = 0; d < D; d++) {\n float diff = x[d] - mix_mus[k * D + d];\n float std = mix_stds[k * D + d];\n log_comp += -0.5f * (diff / std) * (diff / std) - logf(std);\n }\n log_comp += mix_weights_log[k];\n log_p = fmaxf(log_p, log_comp) + logf(1.0f + expf(fminf(log_p, log_comp) - fmaxf(log_p, log_comp))); wait, better use logsumexp properly later.\n }\n // Actually, implement logsumexp correctly in loop.\n}\n\n__global__ void compute_log_w(\n const float* samples,\n float* log_w,\n const float* params, // packed\n int N, int D, int K\n) {\n // To be implemented properly\n}\n'''\n\ncpp_source = '''\ntorch::Tensor compute_importance_weights(torch::Tensor samples, torch::Tensor mix_mus, torch::Tensor mix_stds, torch::Tensor mix_weights, torch::Tensor q_mu, torch::Tensor q_std);\nstd::tuple<torch::Tensor, torch::Tensor> importance_sampling_forward(torch::Tensor samples, torch::Tensor mix_mus, torch::Tensor mix_stds, torch::Tensor mix_weights, torch::Tensor q_mu, torch::Tensor q_std);\n'''\n\n# Note: I'll fix the kernel in next iteration. For now, this is placeholder to test structure."
}
result ← compile_kernel (fail)
failcompile_kernel
compile_kernel FAILED: ModelNew class not found or syntax error prevented execution.
turn 2102.4s turn · 14.1s LLM · 88.3s tools · tokens in=3,959 · out=1,005 · total=4,996
reasoning
I need to provide the full module with ModelNew class.
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 <cmath>\n\n__device__ float logsumexp_device(float a, float b) {\n if (a == -INFINITY) return b;\n if (b == -INFINITY) return a;\n float m = fmaxf(a, b);\n return m + logf(expf(a - m) + expf(b - m));\n}\n\n__global__ void compute_log_w_kernel(\n const float* __restrict__ samples,\n const float* __restrict__ mix_mus,\n const float* __restrict__ mix_stds,\n const float* __restrict__ mix_weights_log,\n const float* __restrict__ q_mu,\n const float* __restrict__ q_std,\n float* __restrict__ log_w,\n int N, int D, int K\n) {\n int i = blockIdx.x * blockDim.x + threadIdx.x;\n if (i >= N) return;\n const float* x = samples + i * D;\n float log_p = -INFINITY;\n for (int k = 0; k < K; ++k) {\n float log_comp = mix_weights_log[k];\n for (int d = 0; d < D; ++d) {\n float diff = x[d] - mix_mus[k*D + d];\n float stdv = mix_stds[k*D + d];\n log_comp += -0.5f * diff*diff / (stdv*stdv) - logf(stdv);\n }\n log_p = logsumexp_device(log_p, log_comp);\n }\n float log_q = 0.0f;\n for (int d = 0; d < D; ++d) {\n float diff = x[d] - q_mu[d];\n float stdv = q_std[d];\n log_q += -0.5f * diff*diff / (stdv*stdv) - logf(stdv);\n }\n log_w[i] = log_p - log_q;\n}\n\ntorch::Tensor compute_log_w_cuda(torch::Tensor samples, torch::Tensor mix_mus, torch::Tensor mix_stds,\n torch::Tensor mix_weights_log, torch::Tensor q_mu, torch::Tensor q_std) {\n int N = samples.size(0);\n int D = samples.size(1);\n int K = mix_mus.size(0);\n auto log_w = torch::empty({N}, samples.options());\n dim3 block(256);\n dim3 grid((N + 255) / 256);\n compute_log_w_kernel<<<grid, block>>>(\n samples.data_ptr<float>(),\n mix_mus.data_ptr<float>(),\n mix_stds.data_ptr<float>(),\n mix_weights_log.data_ptr<float>(),\n q_mu.data_ptr<float>(),\n q_std.data_ptr<float>(),\n log_w.data_ptr<float>(),\n N, D, K\n );\n return log_w;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor compute_log_w_cuda(torch::Tensor samples, torch::Tensor mix_mus, torch::Tensor mix_stds,\n torch::Tensor mix_weights_log, torch::Tensor q_mu, torch::Tensor q_std);\n'''\n\nmodule = load_inline(\n name='importance_sampling',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['compute_log_w_cuda'],\n verbose=False\n)\n'''\n\nclass ModelNew(nn.Module):\n def __init__(self, dim, num_components=3):\n super().__init__()\n self.dim = dim\n self.num_components = num_components\n self.register_buffer(\"mix_weights\", torch.ones(num_components) / num_components)\n mus = torch.randn(num_components, dim) * 2.0\n self.register_buffer(\"mix_mus\", mus)\n self.register_buffer(\"mix_stds\", torch.ones(num_components, dim))\n self.register_buffer(\"q_mu\", torch.zeros(dim))\n self.register_buffer(\"q_std\", torch.ones(dim) * 3.0)\n\n def forward(self, samples: torch.Tensor) -> tuple:\n mix_weights_log = self.mix_weights.log()\n log_w = module.compute_log_w_cuda(\n samples, self.mix_mus, self.mix_stds, mix_weights_log, self.q_mu, self.q_std\n )\n log_w = log_w - torch.logsumexp(log_w, dim=0)\n w = torch.exp(log_w)\n weighted_mean = (w.unsqueeze(-1) * samples).sum(dim=0)\n ess = 1.0 / (w ** 2).sum()\n return weighted_mean, ess\n"
}
result ← compile_kernel (fail)
failcompile_kernel
compile_kernel FAILED: ModelNew class not found or syntax error prevented execution.
Syntax Error in custom generated code or Compilation Error unterminated triple-quoted string literal (detected at line 108) (<string>, line 85)
turn 3161.1s turn · 10.2s LLM · 150.9s tools · tokens in=5,049 · out=1,003 · total=6,083
reasoning
Syntax error because of the r''' and quotes.
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 <cmath>\n\n__device__ float logsumexp_device(float a, float b) {\n if (a == -INFINITY) return b;\n if (b == -INFINITY) return a;\n float m = fmaxf(a, b);\n return m + logf(expf(a - m) + expf(b - m));\n}\n\n__global__ void compute_log_w_kernel(\n const float* __restrict__ samples,\n const float* __restrict__ mix_mus,\n const float* __restrict__ mix_stds,\n const float* __restrict__ mix_weights_log,\n const float* __restrict__ q_mu,\n const float* __restrict__ q_std,\n float* __restrict__ log_w,\n int N, int D, int K\n) {\n int i = blockIdx.x * blockDim.x + threadIdx.x;\n if (i >= N) return;\n const float* x = samples + i * D;\n float log_p = -INFINITY;\n for (int k = 0; k < K; ++k) {\n float log_comp = mix_weights_log[k];\n for (int d = 0; d < D; ++d) {\n float diff = x[d] - mix_mus[k*D + d];\n float stdv = mix_stds[k*D + d];\n log_comp += -0.5f * diff*diff / (stdv*stdv) - logf(stdv);\n }\n log_p = logsumexp_device(log_p, log_comp);\n }\n float log_q = 0.0f;\n for (int d = 0; d < D; ++d) {\n float diff = x[d] - q_mu[d];\n float stdv = q_std[d];\n log_q += -0.5f * diff*diff / (stdv*stdv) - logf(stdv);\n }\n log_w[i] = log_p - log_q;\n}\n\ntorch::Tensor compute_log_w_cuda(torch::Tensor samples, torch::Tensor mix_mus, torch::Tensor mix_stds,\n torch::Tensor mix_weights_log, torch::Tensor q_mu, torch::Tensor q_std) {\n int N = samples.size(0);\n int D = samples.size(1);\n int K = mix_mus.size(0);\n auto log_w = torch::empty({N}, samples.options());\n dim3 block(256);\n dim3 grid((N + 255) / 256);\n compute_log_w_kernel<<<grid, block>>>(\n samples.data_ptr<float>(),\n mix_mus.data_ptr<float>(),\n mix_stds.data_ptr<float>(),\n mix_weights_log.data_ptr<float>(),\n q_mu.data_ptr<float>(),\n q_std.data_ptr<float>(),\n log_w.data_ptr<float>(),\n N, D, K\n );\n return log_w;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor compute_log_w_cuda(torch::Tensor samples, torch::Tensor mix_mus, torch::Tensor mix_stds,\n torch::Tensor mix_weights_log, torch::Tensor q_mu, torch::Tensor q_std);\n'''\n\nmodule = load_inline(\n name='importance_sampling',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['compute_log_w_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, dim, num_components=3):\n super().__init__()\n self.dim = dim\n self.num_components = num_components\n self.register_buffer(\"mix_weights\", torch.ones(num_components) / num_components)\n mus = torch.randn(num_components, dim) * 2.0\n self.register_buffer(\"mix_mus\", mus)\n self.register_buffer(\"mix_stds\", torch.ones(num_components, dim))\n self.register_buffer(\"q_mu\", torch.zeros(dim))\n self.register_buffer(\"q_std\", torch.ones(dim) * 3.0)\n\n def forward(self, samples: torch.Tensor) -> tuple:\n mix_weights_log = self.mix_weights.log()\n log_w = module.compute_log_w_cuda(\n samples, self.mix_mus, self.mix_stds, mix_weights_log, self.q_mu, self.q_std\n )\n log_w = log_w - torch.logsumexp(log_w, dim=0)\n w = torch.exp(log_w)\n weighted_mean = (w.unsqueeze(-1) * samples).sum(dim=0)\n ess = 1.0 / (w ** 2).sum()\n return weighted_mean, ess\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 4152.4s turn · 13.0s LLM · 139.4s tools · tokens in=6,100 · out=1,004 · total=7,111
reasoning
Compiled successfully. Now check 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_runtime.h>\n#include <cmath>\n\n__device__ float logsumexp_device(float a, float b) {\n if (a == -INFINITY) return b;\n if (b == -INFINITY) return a;\n float m = fmaxf(a, b);\n return m + logf(expf(a - m) + expf(b - m));\n}\n\n__global__ void compute_log_w_kernel(\n const float* __restrict__ samples,\n const float* __restrict__ mix_mus,\n const float* __restrict__ mix_stds,\n const float* __restrict__ mix_weights_log,\n const float* __restrict__ q_mu,\n const float* __restrict__ q_std,\n float* __restrict__ log_w,\n int N, int D, int K\n) {\n int i = blockIdx.x * blockDim.x + threadIdx.x;\n if (i >= N) return;\n const float* x = samples + i * D;\n float log_p = -INFINITY;\n for (int k = 0; k < K; ++k) {\n float log_comp = mix_weights_log[k];\n for (int d = 0; d < D; ++d) {\n float diff = x[d] - mix_mus[k*D + d];\n float stdv = mix_stds[k*D + d];\n log_comp += -0.5f * diff*diff / (stdv*stdv) - logf(stdv);\n }\n log_p = logsumexp_device(log_p, log_comp);\n }\n float log_q = 0.0f;\n for (int d = 0; d < D; ++d) {\n float diff = x[d] - q_mu[d];\n float stdv = q_std[d];\n log_q += -0.5f * diff*diff / (stdv*stdv) - logf(stdv);\n }\n log_w[i] = log_p - log_q;\n}\n\ntorch::Tensor compute_log_w_cuda(torch::Tensor samples, torch::Tensor mix_mus, torch::Tensor mix_stds,\n torch::Tensor mix_weights_log, torch::Tensor q_mu, torch::Tensor q_std) {\n int N = samples.size(0);\n int D = samples.size(1);\n int K = mix_mus.size(0);\n auto log_w = torch::empty({N}, samples.options());\n dim3 block(256);\n dim3 grid((N + 255) / 256);\n compute_log_w_kernel<<<grid, block>>>(\n samples.data_ptr<float>(),\n mix_mus.data_ptr<float>(),\n mix_stds.data_ptr<float>(),\n mix_weights_log.data_ptr<float>(),\n q_mu.data_ptr<float>(),\n q_std.data_ptr<float>(),\n log_w.data_ptr<float>(),\n N, D, K\n );\n return log_w;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor compute_log_w_cuda(torch::Tensor samples, torch::Tensor mix_mus, torch::Tensor mix_stds,\n torch::Tensor mix_weights_log, torch::Tensor q_mu, torch::Tensor q_std);\n'''\n\nmodule = load_inline(\n name='importance_sampling',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['compute_log_w_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, dim, num_components=3):\n super().__init__()\n self.dim = dim\n self.num_components = num_components\n self.register_buffer(\"mix_weights\", torch.ones(num_components) / num_components)\n mus = torch.randn(num_components, dim) * 2.0\n self.register_buffer(\"mix_mus\", mus)\n self.register_buffer(\"mix_stds\", torch.ones(num_components, dim))\n self.register_buffer(\"q_mu\", torch.zeros(dim))\n self.register_buffer(\"q_std\", torch.ones(dim) * 3.0)\n\n def forward(self, samples: torch.Tensor) -> tuple:\n mix_weights_log = self.mix_weights.log()\n log_w = module.compute_log_w_cuda(\n samples, self.mix_mus, self.mix_stds, mix_weights_log, self.q_mu, self.q_std\n )\n log_w = log_w - torch.logsumexp(log_w, dim=0)\n w = torch.exp(log_w)\n weighted_mean = (w.unsqueeze(-1) * samples).sum(dim=0)\n ess = 1.0 / (w ** 2).sum()\n return weighted_mean, ess\n"
}
result ← run_correctness (ok)
okrun_correctness
run_correctness PASSED: (3 / 3) trials all matched the reference.
Numerical precision: max_abs_err=7.63e-06 mean_abs_err=1.08e-06 max_rel_err=1.06e-05
turn 595.8s turn · 13.9s LLM · 82.0s tools · tokens in=7,171 · out=1,003 · total=8,204 · FINAL
reasoning
Correct! Now, to check if faster, I should submit or profile.
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_runtime.h>\n#include <cmath>\n\n__device__ float logsumexp_device(float a, float b) {\n if (a == -INFINITY) return b;\n if (b == -INFINITY) return a;\n float m = fmaxf(a, b);\n return m + logf(expf(a - m) + expf(b - m));\n}\n\n__global__ void compute_log_w_kernel(\n const float* __restrict__ samples,\n const float* __restrict__ mix_mus,\n const float* __restrict__ mix_stds,\n const float* __restrict__ mix_weights_log,\n const float* __restrict__ q_mu,\n const float* __restrict__ q_std,\n float* __restrict__ log_w,\n int N, int D, int K\n) {\n int i = blockIdx.x * blockDim.x + threadIdx.x;\n if (i >= N) return;\n const float* x = samples + i * D;\n float log_p = -INFINITY;\n for (int k = 0; k < K; ++k) {\n float log_comp = mix_weights_log[k];\n for (int d = 0; d < D; ++d) {\n float diff = x[d] - mix_mus[k*D + d];\n float stdv = mix_stds[k*D + d];\n log_comp += -0.5f * diff*diff / (stdv*stdv) - logf(stdv);\n }\n log_p = logsumexp_device(log_p, log_comp);\n }\n float log_q = 0.0f;\n for (int d = 0; d < D; ++d) {\n float diff = x[d] - q_mu[d];\n float stdv = q_std[d];\n log_q += -0.5f * diff*diff / (stdv*stdv) - logf(stdv);\n }\n log_w[i] = log_p - log_q;\n}\n\ntorch::Tensor compute_log_w_cuda(torch::Tensor samples, torch::Tensor mix_mus, torch::Tensor mix_stds,\n torch::Tensor mix_weights_log, torch::Tensor q_mu, torch::Tensor q_std) {\n int N = samples.size(0);\n int D = samples.size(1);\n int K = mix_mus.size(0);\n auto log_w = torch::empty({N}, samples.options());\n dim3 block(256);\n dim3 grid((N + 255) / 256);\n compute_log_w_kernel<<<grid, block>>>(\n samples.data_ptr<float>(),\n mix_mus.data_ptr<float>(),\n mix_stds.data_ptr<float>(),\n mix_weights_log.data_ptr<float>(),\n q_mu.data_ptr<float>(),\n q_std.data_ptr<float>(),\n log_w.data_ptr<float>(),\n N, D, K\n );\n return log_w;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor compute_log_w_cuda(torch::Tensor samples, torch::Tensor mix_mus, torch::Tensor mix_stds,\n torch::Tensor mix_weights_log, torch::Tensor q_mu, torch::Tensor q_std);\n'''\n\nmodule = load_inline(\n name='importance_sampling',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['compute_log_w_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, dim, num_components=3):\n super().__init__()\n self.dim = dim\n self.num_components = num_components\n self.register_buffer(\"mix_weights\", torch.ones(num_components) / num_components)\n mus = torch.randn(num_components, dim) * 2.0\n self.register_buffer(\"mix_mus\", mus)\n self.register_buffer(\"mix_stds\", torch.ones(num_components, dim))\n self.register_buffer(\"q_mu\", torch.zeros(dim))\n self.register_buffer(\"q_std\", torch.ones(dim) * 3.0)\n\n def forward(self, samples: torch.Tensor) -> tuple:\n mix_weights_log = self.mix_weights.log()\n log_w = module.compute_log_w_cuda(\n samples, self.mix_mus, self.mix_stds, mix_weights_log, self.q_mu, self.q_std\n )\n log_w = log_w - torch.logsumexp(log_w, dim=0)\n w = torch.exp(log_w)\n weighted_mean = (w.unsqueeze(-1) * samples).sum(dim=0)\n ess = 1.0 / (w ** 2).sum()\n return weighted_mean, ess\n"
}
result ← submit_kernel (ok)
oksubmit_kernel
submit_kernel PASSED: (3 / 3) correctness trials all passed.
Kernel runtime: 0.18 μs
Runtime stats: mean=0.18μs median=?μs std=0.02μs
Numerical precision: max_abs_err=7.63e-06 mean_abs_err=1.08e-06 max_rel_err=1.06e-05
Memory: 33.54 MB peak (ref: 35.51 MB, ratio: 0.9446x)
Kernel launches: 37 (ref: 88, fusion ratio: 2.3784)
SOL score: 0.486
final submitted kernel (106 lines, 3,583 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_runtime.h>
#include <cmath>
__device__ float logsumexp_device(float a, float b) {
if (a == -INFINITY) return b;
if (b == -INFINITY) return a;
float m = fmaxf(a, b);
return m + logf(expf(a - m) + expf(b - m));
}
__global__ void compute_log_w_kernel(
const float* __restrict__ samples,
const float* __restrict__ mix_mus,
const float* __restrict__ mix_stds,
const float* __restrict__ mix_weights_log,
const float* __restrict__ q_mu,
const float* __restrict__ q_std,
float* __restrict__ log_w,
int N, int D, int K
) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i >= N) return;
const float* x = samples + i * D;
float log_p = -INFINITY;
for (int k = 0; k < K; ++k) {
float log_comp = mix_weights_log[k];
for (int d = 0; d < D; ++d) {
float diff = x[d] - mix_mus[k*D + d];
float stdv = mix_stds[k*D + d];
log_comp += -0.5f * diff*diff / (stdv*stdv) - logf(stdv);
}
log_p = logsumexp_device(log_p, log_comp);
}
float log_q = 0.0f;
for (int d = 0; d < D; ++d) {
float diff = x[d] - q_mu[d];
float stdv = q_std[d];
log_q += -0.5f * diff*diff / (stdv*stdv) - logf(stdv);
}
log_w[i] = log_p - log_q;
}
torch::Tensor compute_log_w_cuda(torch::Tensor samples, torch::Tensor mix_mus, torch::Tensor mix_stds,
torch::Tensor mix_weights_log, torch::Tensor q_mu, torch::Tensor q_std) {
int N = samples.size(0);
int D = samples.size(1);
int K = mix_mus.size(0);
auto log_w = torch::empty({N}, samples.options());
dim3 block(256);
dim3 grid((N + 255) / 256);
compute_log_w_kernel<<<grid, block>>>(
samples.data_ptr<float>(),
mix_mus.data_ptr<float>(),
mix_stds.data_ptr<float>(),
mix_weights_log.data_ptr<float>(),
q_mu.data_ptr<float>(),
q_std.data_ptr<float>(),
log_w.data_ptr<float>(),
N, D, K
);
return log_w;
}
'''
cpp_source = '''
torch::Tensor compute_log_w_cuda(torch::Tensor samples, torch::Tensor mix_mus, torch::Tensor mix_stds,
torch::Tensor mix_weights_log, torch::Tensor q_mu, torch::Tensor q_std);
'''
module = load_inline(
name='importance_sampling',
cpp_sources=cpp_source,
cuda_sources=cuda_source,
functions=['compute_log_w_cuda'],
verbose=False
)
class ModelNew(nn.Module):
def __init__(self, dim, num_components=3):
super().__init__()
self.dim = dim
self.num_components = num_components
self.register_buffer("mix_weights", torch.ones(num_components) / num_components)
mus = torch.randn(num_components, dim) * 2.0
self.register_buffer("mix_mus", mus)
self.register_buffer("mix_stds", torch.ones(num_components, dim))
self.register_buffer("q_mu", torch.zeros(dim))
self.register_buffer("q_std", torch.ones(dim) * 3.0)
def forward(self, samples: torch.Tensor) -> tuple:
mix_weights_log = self.mix_weights.log()
log_w = module.compute_log_w_cuda(
samples, self.mix_mus, self.mix_stds, mix_weights_log, self.q_mu, self.q_std
)
log_w = log_w - torch.logsumexp(log_w, dim=0)
w = torch.exp(log_w)
weighted_mean = (w.unsqueeze(-1) * samples).sum(dim=0)
ess = 1.0 / (w ** 2).sum()
return weighted_mean, ess