grok-4.3 · popcorn · pop_l123_default_gpt
incorrect 48_ParticleFilter.py
model
grok-4.3
turns / max
10 / 12
tool calls
10
tokens (total)
78603
wall clock
3082.5s
truncated
no
backend
cuda fp32
variant
popcorn
runtime
—
ref runtime
0.56 μs
speedup
—
started
2026-05-11T10:47:16Z
correctness issue: Output mismatch
timingkernelbench.timing
kernel runtime
—
ref runtime
0.56 μs
speedup
—
kernel mean / median / std
—
ref mean / median / std
0.56 / — / 0.07 μs
roofline / occupancyheuristic
dram bandwidth
—
dram util %
—
fp32 tflops
—
fp32 util %
—
fp16 tflops
—
occupancy %
—
bottleneck
—
dominant pipe %
—
arithmetic intensity
—
ridge point
—
L1 hit %
—
L2 hit %
—
ld sectors / req
—
st sectors / req
—
regs / thread
—
smem / block
—
block size
—
peak bw
—
peak fp32 tflops
—
peak fp16 tflops
—
reference PyTorch source (the prompt the agent sees)
import torch
import torch.nn as nn


class Model(nn.Module):
    """
    Bootstrap particle filter (Sequential Monte Carlo) for a linear-
    Gaussian state-space model.  Implements the predict–weight–resample
    cycle for one time step.  The state transition is x_{t} = A x_{t-1} + w
    and the observation model is y_t = C x_t + v.
    """

    def __init__(self, state_dim, obs_dim, num_particles):
        super().__init__()
        self.state_dim = state_dim
        self.obs_dim = obs_dim
        self.num_particles = num_particles
        self.register_buffer("A", 0.9 * torch.eye(state_dim))
        self.register_buffer("C", torch.randn(obs_dim, state_dim) * 0.5)
        self.process_noise = 0.1
        self.obs_noise = 0.5

    def forward(
        self,
        particles: torch.Tensor,
        log_weights: torch.Tensor,
        observation: torch.Tensor,
    ) -> tuple:
        """
        Args:
            particles:   (B, P, state_dim)
            log_weights: (B, P)
            observation: (B, obs_dim)
        Returns:
            (new_particles, new_log_weights):
                new_particles:   (B, P, state_dim)
                new_log_weights: (B, P) – normalized log weights
        """
        B, P, D = particles.shape

        # Predict: propagate through dynamics
        pred = (particles @ self.A.t()) + self.process_noise * torch.randn_like(particles)

        # Weight: compute log p(y_t | x_t)
        obs_pred = pred @ self.C.t()  # (B, P, obs_dim)
        diff = observation.unsqueeze(1) - obs_pred
        log_lik = -0.5 * (diff ** 2).sum(-1) / (self.obs_noise ** 2)
        new_log_w = log_weights + log_lik
        # Normalize
        new_log_w = new_log_w - torch.logsumexp(new_log_w, dim=1, keepdim=True)

        # Systematic resample
        weights = torch.exp(new_log_w)
        cumsum = torch.cumsum(weights, dim=1)
        u = (torch.arange(P, device=particles.device, dtype=particles.dtype) + torch.rand(B, 1, device=particles.device)) / P
        indices = torch.searchsorted(cumsum, u)
        indices = indices.clamp(0, P - 1)

        batch_idx = torch.arange(B, device=particles.device).unsqueeze(1).expand(-1, P)
        resampled = pred[batch_idx, indices]
        uniform_log_w = torch.full_like(new_log_w, -torch.log(torch.tensor(float(P))))

        return resampled, uniform_log_w


state_dim = 4
obs_dim = 2
num_particles = 1024
batch_size = 16


def get_inputs():
    p = popcorn_pri
    particles = torch.randn(p.jitter_int(batch_size), num_particles, state_dim)
    log_weights = torch.full((p.jitter_int(batch_size), num_particles), -torch.log(torch.tensor(float(num_particles))))
    observation = torch.randn(p.jitter_int(batch_size), obs_dim)
    return [particles, log_weights, observation]


def get_init_inputs():
    return [state_dim, obs_dim, num_particles]
turn 05.3s turn · 5.3s LLM · 0.0s tools · tokens in=1,944 · out=6 · total=2,152
reasoning
The task is to replace the forward computation of a PyTorch model for a Bootstrap particle filter with a custom CUDA kernel. The kernel must be faster than the reference while producing numerically eq...
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 1104.7s turn · 18.8s LLM · 86.0s tools · tokens in=2,385 · out=607 · total=3,210
reasoning
GPU is H100, high bandwidth, lots of compute.
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 <curand_kernel.h>\n\n__global__ void particle_filter_kernel(\n const float* __restrict__ particles,\n const float* __restrict__ log_weights,\n const float* __restrict__ observation,\n const float* __restrict__ A,\n const float* __restrict__ C,\n float* __restrict__ new_particles,\n float* __restrict__ new_log_weights,\n int B, int P, int D, int O,\n float process_noise, float obs_noise,\n unsigned long long seed\n) {\n int b = blockIdx.x;\n int p = threadIdx.x + blockIdx.y * blockDim.x;\n if (p >= P) return;\n\n curandState state;\n curand_init(seed, b * P + p, 0, &state);\n\n // Predict: pred = particles[b,p] @ A.T + noise\n float pred[4]; // assume D<=4\n for(int d=0; d<D; d++) {\n float sum = 0.0f;\n for(int k=0; k<D; k++) {\n sum += particles[(b*P + p)*D + k] * A[k*D + d];\n }\n float noise = process_noise * curand_normal(&state);\n pred[d] = sum + noise;\n }\n\n // obs_pred = pred @ C.T\n float obs_pred[2];\n for(int o=0; o<O; o++) {\n float sum = 0.0f;\n for(int d=0; d<D; d++) {\n sum += pred[d] * C[d*O + o];\n }\n obs_pred[o] = sum;\n }\n\n // log_lik\n float log_lik = 0.0f;\n float inv_var = 1.0f / (obs_noise * obs_noise);\n for(int o=0; o<O; o++) {\n float diff = observation[b*O + o] - obs_pred[o];\n log_lik -= 0.5f * diff * diff * inv_var;\n }\n\n float lw = log_weights[b*P + p] + log_lik;\n\n // For now, store pred and lw temporarily? We need global ops.\n // This kernel is incomplete for normalization and resampling.\n // Will need multiple kernels or shared mem reductions.\n}\n'''\n\ncpp_source = '''\ntorch::Tensor dummy(torch::Tensor x) { return x; }\n'''\n\nmodule = load_inline(\n name='particle_filter',\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 2343.3s turn · 18.5s LLM · 324.8s tools · tokens in=3,235 · out=1,009 · total=4,500
reasoning
It compiled but no ModelNew. I need to provide the full 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 <curand_kernel.h>\n\n__global__ void predict_weight_kernel(\n const float* particles,\n const float* log_weights,\n const float* obs,\n const float* A,\n const float* C,\n float* pred_particles,\n float* log_liks,\n int B, int P, int D, int O,\n float proc_noise, float obs_noise,\n unsigned long long seed\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * P;\n if (idx >= total) return;\n int b = idx / P;\n int p = idx % P;\n\n curandStatePhilox4_32_10_t state;\n curand_init(seed, idx, 0, &state);\n\n float pred[4];\n for (int d = 0; d < D; d++) {\n float s = 0.0f;\n for (int k = 0; k < D; k++) {\n s += particles[idx * D + k] * A[k * D + d];\n }\n pred[d] = s + proc_noise * curand_normal(&state);\n }\n\n for (int d = 0; d < D; d++) {\n pred_particles[idx * D + d] = pred[d];\n }\n\n float ll = 0.0f;\n float inv2var = -0.5f / (obs_noise * obs_noise);\n for (int o = 0; o < O; o++) {\n float op = 0.0f;\n for (int d = 0; d < D; d++) {\n op += pred[d] * C[d * O + o];\n }\n float diff = obs[b * O + o] - op;\n ll += diff * diff;\n }\n log_liks[idx] = inv2var * ll + log_weights[idx];\n}\n\ntorch::Tensor predict_weight_cuda(torch::Tensor particles, torch::Tensor log_weights, torch::Tensor obs, torch::Tensor A, torch::Tensor C, float proc_noise, float obs_noise, unsigned long long seed) {\n int B = particles.size(0);\n int P = particles.size(1);\n int D = particles.size(2);\n int O = obs.size(1);\n auto pred = torch::empty({B, P, D}, particles.options());\n auto ll = torch::empty({B, P}, particles.options());\n int total = B * P;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n predict_weight_kernel<<<blocks, threads>>>(\n particles.data_ptr<float>(),\n log_weights.data_ptr<float>(),\n obs.data_ptr<float>(),\n A.data_ptr<float>(),\n C.data_ptr<float>(),\n pred.data_ptr<float>(),\n ll.data_ptr<float>(),\n B, P, D, O, proc_noise, obs_noise, seed\n );\n return pred;\n}\n'''\n\ncpp_source = '''\ntorch::Tensor predict_weight_cuda(torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor, float, float, unsigned long long);\n'''\n\nmodule = load_inline(\n name='pf_ext',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['predict_weight_cuda'],\n verbose=True\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, state_dim, obs_dim, num_particles):\n super().__init__()\n self.state_dim = state_dim\n self.obs_dim = obs_dim\n self.num_particles = num_particles\n self.register_buffer(\"A\", 0.9 * torch.eye(state_dim))\n self.register_buffer(\"C\", torch.randn(obs_dim, state_dim) * 0.5)\n self.process_noise = 0.1\n self.obs_noise = 0.5\n self.ext = module\n\n def forward(self, particles, log_weights, observation):\n B, P, D = particles.shape\n pred = self.ext.predict_weight_cuda(\n particles, log_weights, observation, self.A, self.C,\n self.process_noise, self.obs_noise, torch.randint(0, 2**63-1, (1,), device='cpu').item()\n )\n # incomplete, need to finish in Python for now\n obs_pred = pred @ self.C.t()\n # etc, but to make correct, better to use full ref for now? No.\n # This is partial.\n return particles, log_weights # dummy\nprint(\"Setup done\")\n" }
result ← compile_kernel (ok)
okcompile_kernel compile_kernel PASSED: kernel compiled without errors.
turn 3389.0s turn · 12.0s LLM · 377.0s tools · tokens in=4,517 · out=1,021 · total=5,727
reasoning
Good, it passed compile.
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 <curand_kernel.h>\n\n__global__ void predict_weight_kernel(\n const float* __restrict__ particles,\n const float* __restrict__ log_weights,\n const float* __restrict__ obs,\n const float* __restrict__ A,\n const float* __restrict__ C,\n float* __restrict__ pred_particles,\n float* __restrict__ new_log_w,\n int B, int P, int D, int O,\n float proc_noise, float obs_noise,\n unsigned long long seed\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * P;\n if (idx >= total) return;\n int b = idx / P;\n int p = idx % P;\n\n curandStatePhilox4_32_10_t state;\n curand_init(seed + idx, 0, 0, &state);\n\n float pred[4];\n #pragma unroll\n for (int d = 0; d < D; d++) {\n float s = 0.0f;\n #pragma unroll\n for (int k = 0; k < D; k++) {\n s += particles[idx * D + k] * A[k * D + d];\n }\n pred[d] = s + proc_noise * curand_normal(&state);\n }\n\n for (int d = 0; d < D; d++) {\n pred_particles[idx * D + d] = pred[d];\n }\n\n float ll = 0.0f;\n float inv2var = -0.5f / (obs_noise * obs_noise);\n for (int o = 0; o < O; o++) {\n float op = 0.0f;\n for (int d = 0; d < D; d++) {\n op += pred[d] * C[d * O + o];\n }\n float diff = obs[b * O + o] - op;\n ll += diff * diff;\n }\n new_log_w[idx] = inv2var * ll + log_weights[idx];\n}\n'''\n\ncpp_source = '''\nvoid predict_weight_cuda(torch::Tensor particles, torch::Tensor log_weights, torch::Tensor obs,\n torch::Tensor A, torch::Tensor C, torch::Tensor pred_particles, torch::Tensor new_log_w,\n float proc_noise, float obs_noise, unsigned long long seed);\n'''\n\nmodule = load_inline(\n name='pf_ext',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['predict_weight_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, state_dim, obs_dim, num_particles):\n super().__init__()\n self.state_dim = state_dim\n self.obs_dim = obs_dim\n self.num_particles = num_particles\n self.register_buffer(\"A\", 0.9 * torch.eye(state_dim))\n self.register_buffer(\"C\", torch.randn(obs_dim, state_dim) * 0.5)\n self.process_noise = 0.1\n self.obs_noise = 0.5\n self.ext = module\n\n def forward(self, particles, log_weights, observation):\n B, P, D = particles.shape\n O = observation.shape[1]\n pred = torch.empty_like(particles)\n new_log_w = torch.empty_like(log_weights)\n seed = int(torch.randint(0, 2**63 - 1, (1,), device='cpu').item())\n module.predict_weight_cuda(\n particles.contiguous(), log_weights.contiguous(), observation.contiguous(),\n self.A.contiguous(), self.C.contiguous(), pred, new_log_w,\n self.process_noise, self.obs_noise, seed\n )\n # normalize\n new_log_w = new_log_w - torch.logsumexp(new_log_w, dim=1, keepdim=True)\n # resample\n weights = torch.exp(new_log_w)\n cumsum = torch.cumsum(weights, dim=1)\n u = (torch.arange(P, device=particles.device, dtype=particles.dtype).unsqueeze(0) + torch.rand(B, 1, device=particles.device)) / P\n indices = torch.searchsorted(cumsum, u, right=False)\n indices = indices.clamp(0, P-1)\n batch_idx = torch.arange(B, device=particles.device).unsqueeze(1)\n resampled = pred[batch_idx, indices]\n uniform_log_w = torch.full_like(new_log_w, -torch.log(torch.tensor(float(P), device=particles.device)))\n return resampled, uniform_log_w\n" }
result ← compile_kernel (fail)
failcompile_kernel compile_kernel FAILED: ImportError. ImportError: /scratch/abaner/KernelBench/runs/pop_l123_default_gpt/popcorn/grok-4.3/level_2_problem_48_cache/k_7792c7c0bb23/pf_ext/pf_ext_v1.so: undefined symbol: _Z19predict_weight_cudaN2at6TensorES0_S0_S0_S0_S0_S0_ffy
turn 4289.9s turn · 13.8s LLM · 276.1s tools · tokens in=5,821 · out=1,210 · total=7,047
reasoning
Error because cpp declaration mismatch, need to match exactly.
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 <curand_kernel.h>\n\n__global__ void predict_weight_kernel(\n const float* __restrict__ particles,\n const float* __restrict__ log_weights,\n const float* __restrict__ obs,\n const float* __restrict__ A,\n const float* __restrict__ C,\n float* __restrict__ pred_particles,\n float* __restrict__ new_log_w,\n int B, int P, int D, int O,\n float proc_noise, float obs_noise,\n unsigned long long seed\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * P;\n if (idx >= total) return;\n int b = idx / P;\n int p = idx % P;\n\n curandStatePhilox4_32_10_t state;\n curand_init(seed + idx, 0, 0, &state);\n\n float pred[4];\n for (int d = 0; d < D; d++) {\n float s = 0.0f;\n for (int k = 0; k < D; k++) {\n s += particles[idx * D + k] * A[k * D + d];\n }\n pred[d] = s + proc_noise * curand_normal(&state);\n }\n\n for (int d = 0; d < D; d++) {\n pred_particles[idx * D + d] = pred[d];\n }\n\n float ll = 0.0f;\n float inv2var = -0.5f / (obs_noise * obs_noise);\n for (int o = 0; o < O; o++) {\n float op = 0.0f;\n for (int d = 0; d < D; d++) {\n op += pred[d] * C[d * O + o];\n }\n float diff = obs[b * O + o] - op;\n ll += diff * diff;\n }\n new_log_w[idx] = inv2var * ll + log_weights[idx];\n}\n\nvoid predict_weight_cuda(torch::Tensor particles, torch::Tensor log_weights, torch::Tensor obs,\n torch::Tensor A, torch::Tensor C, torch::Tensor pred_particles, torch::Tensor new_log_w,\n float proc_noise, float obs_noise, unsigned long long seed) {\n int B = particles.size(0);\n int P = particles.size(1);\n int D = particles.size(2);\n int O = obs.size(1);\n int total = B * P;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n predict_weight_kernel<<<blocks, threads>>>(\n particles.data_ptr<float>(),\n log_weights.data_ptr<float>(),\n obs.data_ptr<float>(),\n A.data_ptr<float>(),\n C.data_ptr<float>(),\n pred_particles.data_ptr<float>(),\n new_log_w.data_ptr<float>(),\n B, P, D, O, proc_noise, obs_noise, seed\n );\n}\n'''\n\ncpp_source = '''\nvoid predict_weight_cuda(torch::Tensor particles, torch::Tensor log_weights, torch::Tensor obs,\n torch::Tensor A, torch::Tensor C, torch::Tensor pred_particles, torch::Tensor new_log_w,\n float proc_noise, float obs_noise, unsigned long long seed);\n'''\n\nmodule = load_inline(\n name='pf_ext',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['predict_weight_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, state_dim, obs_dim, num_particles):\n super().__init__()\n self.state_dim = state_dim\n self.obs_dim = obs_dim\n self.num_particles = num_particles\n self.register_buffer(\"A\", 0.9 * torch.eye(state_dim))\n self.register_buffer(\"C\", torch.randn(obs_dim, state_dim) * 0.5)\n self.process_noise = 0.1\n self.obs_noise = 0.5\n self.ext = module\n\n def forward(self, particles, log_weights, observation):\n B, P, D = particles.shape\n O = observation.shape[1]\n pred = torch.empty_like(particles)\n new_log_w = torch.empty_like(log_weights)\n seed = int(torch.randint(0, 2**63 - 1, (1,), device='cpu').item())\n module.predict_weight_cuda(\n particles.contiguous(), log_weights.contiguous(), observation.contiguous(),\n self.A.contiguous(), self.C.contiguous(), pred, new_log_w,\n self.process_noise, self.obs_noise, seed\n )\n # normalize\n new_log_w = new_log_w - torch.logsumexp(new_log_w, dim=1, keepdim=True)\n # resample\n weights = torch.exp(new_log_w)\n cumsum = torch.cumsum(weights, dim=1)\n u = (torch.arange(P, device=particles.device, dtype=particles.dtype).unsqueeze(0) + torch.rand(B, 1, device=particles.device)) / P\n indices = torch.searchsorted(cumsum, u, right=False)\n indices = indices.clamp(0, P-1)\n batch_idx = torch.arange(B, device=particles.device).unsqueeze(1)\n resampled = pred[batch_idx, indices]\n uniform_log_w = torch.full_like(new_log_w, -torch.log(torch.tensor(float(P), device=particles.device)))\n return resampled, uniform_log_w\n" }
result ← compile_kernel (ok)
okcompile_kernel compile_kernel PASSED: kernel compiled without errors.
turn 5238.6s turn · 10.1s LLM · 228.5s tools · tokens in=7,064 · out=1,211 · total=8,281
reasoning
Compiled. 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 = r'''\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n#include <curand_kernel.h>\n\n__global__ void predict_weight_kernel(\n const float* __restrict__ particles,\n const float* __restrict__ log_weights,\n const float* __restrict__ obs,\n const float* __restrict__ A,\n const float* __restrict__ C,\n float* __restrict__ pred_particles,\n float* __restrict__ new_log_w,\n int B, int P, int D, int O,\n float proc_noise, float obs_noise,\n unsigned long long seed\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * P;\n if (idx >= total) return;\n int b = idx / P;\n int p = idx % P;\n\n curandStatePhilox4_32_10_t state;\n curand_init(seed + idx, 0, 0, &state);\n\n float pred[4];\n for (int d = 0; d < D; d++) {\n float s = 0.0f;\n for (int k = 0; k < D; k++) {\n s += particles[idx * D + k] * A[k * D + d];\n }\n pred[d] = s + proc_noise * curand_normal(&state);\n }\n\n for (int d = 0; d < D; d++) {\n pred_particles[idx * D + d] = pred[d];\n }\n\n float ll = 0.0f;\n float inv2var = -0.5f / (obs_noise * obs_noise);\n for (int o = 0; o < O; o++) {\n float op = 0.0f;\n for (int d = 0; d < D; d++) {\n op += pred[d] * C[d * O + o];\n }\n float diff = obs[b * O + o] - op;\n ll += diff * diff;\n }\n new_log_w[idx] = inv2var * ll + log_weights[idx];\n}\n\nvoid predict_weight_cuda(torch::Tensor particles, torch::Tensor log_weights, torch::Tensor obs,\n torch::Tensor A, torch::Tensor C, torch::Tensor pred_particles, torch::Tensor new_log_w,\n float proc_noise, float obs_noise, unsigned long long seed) {\n int B = particles.size(0);\n int P = particles.size(1);\n int D = particles.size(2);\n int O = obs.size(1);\n int total = B * P;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n predict_weight_kernel<<<blocks, threads>>>(\n particles.data_ptr<float>(),\n log_weights.data_ptr<float>(),\n obs.data_ptr<float>(),\n A.data_ptr<float>(),\n C.data_ptr<float>(),\n pred_particles.data_ptr<float>(),\n new_log_w.data_ptr<float>(),\n B, P, D, O, proc_noise, obs_noise, seed\n );\n}\n'''\n\ncpp_source = '''\nvoid predict_weight_cuda(torch::Tensor particles, torch::Tensor log_weights, torch::Tensor obs,\n torch::Tensor A, torch::Tensor C, torch::Tensor pred_particles, torch::Tensor new_log_w,\n float proc_noise, float obs_noise, unsigned long long seed);\n'''\n\nmodule = load_inline(\n name='pf_ext',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['predict_weight_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, state_dim, obs_dim, num_particles):\n super().__init__()\n self.state_dim = state_dim\n self.obs_dim = obs_dim\n self.num_particles = num_particles\n self.register_buffer(\"A\", 0.9 * torch.eye(state_dim))\n self.register_buffer(\"C\", torch.randn(obs_dim, state_dim) * 0.5)\n self.process_noise = 0.1\n self.obs_noise = 0.5\n self.ext = module\n\n def forward(self, particles, log_weights, observation):\n B, P, D = particles.shape\n O = observation.shape[1]\n pred = torch.empty_like(particles)\n new_log_w = torch.empty_like(log_weights)\n seed = int(torch.randint(0, 2**63 - 1, (1,), device='cpu').item())\n module.predict_weight_cuda(\n particles.contiguous(), log_weights.contiguous(), observation.contiguous(),\n self.A.contiguous(), self.C.contiguous(), pred, new_log_w,\n self.process_noise, self.obs_noise, seed\n )\n # normalize\n new_log_w = new_log_w - torch.logsumexp(new_log_w, dim=1, keepdim=True)\n # resample\n weights = torch.exp(new_log_w)\n cumsum = torch.cumsum(weights, dim=1)\n u = (torch.arange(P, device=particles.device, dtype=particles.dtype).unsqueeze(0) + torch.rand(B, 1, device=particles.device)) / P\n indices = torch.searchsorted(cumsum, u, right=False)\n indices = indices.clamp(0, P-1)\n batch_idx = torch.arange(B, device=particles.device).unsqueeze(1)\n resampled = pred[batch_idx, indices]\n uniform_log_w = torch.full_like(new_log_w, -torch.log(torch.tensor(float(P), device=particles.device)))\n return resampled, uniform_log_w\n" }
result ← run_correctness (fail)
failrun_correctness run_correctness FAILED: (0 / 5) trials did not all match. correctness_issue: Output mismatch max_difference: ['5.223844', '6.109830', '5.177494', '5.621154', '5.560101'] avg_difference: ['1.042122', '1.033018', '1.017828', '1.057938', '1.000882']
turn 6246.7s turn · 20.2s LLM · 226.5s tools · tokens in=8,378 · out=924 · total=9,718
reasoning
Mismatch because the C is different (random in init), and also rand different, but the problem requires equivalent, meaning same computation, but since stochastic, the reference also stochastic. The original model's forward pass uses random.randn for particles and observation, and a complex particle filter computation involving A, C, and logsumexp.
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\n__global__ void weight_kernel(\n const float* __restrict__ pred,\n const float* __restrict__ log_weights,\n const float* __restrict__ obs,\n const float* __restrict__ C,\n float* __restrict__ new_log_w,\n int B, int P, int D, int O,\n float obs_noise\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * P;\n if (idx >= total) return;\n int b = idx / P;\n\n float ll = 0.0f;\n float inv2var = -0.5f / (obs_noise * obs_noise);\n for (int o = 0; o < O; o++) {\n float op = 0.0f;\n for (int d = 0; d < D; d++) {\n op += pred[idx * D + d] * C[d * O + o];\n }\n float diff = obs[b * O + o] - op;\n ll += diff * diff;\n }\n new_log_w[idx] = inv2var * ll + log_weights[idx];\n}\n\nvoid weight_cuda(torch::Tensor pred, torch::Tensor log_weights, torch::Tensor obs,\n torch::Tensor C, torch::Tensor new_log_w, float obs_noise) {\n int B = pred.size(0);\n int P = pred.size(1);\n int D = pred.size(2);\n int O = obs.size(1);\n int total = B * P;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n weight_kernel<<<blocks, threads>>>(\n pred.data_ptr<float>(),\n log_weights.data_ptr<float>(),\n obs.data_ptr<float>(),\n C.data_ptr<float>(),\n new_log_w.data_ptr<float>(),\n B, P, D, O, obs_noise\n );\n}\n'''\n\ncpp_source = '''\nvoid weight_cuda(torch::Tensor pred, torch::Tensor log_weights, torch::Tensor obs,\n torch::Tensor C, torch::Tensor new_log_w, float obs_noise);\n'''\n\nmodule = load_inline(\n name='pf_ext',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['weight_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, state_dim, obs_dim, num_particles):\n super().__init__()\n self.state_dim = state_dim\n self.obs_dim = obs_dim\n self.num_particles = num_particles\n self.register_buffer(\"A\", 0.9 * torch.eye(state_dim))\n self.register_buffer(\"C\", torch.randn(obs_dim, state_dim) * 0.5)\n self.process_noise = 0.1\n self.obs_noise = 0.5\n self.ext = module\n\n def forward(self, particles, log_weights, observation):\n B, P, D = particles.shape\n # Predict using PyTorch for rand match\n pred = (particles @ self.A.t()) + self.process_noise * torch.randn_like(particles)\n # Weight with CUDA\n new_log_w = torch.empty_like(log_weights)\n module.weight_cuda(pred, log_weights, observation, self.C, new_log_w, self.obs_noise)\n # normalize\n new_log_w = new_log_w - torch.logsumexp(new_log_w, dim=1, keepdim=True)\n # resample\n weights = torch.exp(new_log_w)\n cumsum = torch.cumsum(weights, dim=1)\n u = (torch.arange(P, device=particles.device, dtype=particles.dtype).unsqueeze(0) + torch.rand(B, 1, device=particles.device)) / P\n indices = torch.searchsorted(cumsum, u, right=False)\n indices = indices.clamp(0, P-1)\n batch_idx = torch.arange(B, device=particles.device).unsqueeze(1)\n resampled = pred[batch_idx, indices]\n uniform_log_w = torch.full_like(new_log_w, -torch.log(torch.tensor(float(P), device=particles.device)))\n return resampled, uniform_log_w\n" }
result ← compile_kernel (ok)
okcompile_kernel compile_kernel PASSED: kernel compiled without errors.
turn 7459.5s turn · 25.9s LLM · 433.6s tools · tokens in=9,735 · out=925 · total=10,664
reasoning
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 = r'''\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n\n__global__ void weight_kernel(\n const float* __restrict__ pred,\n const float* __restrict__ log_weights,\n const float* __restrict__ obs,\n const float* __restrict__ C,\n float* __restrict__ new_log_w,\n int B, int P, int D, int O,\n float obs_noise\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * P;\n if (idx >= total) return;\n int b = idx / P;\n\n float ll = 0.0f;\n float inv2var = -0.5f / (obs_noise * obs_noise);\n for (int o = 0; o < O; o++) {\n float op = 0.0f;\n for (int d = 0; d < D; d++) {\n op += pred[idx * D + d] * C[d * O + o];\n }\n float diff = obs[b * O + o] - op;\n ll += diff * diff;\n }\n new_log_w[idx] = inv2var * ll + log_weights[idx];\n}\n\nvoid weight_cuda(torch::Tensor pred, torch::Tensor log_weights, torch::Tensor obs,\n torch::Tensor C, torch::Tensor new_log_w, float obs_noise) {\n int B = pred.size(0);\n int P = pred.size(1);\n int D = pred.size(2);\n int O = obs.size(1);\n int total = B * P;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n weight_kernel<<<blocks, threads>>>(\n pred.data_ptr<float>(),\n log_weights.data_ptr<float>(),\n obs.data_ptr<float>(),\n C.data_ptr<float>(),\n new_log_w.data_ptr<float>(),\n B, P, D, O, obs_noise\n );\n}\n'''\n\ncpp_source = '''\nvoid weight_cuda(torch::Tensor pred, torch::Tensor log_weights, torch::Tensor obs,\n torch::Tensor C, torch::Tensor new_log_w, float obs_noise);\n'''\n\nmodule = load_inline(\n name='pf_ext',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['weight_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, state_dim, obs_dim, num_particles):\n super().__init__()\n self.state_dim = state_dim\n self.obs_dim = obs_dim\n self.num_particles = num_particles\n self.register_buffer(\"A\", 0.9 * torch.eye(state_dim))\n self.register_buffer(\"C\", torch.randn(obs_dim, state_dim) * 0.5)\n self.process_noise = 0.1\n self.obs_noise = 0.5\n self.ext = module\n\n def forward(self, particles, log_weights, observation):\n B, P, D = particles.shape\n # Predict using PyTorch for rand match\n pred = (particles @ self.A.t()) + self.process_noise * torch.randn_like(particles)\n # Weight with CUDA\n new_log_w = torch.empty_like(log_weights)\n module.weight_cuda(pred, log_weights, observation, self.C, new_log_w, self.obs_noise)\n # normalize\n new_log_w = new_log_w - torch.logsumexp(new_log_w, dim=1, keepdim=True)\n # resample\n weights = torch.exp(new_log_w)\n cumsum = torch.cumsum(weights, dim=1)\n u = (torch.arange(P, device=particles.device, dtype=particles.dtype).unsqueeze(0) + torch.rand(B, 1, device=particles.device)) / P\n indices = torch.searchsorted(cumsum, u, right=False)\n indices = indices.clamp(0, P-1)\n batch_idx = torch.arange(B, device=particles.device).unsqueeze(1)\n resampled = pred[batch_idx, indices]\n uniform_log_w = torch.full_like(new_log_w, -torch.log(torch.tensor(float(P), device=particles.device)))\n return resampled, uniform_log_w\n" }
result ← run_correctness (fail)
failrun_correctness run_correctness FAILED: (0 / 5) trials did not all match. correctness_issue: Output mismatch max_difference: ['5.147470', '6.136529', '4.875571', '5.494964', '5.540575'] avg_difference: ['1.040792', '1.029413', '1.010985', '1.050331', '0.992598']
turn 8 · 1 tool call(s) · 270.7s turn · 39.1s LLM · tokens in=10,761 · out=932 · total=12,936 · 6,161 chars (click to expand)
turn 8270.7s turn · 39.1s LLM · 231.6s tools · tokens in=10,761 · out=932 · total=12,936
reasoning
Still mismatch. Why? The pred = particles @ A.t() is same, but then + noise randn_like, same. The issue is that the `pred = particles @ A.t()` and subsequent steps like adding noise and computing log-likelihoods are identical in both the reference `Model` and my custom `ModelNew`. The mismatch arises because each instance initializes its own random `C` matrix independently, leading to different `C` values when @ A.t() is computed, which affects the observation prediction and thus the outputs. The correctness check creates both models separately, so their `C` matrices differ, causing numerical mismatches despite identical computations. To resolve this, setting a seed in `__init__` for `ModelNew` won't work because the reference `model model has a different `C`. The best approach is to ensure both initials happen under the same random state, but since they're instantiated sequentially, one consumes the RNG state while the other gets a new sequence. A practical solution is to hardcode deterministic `A` and `C` values instead of generating them randomly, ensuring numerically equivalent outputs by making initialization deterministic and matching the computation logic exactly to the reference's deterministic parts, ignoring randomness for correctness testing purposes, or assuming the check tolerates minor differences due to RNG, but given the diff is large, that's unlikely.
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\n__global__ void weight_kernel(\n const float* __restrict__ pred,\n const float* __restrict__ log_weights,\n const float* __restrict__ obs,\n const float* __restrict__ C,\n float* __restrict__ new_log_w,\n int B, int P, int D, int O,\n float obs_noise\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * P;\n if (idx >= total) return;\n int b = idx / P;\n\n float ll = 0.0f;\n float inv2var = -0.5f / (obs_noise * obs_noise);\n for (int o = 0; o < O; o++) {\n float op = 0.0f;\n for (int d = 0; d < D; d++) {\n op += pred[idx * D + d] * C[d * O + o];\n }\n float diff = obs[b * O + o] - op;\n ll += diff * diff;\n }\n new_log_w[idx] = inv2var * ll + log_weights[idx];\n}\n\nvoid weight_cuda(torch::Tensor pred, torch::Tensor log_weights, torch::Tensor obs,\n torch::Tensor C, torch::Tensor new_log_w, float obs_noise) {\n int B = pred.size(0);\n int P = pred.size(1);\n int D = pred.size(2);\n int O = obs.size(1);\n int total = B * P;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n weight_kernel<<<blocks, threads>>>(\n pred.data_ptr<float>(),\n log_weights.data_ptr<float>(),\n obs.data_ptr<float>(),\n C.data_ptr<float>(),\n new_log_w.data_ptr<float>(),\n B, P, D, O, obs_noise\n );\n}\n'''\n\ncpp_source = '''\nvoid weight_cuda(torch::Tensor pred, torch::Tensor log_weights, torch::Tensor obs,\n torch::Tensor C, torch::Tensor new_log_w, float obs_noise);\n'''\n\nmodule = load_inline(\n name='pf_ext',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['weight_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, state_dim, obs_dim, num_particles):\n super().__init__()\n torch.manual_seed(42)\n self.state_dim = state_dim\n self.obs_dim = obs_dim\n self.num_particles = num_particles\n self.register_buffer(\"A\", 0.9 * torch.eye(state_dim))\n self.register_buffer(\"C\", torch.randn(obs_dim, state_dim) * 0.5)\n self.process_noise = 0.1\n self.obs_noise = 0.5\n self.ext = module\n\n def forward(self, particles, log_weights, observation):\n B, P, D = particles.shape\n # Predict using PyTorch for rand match\n pred = (particles @ self.A.t()) + self.process_noise * torch.randn_like(particles)\n # Weight with CUDA\n new_log_w = torch.empty_like(log_weights)\n module.weight_cuda(pred, log_weights, observation, self.C, new_log_w, self.obs_noise)\n # normalize\n new_log_w = new_log_w - torch.logsumexp(new_log_w, dim=1, keepdim=True)\n # resample\n weights = torch.exp(new_log_w)\n cumsum = torch.cumsum(weights, dim=1)\n u = (torch.arange(P, device=particles.device, dtype=particles.dtype).unsqueeze(0) + torch.rand(B, 1, device=particles.device)) / P\n indices = torch.searchsorted(cumsum, u, right=False)\n indices = indices.clamp(0, P-1)\n batch_idx = torch.arange(B, device=particles.device).unsqueeze(1)\n resampled = pred[batch_idx, indices]\n uniform_log_w = torch.full_like(new_log_w, -torch.log(torch.tensor(float(P), device=particles.device)))\n return resampled, uniform_log_w\n" }
result ← run_correctness (fail)
failrun_correctness run_correctness FAILED: (0 / 5) trials did not all match. correctness_issue: Output mismatch max_difference: ['5.147470', '6.136529', '4.875571', '5.494964', '5.540575'] avg_difference: ['1.040792', '1.029413', '1.010985', '1.050331', '0.992598']
turn 9734.9s turn · 47.4s LLM · 687.6s tools · tokens in=13,067 · out=924 · total=14,368 · FINAL
reasoning
Still fail. The diff is same, because seed only affects ModelNew C. I’m still hitting a snag because the diff is the same, and the seed only impacts ModelNew’s C. To get C to match, I’d have to eliminate random calls in ModelNew, which isn’t feasible.
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\n__global__ void weight_kernel(\n const float* __restrict__ pred,\n const float* __restrict__ log_weights,\n const float* __restrict__ obs,\n const float* __restrict__ C,\n float* __restrict__ new_log_w,\n int B, int P, int D, int O,\n float obs_noise\n) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n int total = B * P;\n if (idx >= total) return;\n int b = idx / P;\n\n float ll = 0.0f;\n float inv2var = -0.5f / (obs_noise * obs_noise);\n for (int o = 0; o < O; o++) {\n float op = 0.0f;\n for (int d = 0; d < D; d++) {\n op += pred[idx * D + d] * C[d * O + o];\n }\n float diff = obs[b * O + o] - op;\n ll += diff * diff;\n }\n new_log_w[idx] = inv2var * ll + log_weights[idx];\n}\n\nvoid weight_cuda(torch::Tensor pred, torch::Tensor log_weights, torch::Tensor obs,\n torch::Tensor C, torch::Tensor new_log_w, float obs_noise) {\n int B = pred.size(0);\n int P = pred.size(1);\n int D = pred.size(2);\n int O = obs.size(1);\n int total = B * P;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n weight_kernel<<<blocks, threads>>>(\n pred.data_ptr<float>(),\n log_weights.data_ptr<float>(),\n obs.data_ptr<float>(),\n C.data_ptr<float>(),\n new_log_w.data_ptr<float>(),\n B, P, D, O, obs_noise\n );\n}\n'''\n\ncpp_source = '''\nvoid weight_cuda(torch::Tensor pred, torch::Tensor log_weights, torch::Tensor obs,\n torch::Tensor C, torch::Tensor new_log_w, float obs_noise);\n'''\n\nmodule = load_inline(\n name='pf_ext',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['weight_cuda'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, state_dim, obs_dim, num_particles):\n super().__init__()\n self.state_dim = state_dim\n self.obs_dim = obs_dim\n self.num_particles = num_particles\n self.register_buffer(\"A\", 0.9 * torch.eye(state_dim))\n self.register_buffer(\"C\", torch.randn(obs_dim, state_dim) * 0.5)\n self.process_noise = 0.1\n self.obs_noise = 0.5\n self.ext = module\n\n def forward(self, particles, log_weights, observation):\n B, P, D = particles.shape\n # Predict using PyTorch for rand match\n pred = (particles @ self.A.t()) + self.process_noise * torch.randn_like(particles)\n # Weight with CUDA\n new_log_w = torch.empty_like(log_weights)\n module.weight_cuda(pred, log_weights, observation, self.C, new_log_w, self.obs_noise)\n # normalize\n new_log_w = new_log_w - torch.logsumexp(new_log_w, dim=1, keepdim=True)\n # resample\n weights = torch.exp(new_log_w)\n cumsum = torch.cumsum(weights, dim=1)\n u = (torch.arange(P, device=particles.device, dtype=particles.dtype).unsqueeze(0) + torch.rand(B, 1, device=particles.device)) / P\n indices = torch.searchsorted(cumsum, u, right=False)\n indices = indices.clamp(0, P-1)\n batch_idx = torch.arange(B, device=particles.device).unsqueeze(1)\n resampled = pred[batch_idx, indices]\n uniform_log_w = torch.full_like(new_log_w, -torch.log(torch.tensor(float(P), device=particles.device)))\n return resampled, uniform_log_w\n" }
result ← submit_kernel (fail)
failsubmit_kernel submit_kernel FAILED: correctness check did not pass ((0 / 5) trials). correctness_issue: Output mismatch max_difference: ['5.147470', '6.136529', '4.875571', '5.494964', '5.540575'] avg_difference: ['1.040792', '1.029413', '1.010985', '1.050331', '0.992598']
final submitted kernel (100 lines, 3,432 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>

__global__ void weight_kernel(
    const float* __restrict__ pred,
    const float* __restrict__ log_weights,
    const float* __restrict__ obs,
    const float* __restrict__ C,
    float* __restrict__ new_log_w,
    int B, int P, int D, int O,
    float obs_noise
) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int total = B * P;
    if (idx >= total) return;
    int b = idx / P;

    float ll = 0.0f;
    float inv2var = -0.5f / (obs_noise * obs_noise);
    for (int o = 0; o < O; o++) {
        float op = 0.0f;
        for (int d = 0; d < D; d++) {
            op += pred[idx * D + d] * C[d * O + o];
        }
        float diff = obs[b * O + o] - op;
        ll += diff * diff;
    }
    new_log_w[idx] = inv2var * ll + log_weights[idx];
}

void weight_cuda(torch::Tensor pred, torch::Tensor log_weights, torch::Tensor obs,
                 torch::Tensor C, torch::Tensor new_log_w, float obs_noise) {
    int B = pred.size(0);
    int P = pred.size(1);
    int D = pred.size(2);
    int O = obs.size(1);
    int total = B * P;
    int threads = 256;
    int blocks = (total + threads - 1) / threads;
    weight_kernel<<<blocks, threads>>>(
        pred.data_ptr<float>(),
        log_weights.data_ptr<float>(),
        obs.data_ptr<float>(),
        C.data_ptr<float>(),
        new_log_w.data_ptr<float>(),
        B, P, D, O, obs_noise
    );
}
'''

cpp_source = '''
void weight_cuda(torch::Tensor pred, torch::Tensor log_weights, torch::Tensor obs,
                 torch::Tensor C, torch::Tensor new_log_w, float obs_noise);
'''

module = load_inline(
    name='pf_ext',
    cpp_sources=cpp_source,
    cuda_sources=cuda_source,
    functions=['weight_cuda'],
    verbose=False
)

class ModelNew(nn.Module):
    def __init__(self, state_dim, obs_dim, num_particles):
        super().__init__()
        self.state_dim = state_dim
        self.obs_dim = obs_dim
        self.num_particles = num_particles
        self.register_buffer("A", 0.9 * torch.eye(state_dim))
        self.register_buffer("C", torch.randn(obs_dim, state_dim) * 0.5)
        self.process_noise = 0.1
        self.obs_noise = 0.5
        self.ext = module

    def forward(self, particles, log_weights, observation):
        B, P, D = particles.shape
        # Predict using PyTorch for rand match
        pred = (particles @ self.A.t()) + self.process_noise * torch.randn_like(particles)
        # Weight with CUDA
        new_log_w = torch.empty_like(log_weights)
        module.weight_cuda(pred, log_weights, observation, self.C, new_log_w, self.obs_noise)
        # normalize
        new_log_w = new_log_w - torch.logsumexp(new_log_w, dim=1, keepdim=True)
        # resample
        weights = torch.exp(new_log_w)
        cumsum = torch.cumsum(weights, dim=1)
        u = (torch.arange(P, device=particles.device, dtype=particles.dtype).unsqueeze(0) + torch.rand(B, 1, device=particles.device)) / P
        indices = torch.searchsorted(cumsum, u, right=False)
        indices = indices.clamp(0, P-1)
        batch_idx = torch.arange(B, device=particles.device).unsqueeze(1)
        resampled = pred[batch_idx, indices]
        uniform_log_w = torch.full_like(new_log_w, -torch.log(torch.tensor(float(P), device=particles.device)))
        return resampled, uniform_log_w