grok-4.3 · popcorn · pop_l123_all_gpt
correct 26_ImportanceSampling.py
model
grok-4.3
turns / max
6 / 15
tool calls
6
tokens (total)
33181
wall clock
606.7s
truncated
no
backend
cuda fp32
variant
popcorn
runtime
0.18 μs
ref runtime
0.43 μs
speedup
2.38x
started
2026-05-11T17:36:19Z
timingkernelbench.timing
kernel runtime
0.18 μs
ref runtime
0.43 μs
speedup
2.38x
kernel mean / median / std
0.18 / — / 0.02 μs
ref mean / median / std
0.43 / — / 0.06 μs
numerical precisioneval_kernel_against_ref
max abs error
7.63e-06
mean abs error
1.08e-06
max rel error
1.06e-05
mean rel error
1.15e-06
gpu memoryextended_metrics.measure_memory
kernel peak
33.54 MB
ref peak
35.51 MB
kernel / ref
0.94x
kernel launches / fusiontorch.profiler
kernel launches
37
ref launches
88
fusion ratio (ref / kernel)
2.38x
total cuda time
1429 μs
SOL (speed-of-light)nsight_hardware_counters
SOL score (0–1)
0.486
dram util %
6.3%
compute util %
48.6%
bottleneck
compute-bound
dominant pipe
FP32
arithmetic intensity
—
ridge point
20.00
roofline / occupancynsight_hardware_counters
dram bandwidth
210.96 GB/s
dram util %
6.3%
fp32 tflops
—
fp32 util %
—
fp16 tflops
—
occupancy %
174.0%
bottleneck
compute-bound
dominant pipe %
48.6%
arithmetic intensity
—
ridge point
20.00
L1 hit %
—
L2 hit %
1789.0%
ld sectors / req
138.67
st sectors / req
—
regs / thread
585
smem / block
28160 B
block size
3712
peak bw
3350.00 GB/s
peak fp32 tflops
67.00
peak fp16 tflops
1979.00
warp stalls (top 5)nsight
long_scoreboard
13841.1%
short_scoreboard
2431.3%
mio_throttle
1.9%
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