call → compile_kernel
callcompile_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nimport torch.nn.functional as F\nfrom torch.utils.cpp_extension import load_inline\n\n_cpp_src = r'''\n#include <torch/extension.h>\ntorch::Tensor channel_shuffle_cuda(torch::Tensor input, int64_t groups);\n'''\n\n_cuda_src = r'''\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\ntemplate<int C>\n__global__ void channel_shuffle_g3_kernel(const float* __restrict__ x, float* __restrict__ y, long pixels) {\n long p = blockIdx.x * (long)blockDim.x + threadIdx.x;\n long stride = (long)blockDim.x * gridDim.x;\n constexpr int cpg = C / 3;\n for (; p < pixels; p += stride) {\n long base = p * (long)C;\n #pragma unroll\n for (int c = 0; c < C; ++c) {\n int ci = (c % 3) * cpg + (c / 3);\n y[base + c] = x[base + ci];\n }\n }\n}\n\ntemplate <typename scalar_t>\n__global__ void channel_shuffle_generic_kernel(const scalar_t* __restrict__ x, scalar_t* __restrict__ y, long total, int C, int groups) {\n long idx = blockIdx.x * (long)blockDim.x + threadIdx.x;\n long stride = (long)blockDim.x * gridDim.x;\n int cpg = C / groups;\n for (; idx < total; idx += stride) {\n int c = idx % C;\n int ci = (c % groups) * cpg + (c / groups);\n y[idx] = x[(idx - c) + ci];\n }\n}\n\ntorch::Tensor channel_shuffle_cuda(torch::Tensor input, int64_t groups64) {\n TORCH_CHECK(input.is_cuda(), \"input must be CUDA\");\n TORCH_CHECK(input.dim() == 4, \"input must be NCHW\");\n int C = (int)input.size(1);\n int groups = (int)groups64;\n auto x = input.contiguous(at::MemoryFormat::ChannelsLast);\n auto y = torch::empty_like(x, x.options(), at::MemoryFormat::ChannelsLast);\n long pixels = (long)input.size(0) * input.size(2) * input.size(3);\n if (x.scalar_type() == at::kFloat && groups == 3 && (C == 60 || C == 120 || C == 240)) {\n int threads = 128;\n int blocks = (int)((pixels + threads - 1) / threads);\n if (blocks > 65535) blocks = 65535;\n if (C == 60) channel_shuffle_g3_kernel<60><<<blocks, threads>>>(x.data_ptr<float>(), y.data_ptr<float>(), pixels);\n else if (C == 120) channel_shuffle_g3_kernel<120><<<blocks, threads>>>(x.data_ptr<float>(), y.data_ptr<float>(), pixels);\n else channel_shuffle_g3_kernel<240><<<blocks, threads>>>(x.data_ptr<float>(), y.data_ptr<float>(), pixels);\n } else {\n long total = pixels * (long)C;\n int threads = 256;\n int blocks = (int)((total + threads - 1) / threads);\n if (blocks > 65535) blocks = 65535;\n AT_DISPATCH_FLOATING_TYPES_AND_HALF(input.scalar_type(), \"channel_shuffle_cuda\", [&] {\n channel_shuffle_generic_kernel<scalar_t><<<blocks, threads>>>(x.data_ptr<scalar_t>(), y.data_ptr<scalar_t>(), total, C, groups);\n });\n }\n return y;\n}\n'''\n\n_ext = load_inline(name='shufflenet_kernels_final_a', cpp_sources=[_cpp_src], cuda_sources=[_cuda_src], functions=['channel_shuffle_cuda'], verbose=False, extra_cuda_cflags=['-O3'], extra_cflags=['-O3'])\n\ndef _fuse_conv_bn(conv, bn):\n scale = bn.weight / torch.sqrt(bn.running_var + bn.eps)\n w = (conv.weight * scale.reshape(-1, 1, 1, 1)).contiguous(memory_format=torch.channels_last)\n b = (bn.bias - bn.running_mean * scale).contiguous()\n return w, b\n\ndef _conv_bn(x, conv, bn, cache, key, relu=False):\n if conv.training or bn.training:\n y = bn(conv(x))\n else:\n item = cache.get(key)\n if item is None or item[0].device != conv.weight.device or item[0].dtype != conv.weight.dtype:\n with torch.no_grad():\n item = _fuse_conv_bn(conv, bn)\n cache[key] = item\n y = F.conv2d(x, item[0], item[1], conv.stride, conv.padding, conv.dilation, conv.groups)\n if relu:\n return torch.relu_(y)\n return y\n\nclass ChannelShuffleNew(nn.Module):\n def __init__(self, groups):\n super().__init__()\n self.groups = groups\n def forward(self, x):\n return _ext.channel_shuffle_cuda(x, self.groups)\n\nclass ShuffleNetUnitNew(nn.Module):\n def __init__(self, in_channels, out_channels, groups=3):\n super().__init__()\n assert out_channels % 4 == 0\n mid_channels = out_channels // 4\n self.conv1 = nn.Conv2d(in_channels, mid_channels, kernel_size=1, stride=1, padding=0, groups=groups, bias=False)\n self.bn1 = nn.BatchNorm2d(mid_channels)\n self.conv2 = nn.Conv2d(mid_channels, mid_channels, kernel_size=3, stride=1, padding=1, groups=mid_channels, bias=False)\n self.bn2 = nn.BatchNorm2d(mid_channels)\n self.conv3 = nn.Conv2d(mid_channels, out_channels, kernel_size=1, stride=1, padding=0, groups=groups, bias=False)\n self.bn3 = nn.BatchNorm2d(out_channels)\n self.shuffle = ChannelShuffleNew(groups)\n if in_channels == out_channels:\n self.shortcut = nn.Sequential()\n else:\n self.shortcut = nn.Sequential(\n nn.Conv2d(in_channels, out_channels, kernel_size=1, stride=1, padding=0, bias=False),\n nn.BatchNorm2d(out_channels)\n )\n self._cache = {}\n def forward(self, x):\n out = _conv_bn(x, self.conv1, self.bn1, self._cache, '1', True)\n out = _conv_bn(out, self.conv2, self.bn2, self._cache, '2', False)\n out = self.shuffle(out)\n out = _conv_bn(out, self.conv3, self.bn3, self._cache, '3', True)\n if len(self.shortcut) == 0:\n return out.add_(x)\n sc = _conv_bn(x, self.shortcut[0], self.shortcut[1], self._cache, 's', False)\n return out.add_(sc)\n\nclass ModelNew(nn.Module):\n def __init__(self, num_classes=1000, groups=3, stages_repeats=[3, 7, 3], stages_out_channels=[24, 240, 480, 960]):\n super().__init__()\n self.conv1 = nn.Conv2d(3, stages_out_channels[0], kernel_size=3, stride=2, padding=1, bias=False)\n self.bn1 = nn.BatchNorm2d(stages_out_channels[0])\n self.maxpool = nn.MaxPool2d(kernel_size=3, stride=2, padding=1)\n self.stage2 = self._make_stage(stages_out_channels[0], stages_out_channels[1], stages_repeats[0], groups)\n self.stage3 = self._make_stage(stages_out_channels[1], stages_out_channels[2], stages_repeats[1], groups)\n self.stage4 = self._make_stage(stages_out_channels[2], stages_out_channels[3], stages_repeats[2], groups)\n self.conv5 = nn.Conv2d(stages_out_channels[3], 1024, kernel_size=1, stride=1, padding=0, bias=False)\n self.bn5 = nn.BatchNorm2d(1024)\n self.fc = nn.Linear(1024, num_classes)\n self._cache = {}\n self.to(memory_format=torch.channels_last)\n def _make_stage(self, in_channels, out_channels, repeats, groups):\n layers = [ShuffleNetUnitNew(in_channels, out_channels, groups)]\n for _ in range(1, repeats):\n layers.append(ShuffleNetUnitNew(out_channels, out_channels, groups))\n return nn.Sequential(*layers)\n def _clear_cache(self):\n self._cache.clear()\n for m in self.modules():\n if isinstance(m, ShuffleNetUnitNew):\n m._cache.clear()\n def load_state_dict(self, state_dict, strict=True, assign=False):\n res = super().load_state_dict(state_dict, strict=strict, assign=assign)\n self.to(memory_format=torch.channels_last)\n self._clear_cache()\n return res\n def train(self, mode=True):\n res = super().train(mode)\n self._clear_cache()\n return res\n def forward(self, x):\n x = x.contiguous(memory_format=torch.channels_last)\n x = _conv_bn(x, self.conv1, self.bn1, self._cache, 'stem', True)\n x = self.maxpool(x)\n x = self.stage2(x)\n x = self.stage3(x)\n x = self.stage4(x)\n x = _conv_bn(x, self.conv5, self.bn5, self._cache, 'conv5', True)\n x = F.adaptive_avg_pool2d(x, (1, 1))\n x = torch.flatten(x, 1)\n return self.fc(x)\n"
}