FW-GLM-5-1 · popcorn · pop_l123_st_gpt
compile_fail 2_FusedMultiStepDiffusionSolver.py
model
FW-GLM-5-1
turns / max
1 / 1
tool calls
1
tokens (total)
7524
wall clock
80.8s
truncated
yes
backend
cuda fp32
variant
popcorn
runtime
—
ref runtime
-1.00 μs
speedup
—
started
2026-05-11T09:35:50Z
compile error: /scratch/abaner/KernelBench/runs/pop_l123_st_gpt/popcorn/FW-GLM-5-1/level_3_problem_2_cache/k_c18d1a92d8e4/fused_diffusion_kernel/fused_diffusion_kernel.so: cannot open shared object file: No such file or directory
timingkernelbench.timing
kernel runtime
—
ref runtime
—
speedup
—
kernel mean / median / std
—
ref mean / median / std
—
reference PyTorch source (the prompt the agent sees)
import torch
import torch.nn as nn
import torch.nn.functional as F

# Default ModelNew is PyTorch-only. run_and_check: use check_kernel=False with backend=cuda
# until the candidate file contains a real CUDA kernel (see kernel_static_checker).


def _linear_beta_schedule(num_timesteps: int, beta_start: float, beta_end: float) -> torch.Tensor:
    return torch.linspace(beta_start, beta_end, num_timesteps)


class Model(nn.Module):
    """
    Fused **multi-step** diffusion denoising (VP-DDPM-style) in a single ``forward``.

    Instead of one scheduler step per Python call, this reference unrolls
    ``num_fused_steps`` consecutive **deterministic** reverse updates (posterior mean,
    no sampler noise)—a typical fusion target for robotics / world-model diffusion
    heads where latency matters.

    Each inner step: predict ``ε_θ(x, t)`` with a small MLP, then apply the standard
    DDPM mean update from discrete time ``t`` to ``t-1``. Batch elements may start
    from different ``t_start``; inputs must satisfy ``t_start >= num_fused_steps`` so
    indices stay valid.
    """

    def __init__(self, dim: int, num_fused_steps: int, total_timesteps: int = 1000):
        super().__init__()
        assert num_fused_steps >= 1
        assert total_timesteps >= num_fused_steps + 1
        self.dim = dim
        self.num_fused_steps = num_fused_steps
        self.total_timesteps = total_timesteps

        self.eps_net = nn.Sequential(
            nn.Linear(dim + 1, dim * 4),
            nn.SiLU(),
            nn.Linear(dim * 4, dim * 4),
            nn.SiLU(),
            nn.Linear(dim * 4, dim),
        )

        betas = _linear_beta_schedule(total_timesteps, 1e-4, 0.02)
        alphas = 1.0 - betas
        alphas_cumprod = torch.cumprod(alphas, dim=0)
        self.register_buffer("betas", betas)
        self.register_buffer("alphas", alphas)
        self.register_buffer("alphas_cumprod", alphas_cumprod)
        # x_{t-1} = (1/sqrt(alpha_t)) * (x_t - (beta_t / sqrt(1-alpha_bar_t)) * eps)
        self.register_buffer("sqrt_recip_alphas", torch.sqrt(1.0 / alphas))
        self.register_buffer(
            "coef_eps",
            betas / torch.sqrt(1.0 - alphas_cumprod),
        )

    def forward(self, x: torch.Tensor, t_start: torch.Tensor) -> torch.Tensor:
        """
        Args:
            x: (B, dim) noisy latent at discrete time ``t_start`` (same semantics as DDPM).
            t_start: (B,) int64 timesteps in ``[num_fused_steps, total_timesteps - 1]``.
        Returns:
            (B, dim) latent after ``num_fused_steps`` fused denoise steps toward 0.
        """
        B, d = x.shape
        assert d == self.dim
        assert t_start.shape == (B,)
        # Harness paths (measure_ref_program_time / eval) cast activations to fp32;
        # discrete timesteps must survive as whole-valued floats → long indices.
        t_start = t_start.long()

        x_cur = x
        for k in range(self.num_fused_steps):
            t_idx = t_start - k
            # ``get_inputs`` samples t_start >= num_fused_steps so t_idx >= 1 always.
            t_norm = t_idx.float() / float(self.total_timesteps)
            eps = self.eps_net(torch.cat([x_cur, t_norm.unsqueeze(-1)], dim=-1))

            c1 = self.sqrt_recip_alphas[t_idx].view(B, 1)
            c2 = self.coef_eps[t_idx].view(B, 1)
            x_cur = c1 * (x_cur - c2 * eps)

        return x_cur


class ModelNew(nn.Module):
    """KernelBench candidate entry point; default delegates to ``Model``."""

    def __init__(self, dim: int, num_fused_steps: int, total_timesteps: int = 1000):
        super().__init__()
        self._impl = Model(dim, num_fused_steps, total_timesteps=total_timesteps)

    def forward(self, x: torch.Tensor, t_start: torch.Tensor) -> torch.Tensor:
        return self._impl(x, t_start)


batch_size = 8
dim = 64
num_fused_steps = 16
total_timesteps = 1000


def get_inputs():
    p = popcorn_pri
    x = torch.randn(p.jitter_int(batch_size), dim)
    low = num_fused_steps
    high = total_timesteps
    t_start = torch.randint(low, high, (p.jitter_int(batch_size),), dtype=torch.long)
    return [x, t_start]


def get_init_inputs():
    return [dim, num_fused_steps, total_timesteps]
turn 0 · FINAL · 1 tool call(s) · 80.8s turn · 68.7s LLM · tokens in=1,909 · out=5,615 · total=7,524 · 27,522 chars (click to expand)
turn 080.8s turn · 68.7s LLM · 12.1s tools · tokens in=1,909 · out=5,615 · total=7,524 · FINAL · reasoning clipped
reasoning
reasoning (truncated 2731 chars): Let me analyze this task carefully. I need to create a custom CUDA kernel for a multi-step diffusion denoising operation. Let me understand the reference implementation: 1. The model has an epsilon network (MLP) that predicts noise 2. For each of `num_fused_steps` steps, it: - Gets the timestep index `t_idx = t_start - k` - Normalizes it: `t_norm = t_idx / total_timesteps` - Predicts noise using the MLP: `eps = eps_net(concat(x_cur, t_norm))` - Applies DDPM mean update: `x_cur = c1 * (x_cur - c2 * eps)` where `c1 = sqrt_recip_alphas[t_idx]` and `c2 = coef_eps[t_idx]` The key operations to fuse: 1. MLP forward pass (Linear + SiLU + Linear + SiLU + Linear) 2. Coefficient lookup and application 3. The iterative loop over `num_fused_steps` To make this faster, I should: 1. Fuse the entire multi-step loop into a single kernel 2. Fuse the MLP operations (Linear + SiLU activation) 3. Minimize memory transfers Let me think about the dimensions: - Batch size: 8 - Dim: 64 - num_fused_steps: 16 - total_timesteps: 1000 The MLP structure: - Input: (dim + 1) = 65 - Hidden 1: dim * 4 = 256 (with SiLU) - Hidden 2: dim * 4 = 256 (with SiLU) - Output: dim = 64 For each step, we need to: 1. Normalize t_idx: `t_norm = t_idx / total_timesteps` 2. Concatenate x_cur and t_norm 3. Pass through MLP 4. Look up coefficients 5. Update x_cur 6. Leveraging shared memory for intermediate MLP results could significantly reduce global memory traffic, especially given the small dimensions. I'm considering a custom CUDA implementation that fuses the entire multi-step process into a single kernel, optimizing memory usage and computational efficiency for this specific diffusion denoising operation. The computation strategy involves processing batch elements across multiple diffusion steps, with careful attention to memory access patterns and coefficient retrieval. Matrix-vector multiplications within the MLP layers will be key to performance optimization. I'll focus on maximizing parallelism by ensuring each thread block handles a single batch element, allowing for complete step processing while maintaining efficient memory utilization. The compact dimensions (batch size 8, dimension 64) suggest a targeted, specialized kernel approach will yield the most significant performance gains. The implementation will prioritize coalesced memory access and minimize data movement between global and shared memory spaces during the denoising process. I'll continue the partial thought about CUDA kernel implementation, focusing on the forward function and kernel launch details. The code will complete the MLP network operations with a specialized CUDA kernel for efficient batch processing of the denoising steps. The forward method prepares CUDA kernel launch parameters, setting up grid and block dimensions to process the input tensor across multiple diffusion steps. It ensures the input tensor is contiguous on the GPU and calls a custom kernel that handles the entire denoising process in a single operation. The kernel will leverage parallel computation to process batch elements simultaneously, optimizing memory access and computational efficiency for the multiefficient MLP denoising network. I'll normalize the time index by dividing by total timesteps. The forward kernel prepares to process each batch element, with threads collaborating across the specified dimension. The kernel seems designed to handle a sequential processing of fused steps, with careful thread and block configuration to manage the computational load efficiently. The next section appears to begin handling batch element processing, focusing on the first fusion step and preparing for time index normalization. I continue processing the SiLU activation for the first layer's hidden layer, calculating the sigmoid and applying the element-wise multiplication. I then prepare for the second linear layer's matrix multiplication, loading the appropriate weights from the weight tensor. The computation continues with a partial matrix multiplication operation, carefully tracking the thread index and weight offsets. The next iteration begins, accumulating the matrix multiplication result by multiplying the hidden layer activation with the corresponding weight, adding to the running sum of the output. This process systematically computes the linear transformation across the layer's weights and hidden activations. The method progressively builds the output through iterative dot product calculations, leveraging the hidden layer's sigmoid-weighted linear unit (SiLU) activation and the predefined weight matrix for the layer. I'll complete the partial thought by adding the next thinking to finish the current computation sequence. Final processing for hidden1 involves writing the accumulated sum with bias, preparing for subsequent layer transformations. The kernel then advances to processing subsequent linear layers, systematically computing matrix multiplications across hidden representations while maintaining thread-level parallelism. The computation continues with loading weights from subsequent layer matrices, initiating dot product calculations for the next neural network layer transformation. I'm computing matrix-vector multiplication efficiently by leveraging shared memory for intermediate activations and iterating through weight dimensions with parallel thread processing. The computation follows a standard neural network layer transformation, accumulating weighted sums across layer weights. Each thread contributes to the parallel computation by calculating partial dot products, progressively building the layer's output representation through systematic weight and activation interactions. I'll continue the thought by processing the next part of the linear layer computation, focusing on writing the output and preparing for the subsequent steps: Completing the output computation for the second linear layer, I write the result to the output buffer. The kernel then moves into a prediction step, retrieving time indices and preparing for a scaling operation. The process involves carefully managing thread synchronization and preparing data for subsequent transformations across multiple fused steps. The computation involves calculating scaled values using precomputed coefficients, with each thread handling a specific dimension index. The method ensures efficient parallel processing while maintaining data consistency across different time steps and prediction stages. I'll launch the kernel with carefully configured grid and block dimensions, ensuring the output tensor is initialized and returned. The kernel will process the input data across multiple time steps, applying the denoising coefficients and weights efficiently on the GPU. The launch parameters are set to match the batch size and dimension, with a specified number of threads per block. After kernel execution, the output tensor is returned, representing the denoised result of the multi-step diffusion process. Optimization strategies include leveraging shared memory for faster intermediate computations, maintaining coalesced memory access patterns, and minimizing synchronization overhead. The kernel design explores potential performance improvements through careful thread block configuration and parallel computation strategies. The computational complexity involves significant matrix transformations across multiple layers, with weight dimensions carefully managed to balance memory usage and computational efficiency. Careful consideration of shared memory allocation becomes critical for optimal kernel performance. Efficient thread block configuration remains a key challenge, with potential approaches exploring thread allocation strategies that maximize parallel processing while respecting hardware constraints. The initial implementation leverages standard CUDA optimization techniques like vectorized memory access and strategic shared memory utilization, ensur ...[ omitted 2731 characters ]... ach element in the shared memory array, transforming the hidden representation before the next layer's computation. This prepares the data for subsequent linear transformations and feature extraction. The forward method prepares for a CUDA kernel launch, setting up the grid and block dimensions based on the input tensor's batch size and dimensionality. It ensures the input is contiguous on CUDA devices, preparing for efficient parallel computation across multiple fused steps. The kernel will leverage shared memory for storing intermediate layer weights and activations, optimizing memory access and computational efficiency during the neural network's processing. I'll need to verify the weight tensor dimensions and ensure proper memory allocation for the CUDA kernel parameters. The MLP architecture suggests specific weight shapes that must align with the kernel's computational requirements. Weight tensor dimensions confirm: - First linear layer: 65x256 - Second linear layer: 256x256 - Third linear layer: 256x64 The forward method appears to be preparing a CUDA kernel launch, with grid and block dimensions being configured for parallel processing. Shared memory allocation seems carefully calculated for each layer's weights and activations, ensuring efficient GPU memory utilization. The kernel is set to use 256 threads per block, optimizing computational performance across different network dimensions. Parameter validation confirms expected tensor shapes: input data matches batch and dimension specifications, timestep start tensor aligns with batch size. The implementation ensures contiguous CUDA memory for input tensor, preparing for parallel processing. I'll carefully track coefficient tensor dimensions, noting potential broadcasting challenges with timestep indexing. The sqrt reciprocal alphas and coefficient tensors require precise handling during kernel operations. The view reshaping method transforms 1D coefficient tensors into broadcastable formats, enabling element-wise multiplication across batch dimensions. Kernel implementation will need manual broadcasting logic to ensure correct computational flow. The coefficient tensors have specific shapes matching total timestep count, requiring careful selection and expansion during GPU processing. I'll need to verify coefficient application matches PyTorch's tensor manipulation strategy, particularly focusing on how scalar values propagate through batch computations. The element-wise broadcasting mechanism ensures consistent coefficient scaling across each dimension, allowing precise timestep-dependent transformations without explicit reshaping. When accessing coefficients from a contiguous tensor, thread-level indexing enables parallel coefficient retrieval. Each thread independently loads its coefficient, enabling efficient GPU computation while maintaining consistent coefficient application across batch dimensions. The implementation guarantees memory layout compatibility, with coefficients strategically positioned at precise tensor offsets for each timestep index. I'll verify potential numerical precision challenges involving SiLU activation and potential floating-point variability across different computational implementations. To ensure robustness, I'll explicitly handle edge cases like dimension divisibility and carefully manage CUDA kernel launch configurations. The approach involves comprehensive parameter validation and creating a flexible implementation that can accommodate varied input tensor characteristics. The strategy includes developing a resilient kernel that dynamically adapts to input tensor shapes, with careful attention to thread synchronization and memory access patterns. I'm optimizing the kernel for typical scenarios while maintaining flexibility for edge cases. The standard configuration of dim=64 aligns well with thread block sizing, allowing efficient parallel processing. My implementation will focus on performance for common use cases, accepting potential performance variations for less typical input dimensions. The key is balancing specialized optimization with broader compatibility, ensuring the kernel performs well across most expected input scenarios. I'll carefully examine the weight matrix storage and computation strategy. The weight matrices are stored with standard PyTorch conventions: W1 at (256, 65), W2 at (256, 256), and W3 at (64, 256). For matrix-vector multiplication, we compute output by multiplying transposed weights with the input vector. The dot product calculation involves iterating through input dimensions, using row-major storage to efficiently access weight elements. Each output element is computed by summing weighted contributions across input dimensions. The index calculation ensures precise memory access, mapping thread and input indices to the correct weight matrix locations. The weight indexing checks out perfectly. For the second linear layer, I'll verify the dot product calculation with row and column indices. The matrix multiplication logic follows the same pattern, ensuring each weight element is correctly accessed during the computation. I'll continue processing the second layer's computation systematically. The accumulation method ensures each thread computes its specific output element efficiently. The shared memory approach allows collaborative loading of hidden layer values, minimizing redundant memory accesses. This strategy optimizes parallel computation across threads. The weight indexing mechanism enables precise matrix-vector multiplication by carefully mapping thread indices to weight matrix coordinates. By leveraging shared memory and cooperative loading, the kernel maximizes computational efficiency for the linear layer transformation. The parallel computation strategy allows each thread to independently calculate dot products, accumulating partial results across hidden layer dimensions. The loop structure ensures comprehensive matrix multiplication while maintaining thread-level independence. I'm carefully tracking weight matrix offsets and verifying dimensional alignment between network layers. The weight storage requires precise indexing to ensure correct computational mapping across network transformations. The intermediate weight matrices connect network stages with specific dimensionality transitions, requiring meticulous offset calculations to maintain computational integrity. I've calculated the weight offsets and verified the weight shapes for the neural network layers. The biases are confirmed to be 1D tensors matching the output dimensions. I'm tracking the MLP computation steps, noting how the input concatenates the current state and normalized time, then passes through linear layers with SiLU activations. The computation seems to be mid-description when the cutoff occurs. The key insight is the step iteration process: each step updates the current state using predicted noise from the neural network, gradually transforming the input over multiple iterations. I'll implement a custom CUDA extension that loads the noise prediction kernel inline, preparing to accelerate the multi-step noise prediction process efficiently. The forward method prepares CUDA kernel launch configurations, mapping batch and dimension parameters to grid and block dimensions. It ensures input tensors are contiguous and device-compatible, preparing for parallel noise prediction computation. Shared memory allocation enables efficient data sharing across thread blocks during noise prediction operations. The kernel launch synchronizes to guarantee computational completion before returning processed results. The implementation strategically uses inline CUDA compilation, optimizing performance through custom kernel design and careful memory management. I'm considering vectorized memory access techniques to potentially accelerate data loading and storing operations, which could improve overall computational efficiency.
call → submit_kernel
callsubmit_kernel { "kernel_code": "import torch\nimport torch.nn as nn\nimport torch.utils.cpp_extension as cpp_ext\n\n# Compile CUDA kernel at module load time\nfused_diffusion_ext = cpp_ext.load_inline(\n name='fused_diffusion_kernel',\n cpp_sources=r'''\n #include <torch/extension.h>\n #include <cuda.h>\n #include <cuda_runtime.h>\n\n __device__ __forceinline__ float silu(float x) {\n return x * (1.0f / (1.0f + expf(-x)));\n }\n\n __global__ void fused_diffusion_kernel(\n const float* __restrict__ x,\n const int64_t* __restrict__ t_start,\n const float* __restrict__ w1, const float* __restrict__ b1,\n const float* __restrict__ w2, const float* __restrict__ b2,\n const float* __restrict__ w3, const float* __restrict__ b3,\n const float* __restrict__ sqrt_recip_alphas,\n const float* __restrict__ coef_eps,\n float* __restrict__ output,\n int batch_size, int dim, int num_fused_steps, int total_timesteps\n ) {\n int batch_idx = blockIdx.x;\n int dim_idx = threadIdx.x;\n \n if (batch_idx >= batch_size) return;\n \n int base_t = t_start[batch_idx];\n \n // Load input x for this batch element into registers\n // dim = 64, blockDim = 64, each thread loads one element\n float x_cur;\n if (dim_idx < dim) {\n x_cur = x[batch_idx * dim + dim_idx];\n }\n \n // Loop over fusion steps\n for (int step = 0; step < num_fused_steps; step++) {\n int t_idx = base_t - step;\n float t_norm = float(t_idx) / float(total_timesteps);\n \n // MLP: Linear1 (65 -> 256) + SiLU\n float h1[4] = {b1[0], b1[1], b1[2], b1[3]}; // Start with bias\n \n // Each thread handles 4 consecutive output elements\n // Process input: [x_cur (dim=64), t_norm (1)]\n // Thread dim_idx handles output elements [dim_idx*4, dim_idx*4+1, dim_idx*4+2, dim_idx*4+3]\n \n // Dot product for 4 output neurons\n // Weight matrix is (256, 65) - row-major\n // Output element j uses row j of weights\n \n for (int i = 0; i < 64; i++) {\n // Broadcast x_cur[i] from thread i to all threads\n // Use warp shuffle for broadcast\n float x_i = __shfl_sync(0xffffffff, x_cur, i);\n \n // Each thread computes contribution to its 4 output elements\n #pragma unroll\n for (int k = 0; k < 4; k++) {\n int j = dim_idx * 4 + k; // output index\n if (j < 256) {\n h1[k] += w1[j * 65 + i] * x_i;\n }\n }\n }\n \n // Add t_norm contribution\n #pragma unroll\n for (int k = 0; k < 4; k++) {\n int j = dim_idx * 4 + k;\n if (j < 256) {\n h1[k] += w1[j * 65 + 64] * t_norm;\n }\n }\n \n // Apply SiLU\n #pragma unroll\n for (int k = 0; k < 4; k++) {\n h1[k] = silu(h1[k]);\n }\n \n // MLP: Linear2 (256 -> 256) + SiLU\n float h2[4] = {b2[dim_idx*4], b2[dim_idx*4+1], b2[dim_idx*4+2], b2[dim_idx*4+3]};\n \n // For each of our 4 output elements, compute dot product with all 256 inputs\n for (int i = 0; i < 256; i++) {\n // Get h1[i] from the thread that computed it\n // Thread (i/4) computed h1[i%4]\n int src_thread = i / 4;\n int src_idx = i % 4;\n \n // Broadcast h1[i] to all threads\n float h1_i;\n if (src_idx == 0) h1_i = __shfl_sync(0xffffffff, h1[0], src_thread);\n else if (src_idx == 1) h1_i = __shfl_sync(0xffffffff, h1[1], src_thread);\n else if (src_idx == 2) h1_i = __shfl_sync(0xffffffff, h1[2], src_thread);\n else h1_i = __shfl_sync(0xffffffff, h1[3], src_thread);\n \n #pragma unroll\n for (int k = 0; k < 4; k++) {\n int j = dim_idx * 4 + k;\n if (j < 256) {\n h2[k] += w2[j * 256 + i] * h1_i;\n }\n }\n }\n \n // Apply SiLU\n #pragma unroll\n for (int k = 0; k < 4; k++) {\n h2[k] = silu(h2[k]);\n }\n \n // MLP: Linear3 (256 -> 64)\n float eps = b3[dim_idx];\n \n for (int i = 0; i < 256; i++) {\n int src_thread = i / 4;\n int src_idx = i % 4;\n \n float h2_i;\n if (src_idx == 0) h2_i = __shfl_sync(0xffffffff, h2[0], src_thread);\n else if (src_idx == 1) h2_i = __shfl_sync(0xffffffff, h2[1], src_thread);\n else if (src_idx == 2) h2_i = __shfl_sync(0xffffffff, h2[2], src_thread);\n else h2_i = __shfl_sync(0xffffffff, h2[3], src_thread);\n \n eps += w3[dim_idx * 256 + i] * h2_i;\n }\n \n // Apply denoising step: x_cur = c1 * (x_cur - c2 * eps)\n float c1 = sqrt_recip_alphas[t_idx];\n float c2 = coef_eps[t_idx];\n \n x_cur = c1 * (x_cur - c2 * eps);\n }\n \n // Write output\n if (dim_idx < dim) {\n output[batch_idx * dim + dim_idx] = x_cur;\n }\n }\n\n torch::Tensor fused_diffusion_forward(\n torch::Tensor x,\n torch::Tensor t_start,\n torch::Tensor w1, torch::Tensor b1,\n torch::Tensor w2, torch::Tensor b2,\n torch::Tensor w3, torch::Tensor b3,\n torch::Tensor sqrt_recip_alphas,\n torch::Tensor coef_eps,\n int num_fused_steps,\n int total_timesteps\n ) {\n int batch_size = x.size(0);\n int dim = x.size(1);\n \n auto output = torch::empty_like(x);\n \n int threads_per_block = dim; // 64 threads per block\n int blocks = batch_size;\n \n fused_diffusion_kernel<<<blocks, threads_per_block>>>(\n x.data_ptr<float>(),\n t_start.data_ptr<int64_t>(),\n w1.data_ptr<float>(), b1.data_ptr<float>(),\n w2.data_ptr<float>(), b2.data_ptr<float>(),\n w3.data_ptr<float>(), b3.data_ptr<float>(),\n sqrt_recip_alphas.data_ptr<float>(),\n coef_eps.data_ptr<float>(),\n output.data_ptr<float>(),\n batch_size, dim, num_fused_steps, total_timesteps\n );\n \n return output;\n }\n ''',\n cuda_sources='',\n functions=['fused_diffusion_forward'],\n with_cuda=True,\n extra_cuda_cflags=['-O3', '-use_fast_math'],\n)\n\n\nclass ModelNew(nn.Module):\n def __init__(self, dim: int, num_fused_steps: int, total_timesteps: int = 1000):\n super().__init__()\n self.dim = dim\n self.num_fused_steps = num_fused_steps\n self.total_timesteps = total_timesteps\n \n # Create the same eps_net as reference\n self.eps_net = nn.Sequential(\n nn.Linear(dim + 1, dim * 4),\n nn.SiLU(),\n nn.Linear(dim * 4, dim * 4),\n nn.SiLU(),\n nn.Linear(dim * 4, dim),\n )\n \n # Create the same schedule buffers\n betas = torch.linspace(1e-4, 0.02, total_timesteps)\n alphas = 1.0 - betas\n alphas_cumprod = torch.cumprod(alphas, dim=0)\n \n self.register_buffer(\"betas\", betas)\n self.register_buffer(\"alphas\", alphas)\n self.register_buffer(\"alphas_cumprod\", alphas_cumprod)\n self.register_buffer(\"sqrt_recip_alphas\", torch.sqrt(1.0 / alphas))\n self.register_buffer(\"coef_eps\", betas / torch.sqrt(1.0 - alphas_cumprod))\n \n def forward(self, x: torch.Tensor, t_start: torch.Tensor) -> torch.Tensor:\n # Fall back to reference if not on CUDA or dim doesn't match expected\n if not x.is_cuda or self.dim != 64:\n # Reference implementation\n B, d = x.shape\n t_start = t_start.long()\n x_cur = x\n for k in range(self.num_fused_steps):\n t_idx = t_start - k\n t_norm = t_idx.float() / float(self.total_timesteps)\n eps = self.eps_net(torch.cat([x_cur, t_norm.unsqueeze(-1)], dim=-1))\n c1 = self.sqrt_recip_alphas[t_idx].view(B, 1)\n c2 = self.coef_eps[t_idx].view(B, 1)\n x_cur = c1 * (x_cur - c2 * eps)\n return x_cur\n \n # Extract weights from eps_net\n # eps_net structure: Linear, SiLU, Linear, SiLU, Linear\n # Indices: 0=Linear1, 2=Linear2, 4=Linear3\n w1 = self.eps_net[0].weight # (256, 65)\n b1 = self.eps_net[0].bias # (256,)\n w2 = self.eps_net[2].weight # (256, 256)\n b2 = self.eps_net[2].bias # (256,)\n w3 = self.eps_net[4].weight # (64, 256)\n b3 = self.eps_net[4].bias # (64,)\n \n # Ensure t_start is long\n t_start = t_start.long()\n \n # Call the fused CUDA kernel\n output = fused_diffusion_ext.fused_diffusion_forward(\n x.contiguous(),\n t_start.contiguous(),\n w1.contiguous(), b1.contiguous(),\n w2.contiguous(), b2.contiguous(),\n w3.contiguous(), b3.contiguous(),\n self.sqrt_recip_alphas.contiguous(),\n self.coef_eps.contiguous(),\n self.num_fused_steps,\n self.total_timesteps\n )\n \n return output" }
result ← submit_kernel (fail)
failsubmit_kernel submit_kernel FAILED: kernel did not compile. /scratch/abaner/KernelBench/runs/pop_l123_st_gpt/popcorn/FW-GLM-5-1/level_3_problem_2_cache/k_c18d1a92d8e4/fused_diffusion_kernel/fused_diffusion_kernel.so: cannot open shared object file: No such file or directory
final submitted kernel (259 lines, 9,740 chars) — full source, untruncated
import torch
import torch.nn as nn
import torch.utils.cpp_extension as cpp_ext

# Compile CUDA kernel at module load time
fused_diffusion_ext = cpp_ext.load_inline(
    name='fused_diffusion_kernel',
    cpp_sources=r'''
    #include <torch/extension.h>
    #include <cuda.h>
    #include <cuda_runtime.h>

    __device__ __forceinline__ float silu(float x) {
        return x * (1.0f / (1.0f + expf(-x)));
    }

    __global__ void fused_diffusion_kernel(
        const float* __restrict__ x,
        const int64_t* __restrict__ t_start,
        const float* __restrict__ w1, const float* __restrict__ b1,
        const float* __restrict__ w2, const float* __restrict__ b2,
        const float* __restrict__ w3, const float* __restrict__ b3,
        const float* __restrict__ sqrt_recip_alphas,
        const float* __restrict__ coef_eps,
        float* __restrict__ output,
        int batch_size, int dim, int num_fused_steps, int total_timesteps
    ) {
        int batch_idx = blockIdx.x;
        int dim_idx = threadIdx.x;
        
        if (batch_idx >= batch_size) return;
        
        int base_t = t_start[batch_idx];
        
        // Load input x for this batch element into registers
        // dim = 64, blockDim = 64, each thread loads one element
        float x_cur;
        if (dim_idx < dim) {
            x_cur = x[batch_idx * dim + dim_idx];
        }
        
        // Loop over fusion steps
        for (int step = 0; step < num_fused_steps; step++) {
            int t_idx = base_t - step;
            float t_norm = float(t_idx) / float(total_timesteps);
            
            // MLP: Linear1 (65 -> 256) + SiLU
            float h1[4] = {b1[0], b1[1], b1[2], b1[3]}; // Start with bias
            
            // Each thread handles 4 consecutive output elements
            // Process input: [x_cur (dim=64), t_norm (1)]
            // Thread dim_idx handles output elements [dim_idx*4, dim_idx*4+1, dim_idx*4+2, dim_idx*4+3]
            
            // Dot product for 4 output neurons
            // Weight matrix is (256, 65) - row-major
            // Output element j uses row j of weights
            
            for (int i = 0; i < 64; i++) {
                // Broadcast x_cur[i] from thread i to all threads
                // Use warp shuffle for broadcast
                float x_i = __shfl_sync(0xffffffff, x_cur, i);
                
                // Each thread computes contribution to its 4 output elements
                #pragma unroll
                for (int k = 0; k < 4; k++) {
                    int j = dim_idx * 4 + k; // output index
                    if (j < 256) {
                        h1[k] += w1[j * 65 + i] * x_i;
                    }
                }
            }
            
            // Add t_norm contribution
            #pragma unroll
            for (int k = 0; k < 4; k++) {
                int j = dim_idx * 4 + k;
                if (j < 256) {
                    h1[k] += w1[j * 65 + 64] * t_norm;
                }
            }
            
            // Apply SiLU
            #pragma unroll
            for (int k = 0; k < 4; k++) {
                h1[k] = silu(h1[k]);
            }
            
            // MLP: Linear2 (256 -> 256) + SiLU
            float h2[4] = {b2[dim_idx*4], b2[dim_idx*4+1], b2[dim_idx*4+2], b2[dim_idx*4+3]};
            
            // For each of our 4 output elements, compute dot product with all 256 inputs
            for (int i = 0; i < 256; i++) {
                // Get h1[i] from the thread that computed it
                // Thread (i/4) computed h1[i%4]
                int src_thread = i / 4;
                int src_idx = i % 4;
                
                // Broadcast h1[i] to all threads
                float h1_i;
                if (src_idx == 0) h1_i = __shfl_sync(0xffffffff, h1[0], src_thread);
                else if (src_idx == 1) h1_i = __shfl_sync(0xffffffff, h1[1], src_thread);
                else if (src_idx == 2) h1_i = __shfl_sync(0xffffffff, h1[2], src_thread);
                else h1_i = __shfl_sync(0xffffffff, h1[3], src_thread);
                
                #pragma unroll
                for (int k = 0; k < 4; k++) {
                    int j = dim_idx * 4 + k;
                    if (j < 256) {
                        h2[k] += w2[j * 256 + i] * h1_i;
                    }
                }
            }
            
            // Apply SiLU
            #pragma unroll
            for (int k = 0; k < 4; k++) {
                h2[k] = silu(h2[k]);
            }
            
            // MLP: Linear3 (256 -> 64)
            float eps = b3[dim_idx];
            
            for (int i = 0; i < 256; i++) {
                int src_thread = i / 4;
                int src_idx = i % 4;
                
                float h2_i;
                if (src_idx == 0) h2_i = __shfl_sync(0xffffffff, h2[0], src_thread);
                else if (src_idx == 1) h2_i = __shfl_sync(0xffffffff, h2[1], src_thread);
                else if (src_idx == 2) h2_i = __shfl_sync(0xffffffff, h2[2], src_thread);
                else h2_i = __shfl_sync(0xffffffff, h2[3], src_thread);
                
                eps += w3[dim_idx * 256 + i] * h2_i;
            }
            
            // Apply denoising step: x_cur = c1 * (x_cur - c2 * eps)
            float c1 = sqrt_recip_alphas[t_idx];
            float c2 = coef_eps[t_idx];
            
            x_cur = c1 * (x_cur - c2 * eps);
        }
        
        // Write output
        if (dim_idx < dim) {
            output[batch_idx * dim + dim_idx] = x_cur;
        }
    }

    torch::Tensor fused_diffusion_forward(
        torch::Tensor x,
        torch::Tensor t_start,
        torch::Tensor w1, torch::Tensor b1,
        torch::Tensor w2, torch::Tensor b2,
        torch::Tensor w3, torch::Tensor b3,
        torch::Tensor sqrt_recip_alphas,
        torch::Tensor coef_eps,
        int num_fused_steps,
        int total_timesteps
    ) {
        int batch_size = x.size(0);
        int dim = x.size(1);
        
        auto output = torch::empty_like(x);
        
        int threads_per_block = dim; // 64 threads per block
        int blocks = batch_size;
        
        fused_diffusion_kernel<<<blocks, threads_per_block>>>(
            x.data_ptr<float>(),
            t_start.data_ptr<int64_t>(),
            w1.data_ptr<float>(), b1.data_ptr<float>(),
            w2.data_ptr<float>(), b2.data_ptr<float>(),
            w3.data_ptr<float>(), b3.data_ptr<float>(),
            sqrt_recip_alphas.data_ptr<float>(),
            coef_eps.data_ptr<float>(),
            output.data_ptr<float>(),
            batch_size, dim, num_fused_steps, total_timesteps
        );
        
        return output;
    }
    ''',
    cuda_sources='',
    functions=['fused_diffusion_forward'],
    with_cuda=True,
    extra_cuda_cflags=['-O3', '-use_fast_math'],
)


class ModelNew(nn.Module):
    def __init__(self, dim: int, num_fused_steps: int, total_timesteps: int = 1000):
        super().__init__()
        self.dim = dim
        self.num_fused_steps = num_fused_steps
        self.total_timesteps = total_timesteps
        
        # Create the same eps_net as reference
        self.eps_net = nn.Sequential(
            nn.Linear(dim + 1, dim * 4),
            nn.SiLU(),
            nn.Linear(dim * 4, dim * 4),
            nn.SiLU(),
            nn.Linear(dim * 4, dim),
        )
        
        # Create the same schedule buffers
        betas = torch.linspace(1e-4, 0.02, total_timesteps)
        alphas = 1.0 - betas
        alphas_cumprod = torch.cumprod(alphas, dim=0)
        
        self.register_buffer("betas", betas)
        self.register_buffer("alphas", alphas)
        self.register_buffer("alphas_cumprod", alphas_cumprod)
        self.register_buffer("sqrt_recip_alphas", torch.sqrt(1.0 / alphas))
        self.register_buffer("coef_eps", betas / torch.sqrt(1.0 - alphas_cumprod))
    
    def forward(self, x: torch.Tensor, t_start: torch.Tensor) -> torch.Tensor:
        # Fall back to reference if not on CUDA or dim doesn't match expected
        if not x.is_cuda or self.dim != 64:
            # Reference implementation
            B, d = x.shape
            t_start = t_start.long()
            x_cur = x
            for k in range(self.num_fused_steps):
                t_idx = t_start - k
                t_norm = t_idx.float() / float(self.total_timesteps)
                eps = self.eps_net(torch.cat([x_cur, t_norm.unsqueeze(-1)], dim=-1))
                c1 = self.sqrt_recip_alphas[t_idx].view(B, 1)
                c2 = self.coef_eps[t_idx].view(B, 1)
                x_cur = c1 * (x_cur - c2 * eps)
            return x_cur
        
        # Extract weights from eps_net
        # eps_net structure: Linear, SiLU, Linear, SiLU, Linear
        # Indices: 0=Linear1, 2=Linear2, 4=Linear3
        w1 = self.eps_net[0].weight   # (256, 65)
        b1 = self.eps_net[0].bias     # (256,)
        w2 = self.eps_net[2].weight   # (256, 256)
        b2 = self.eps_net[2].bias     # (256,)
        w3 = self.eps_net[4].weight   # (64, 256)
        b3 = self.eps_net[4].bias     # (64,)
        
        # Ensure t_start is long
        t_start = t_start.long()
        
        # Call the fused CUDA kernel
        output = fused_diffusion_ext.fused_diffusion_forward(
            x.contiguous(),
            t_start.contiguous(),
            w1.contiguous(), b1.contiguous(),
            w2.contiguous(), b2.contiguous(),
            w3.contiguous(), b3.contiguous(),
            self.sqrt_recip_alphas.contiguous(),
            self.coef_eps.contiguous(),
            self.num_fused_steps,
            self.total_timesteps
        )
        
        return output