call → run_correctness
callrun_correctness
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\n_cpp_src = r'''\n#include <torch/extension.h>\nvoid fused_ddpm_launch(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, torch::Tensor out, int num_steps, int total_timesteps);\ntorch::Tensor fused_ddpm_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_steps, int total_timesteps) {\n auto out = torch::empty_like(x);\n fused_ddpm_launch(x, t_start, w1, b1, w2, b2, w3, b3, sqrt_recip_alphas, coef_eps, out, num_steps, total_timesteps);\n return out;\n}\n'''\n\n_cuda_src = r'''\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <stdint.h>\n\n__inline__ __device__ float warp_sum(float v) { for (int off=16; off>0; off>>=1) v += __shfl_down_sync(0xffffffff, v, off); return v; }\n__device__ __forceinline__ float silu_f(float x) { return x / (1.0f + expf(-x)); }\n\n__global__ void fused_ddpm_kernel(const float* __restrict__ x_in, const int64_t* __restrict__ ts_l, const float* __restrict__ ts_f, int ts_is_float,\n const float* __restrict__ w1, const float* __restrict__ b1, const float* __restrict__ w2, const float* __restrict__ b2,\n const float* __restrict__ w3, const float* __restrict__ b3, const float* __restrict__ sqrt_recip_alphas,\n const float* __restrict__ coef_eps, float* __restrict__ out, int dim, int num_steps, int total_timesteps) {\n int b = blockIdx.x, tid = threadIdx.x, lane = tid & 31, warp = tid >> 5, nwarps = blockDim.x >> 5;\n extern __shared__ float smem[];\n float* x = smem; float* h1 = x + dim; float* h2 = h1 + 4*dim; float* eps = h2 + 4*dim;\n for (int i=tid; i<dim; i+=blockDim.x) x[i] = x_in[b*dim+i];\n __syncthreads();\n int hidden = 4*dim, in_dim = dim + 1;\n int ts = ts_is_float ? (int)ts_f[b] : (int)ts_l[b];\n for (int k=0; k<num_steps; ++k) {\n int t_idx = ts - k;\n float tnorm = (float)t_idx / (float)total_timesteps;\n for (int o=warp; o<hidden; o+=nwarps) {\n float sum = 0.0f; const float* ww = w1 + o*in_dim;\n for (int i=lane; i<dim; i+=32) sum = fmaf(x[i], ww[i], sum);\n if (lane == 0) sum = fmaf(tnorm, ww[dim], sum);\n sum = warp_sum(sum);\n if (lane == 0) h1[o] = silu_f(sum + b1[o]);\n }\n __syncthreads();\n for (int o=warp; o<hidden; o+=nwarps) {\n float sum = 0.0f; const float* ww = w2 + o*hidden;\n for (int i=lane; i<hidden; i+=32) sum = fmaf(h1[i], ww[i], sum);\n sum = warp_sum(sum);\n if (lane == 0) h2[o] = silu_f(sum + b2[o]);\n }\n __syncthreads();\n for (int o=warp; o<dim; o+=nwarps) {\n float sum = 0.0f; const float* ww = w3 + o*hidden;\n for (int i=lane; i<hidden; i+=32) sum = fmaf(h2[i], ww[i], sum);\n sum = warp_sum(sum);\n if (lane == 0) eps[o] = sum + b3[o];\n }\n __syncthreads();\n float c1 = sqrt_recip_alphas[t_idx], c2 = coef_eps[t_idx];\n for (int i=tid; i<dim; i+=blockDim.x) x[i] = c1 * (x[i] - c2 * eps[i]);\n __syncthreads();\n }\n for (int i=tid; i<dim; i+=blockDim.x) out[b*dim+i] = x[i];\n}\n\nvoid fused_ddpm_launch(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, torch::Tensor out, int num_steps, int total_timesteps) {\n int B=x.size(0), dim=x.size(1); int threads=1024; size_t shmem=(dim+4*dim+4*dim+dim)*sizeof(float);\n bool is_float=(t_start.scalar_type()==at::kFloat); const int64_t* p_l=is_float?nullptr:t_start.data_ptr<int64_t>(); const float* p_f=is_float?t_start.data_ptr<float>():nullptr;\n fused_ddpm_kernel<<<B, threads, shmem, at::cuda::getCurrentCUDAStream()>>>(x.data_ptr<float>(), p_l, p_f, is_float?1:0, 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>(), out.data_ptr<float>(), dim, num_steps, total_timesteps);\n}\n'''\n_ext = load_inline(name='fused_ddpm_ext_v4c', cpp_sources=_cpp_src, cuda_sources=_cuda_src, functions=['fused_ddpm_forward'], verbose=False, extra_cflags=['-O3'], extra_cuda_cflags=['-O3'])\n\ndef _linear_beta_schedule(num_timesteps:int,beta_start:float,beta_end:float)->torch.Tensor:\n return torch.linspace(beta_start,beta_end,num_timesteps)\nclass ModelNew(nn.Module):\n def __init__(self, dim:int, num_fused_steps:int, total_timesteps:int=1000):\n super().__init__(); self.dim=dim; self.num_fused_steps=num_fused_steps; self.total_timesteps=total_timesteps\n 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))\n betas=_linear_beta_schedule(total_timesteps,1e-4,0.02); alphas=1.0-betas; ac=torch.cumprod(alphas,dim=0)\n self.register_buffer('betas',betas); self.register_buffer('alphas',alphas); self.register_buffer('alphas_cumprod',ac)\n self.register_buffer('sqrt_recip_alphas',torch.sqrt(1.0/alphas)); self.register_buffer('coef_eps',betas/torch.sqrt(1.0-ac))\n def forward(self,x,t_start):\n return _ext.fused_ddpm_forward(x.contiguous(),t_start.contiguous(),self.eps_net[0].weight,self.eps_net[0].bias,self.eps_net[2].weight,self.eps_net[2].bias,self.eps_net[4].weight,self.eps_net[4].bias,self.sqrt_recip_alphas,self.coef_eps,self.num_fused_steps,self.total_timesteps)\n"
}