call → compile_kernel
callcompile_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncpp_src = r'''\n#include <torch/extension.h>\n\ntorch::Tensor convt3d_forward(torch::Tensor x, torch::Tensor w, int64_t stride, int64_t padding, int64_t groups);\n'''\n\ncuda_src = r'''\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\n__global__ void convt3d_s2k3_kernel(const float* __restrict__ x,\n const float* __restrict__ w,\n float* __restrict__ y,\n long N, long Cin, long D, long H, long W,\n long Cout, long G,\n long Do, long Ho, long Wo) {\n long idx = blockIdx.x * blockDim.x + threadIdx.x;\n long total = N * Cout * Do * Ho * Wo;\n if (idx >= total) return;\n\n long t = idx;\n long ow = t % Wo; t /= Wo;\n long oh = t % Ho; t /= Ho;\n long od = t % Do; t /= Do;\n long oc = t % Cout; t /= Cout;\n long n = t;\n\n const long out_per_g = Cout / G;\n const long in_per_g = Cin / G;\n const long g = oc / out_per_g;\n const long ocg = oc - g * out_per_g;\n const long ic_start = g * in_per_g;\n\n float acc = 0.0f;\n\n // stride=2, padding=1, kernel=3 specialized. For each dim, only k with (o+1-k) even.\n int kd0 = ((od + 1) & 1) ? 0 : 1; // parity matching first kd\n #pragma unroll\n for (int kd = kd0; kd < 3; kd += 2) {\n long id = (od + 1 - kd) >> 1;\n if ((unsigned long)id >= (unsigned long)D) continue;\n int kh0 = ((oh + 1) & 1) ? 0 : 1;\n #pragma unroll\n for (int kh = kh0; kh < 3; kh += 2) {\n long ih = (oh + 1 - kh) >> 1;\n if ((unsigned long)ih >= (unsigned long)H) continue;\n int kw0 = ((ow + 1) & 1) ? 0 : 1;\n #pragma unroll\n for (int kw = kw0; kw < 3; kw += 2) {\n long iw = (ow + 1 - kw) >> 1;\n if ((unsigned long)iw >= (unsigned long)W) continue;\n long x_base = ((n * Cin + ic_start) * D + id) * H * W + ih * W + iw;\n long w_base = (((ic_start) * out_per_g + ocg) * 3 + kd) * 3 * 3 + kh * 3 + kw;\n #pragma unroll 1\n for (long ii = 0; ii < in_per_g; ++ii) {\n acc += x[x_base + ii * D * H * W] * w[w_base + ii * out_per_g * 27];\n }\n }\n }\n }\n y[idx] = acc;\n}\n\n__global__ void convt3d_generic_kernel(const float* __restrict__ x,\n const float* __restrict__ w,\n float* __restrict__ y,\n long N, long Cin, long D, long H, long W,\n long Cout, long K, long stride, long padding, long G,\n long Do, long Ho, long Wo) {\n long idx = blockIdx.x * blockDim.x + threadIdx.x;\n long total = N * Cout * Do * Ho * Wo;\n if (idx >= total) return;\n long t=idx; long ow=t%Wo; t/=Wo; long oh=t%Ho; t/=Ho; long od=t%Do; t/=Do; long oc=t%Cout; t/=Cout; long n=t;\n long out_per_g=Cout/G, in_per_g=Cin/G, g=oc/out_per_g, ocg=oc-g*out_per_g, ic_start=g*in_per_g;\n float acc=0.0f;\n for(long kd=0; kd<K; ++kd){ long zd=od+padding-kd; if(zd%stride) continue; long id=zd/stride; if((unsigned long)id>=(unsigned long)D) continue;\n for(long kh=0; kh<K; ++kh){ long zh=oh+padding-kh; if(zh%stride) continue; long ih=zh/stride; if((unsigned long)ih>=(unsigned long)H) continue;\n for(long kw=0; kw<K; ++kw){ long zw=ow+padding-kw; if(zw%stride) continue; long iw=zw/stride; if((unsigned long)iw>=(unsigned long)W) continue;\n long xb=((n*Cin+ic_start)*D+id)*H*W+ih*W+iw;\n long wb=(((ic_start)*out_per_g+ocg)*K+kd)*K*K+kh*K+kw;\n for(long ii=0; ii<in_per_g; ++ii) acc += x[xb+ii*D*H*W]*w[wb+ii*out_per_g*K*K*K];\n }\n }\n }\n y[idx]=acc;\n}\n\ntorch::Tensor convt3d_forward(torch::Tensor x, torch::Tensor w, int64_t stride, int64_t padding, int64_t groups) {\n auto N = x.size(0), Cin = x.size(1), D = x.size(2), H = x.size(3), W = x.size(4);\n auto K = w.size(2);\n auto Cout = w.size(1) * groups;\n auto Do = (D - 1) * stride - 2 * padding + K;\n auto Ho = (H - 1) * stride - 2 * padding + K;\n auto Wo = (W - 1) * stride - 2 * padding + K;\n auto y = torch::empty({N, Cout, Do, Ho, Wo}, x.options());\n long total = N*Cout*Do*Ho*Wo;\n int threads = 256;\n int blocks = (total + threads - 1) / threads;\n if (x.scalar_type() == torch::kFloat32 && K == 3 && stride == 2 && padding == 1) {\n convt3d_s2k3_kernel<<<blocks, threads>>>(x.data_ptr<float>(), w.data_ptr<float>(), y.data_ptr<float>(), N,Cin,D,H,W,Cout,groups,Do,Ho,Wo);\n } else {\n convt3d_generic_kernel<<<blocks, threads>>>(x.data_ptr<float>(), w.data_ptr<float>(), y.data_ptr<float>(), N,Cin,D,H,W,Cout,K,stride,padding,groups,Do,Ho,Wo);\n }\n return y;\n}\n'''\n\n_ext = load_inline(name='convt3d_ext_v1', cpp_sources=cpp_src, cuda_sources=cuda_src, functions=['convt3d_forward'], extra_cuda_cflags=['-O3', '--use_fast_math'], verbose=False)\n\nclass ModelNew(nn.Module):\n def __init__(self, in_channels: int, out_channels: int, kernel_size: int, stride: int = 1, padding: int = 0, output_padding: int = 0, groups: int = 1, bias: bool = False):\n super(ModelNew, self).__init__()\n self.conv_transpose3d = nn.ConvTranspose3d(in_channels, out_channels, kernel_size=(kernel_size, kernel_size, kernel_size), stride=stride, padding=padding, groups=groups, bias=bias)\n self.stride = stride\n self.padding = padding\n self.groups = groups\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n return _ext.convt3d_forward(x.contiguous(), self.conv_transpose3d.weight.contiguous(), int(self.stride), int(self.padding), int(self.groups))\n"
}