call → run_correctness
callrun_correctness
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n_cpp_src = r'''\n#include <torch/extension.h>\ntorch::Tensor avgpool2d_forward(torch::Tensor x, int64_t kernel_size, int64_t stride, int64_t padding);\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#define CHECK_CUDA(x) TORCH_CHECK(x.is_cuda(), #x \" must be a CUDA tensor\")\n\ntemplate<int WARPS_PER_BLOCK>\n__global__ __launch_bounds__(WARPS_PER_BLOCK*32, 4)\nvoid avgpool_k11_triplewarp_kernel(const float* __restrict__ x, float* __restrict__ y,\n int total_groups, int OH, int OW, int H, int W) {\n int warp_in_block = threadIdx.x >> 5;\n int lane = threadIdx.x & 31;\n int group_idx = (int)blockIdx.x * WARPS_PER_BLOCK + warp_in_block;\n if (group_idx >= total_groups) return;\n int groups_per_row = (OW + 2) / 3;\n int group_ow = group_idx % groups_per_row;\n int t = group_idx / groups_per_row;\n int oh = t % OH;\n int nc = t / OH;\n int ow0 = group_ow * 3;\n const float* base = x + ((long long)nc * H + oh * 11) * W + ow0 * 11;\n float s = 0.0f;\n if (ow0 + 2 < OW) {\n#pragma unroll\n for (int r = 0; r < 11; ++r) {\n const float* row = base + r * W;\n s += row[lane];\n if (lane == 31) s += row[32];\n }\n } else {\n // generic tail for non-multiple of 3 output widths\n int maxcols = (OW - ow0) * 11;\n if (lane < maxcols) {\n#pragma unroll\n for (int r = 0; r < 11; ++r) s += base[r * W + lane];\n }\n if (lane == 31 && maxcols > 32) {\n#pragma unroll\n for (int r = 0; r < 11; ++r) s += base[r * W + 32];\n }\n }\n long long outbase = ((long long)nc * OH + oh) * OW + ow0;\n if (lane < 11) {\n unsigned mask = 0x000007ffu;\n float v;\n v = __shfl_down_sync(mask, s, 8); if (lane + 8 < 11) s += v;\n v = __shfl_down_sync(mask, s, 4); if (lane + 4 < 11) s += v;\n v = __shfl_down_sync(mask, s, 2); if (lane + 2 < 11) s += v;\n v = __shfl_down_sync(mask, s, 1); if (lane + 1 < 11) s += v;\n if (lane == 0) y[outbase] = s * 0.0082644628099173553719f;\n } else if (lane < 22 && ow0 + 1 < OW) {\n unsigned mask = 0x003ff800u;\n float v;\n v = __shfl_down_sync(mask, s, 8); if (lane + 8 < 22) s += v;\n v = __shfl_down_sync(mask, s, 4); if (lane + 4 < 22) s += v;\n v = __shfl_down_sync(mask, s, 2); if (lane + 2 < 22) s += v;\n v = __shfl_down_sync(mask, s, 1); if (lane + 1 < 22) s += v;\n if (lane == 11) y[outbase + 1] = s * 0.0082644628099173553719f;\n } else if (ow0 + 2 < OW) {\n unsigned mask = 0xffc00000u;\n float v;\n v = __shfl_down_sync(mask, s, 8); if (lane + 8 < 32) s += v;\n v = __shfl_down_sync(mask, s, 4); if (lane + 4 < 32) s += v;\n v = __shfl_down_sync(mask, s, 2); if (lane + 2 < 32) s += v;\n v = __shfl_down_sync(mask, s, 1); if (lane + 1 < 32) s += v;\n if (lane == 22) y[outbase + 2] = s * 0.0082644628099173553719f;\n }\n}\n\n__global__ void avgpool_generic_kernel(const float* __restrict__ x, float* __restrict__ y,\n int total_out, int C, int H, int W, int OH, int OW,\n int K, int S, int P) {\n int idx = blockIdx.x * blockDim.x + threadIdx.x;\n if (idx >= total_out) return;\n int ow = idx % OW; int tmp = idx / OW; int oh = tmp % OH; tmp /= OH; int c = tmp % C; int n = tmp / C;\n int ih0 = oh * S - P; int iw0 = ow * S - P; float sum=0.0f;\n for(int kh=0; kh<K; ++kh){ int ih=ih0+kh; if((unsigned)ih<(unsigned)H){ const float* row=x+(((long long)n*C+c)*H+ih)*W; for(int kw=0; kw<K; ++kw){ int iw=iw0+kw; if((unsigned)iw<(unsigned)W) sum += row[iw]; }}}\n y[idx] = sum / (float)(K*K);\n}\n\ntorch::Tensor avgpool2d_forward(torch::Tensor x, int64_t kernel_size, int64_t stride, int64_t padding) {\n CHECK_CUDA(x); TORCH_CHECK(x.scalar_type()==at::kFloat, \"only float32 supported\"); TORCH_CHECK(x.dim()==4, \"input must be NCHW\");\n auto xc=x.contiguous(); int N=xc.size(0), C=xc.size(1), H=xc.size(2), W=xc.size(3); int K=kernel_size, S=stride, P=padding;\n int OH=(H+2*P-K)/S+1, OW=(W+2*P-K)/S+1; auto y=torch::empty({N,C,OH,OW}, xc.options()); cudaStream_t stream=at::cuda::getCurrentCUDAStream(); int total=N*C*OH*OW;\n if(K==11 && S==11 && P==0){ constexpr int WARPS=8; int groups=N*C*OH*((OW+2)/3); int blocks=(groups+WARPS-1)/WARPS; avgpool_k11_triplewarp_kernel<WARPS><<<blocks,WARPS*32,0,stream>>>(xc.data_ptr<float>(), y.data_ptr<float>(), groups, OH, OW, H, W); }\n else { int threads=256, blocks=(total+threads-1)/threads; avgpool_generic_kernel<<<blocks,threads,0,stream>>>(xc.data_ptr<float>(), y.data_ptr<float>(), total, C,H,W,OH,OW,K,S,P); }\n return y;\n}\n'''\n_ext=load_inline(name='avgpool2d_ext_tri_v1', cpp_sources=_cpp_src, cuda_sources=_cuda_src, functions=['avgpool2d_forward'], extra_cuda_cflags=['-O3','--use_fast_math'], extra_cflags=['-O3'], verbose=False)\nclass ModelNew(nn.Module):\n def __init__(self, kernel_size:int, stride:int=None, padding:int=0):\n super(ModelNew,self).__init__(); self.kernel_size=int(kernel_size); self.stride=int(kernel_size if stride is None else stride); self.padding=int(padding)\n def forward(self,x:torch.Tensor)->torch.Tensor:\n return _ext.avgpool2d_forward(x,self.kernel_size,self.stride,self.padding)\n"
}