call → compile_kernel
callcompile_kernel
{
"kernel_code": "import math\nimport torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\n_cpp_src = r\"\"\"\n#include <torch/extension.h>\ntorch::Tensor netvlad_forward(torch::Tensor x, torch::Tensor clusters, torch::Tensor bn_weight, torch::Tensor bn_bias,\n torch::Tensor running_mean, torch::Tensor running_var, torch::Tensor clusters2,\n bool training, double eps, double momentum, int64_t cluster_size, int64_t ghost_clusters);\n\"\"\"\n\n_cuda_src = r\"\"\"\n#include <torch/extension.h>\n#include <ATen/ATen.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <cmath>\n\n#define CHECK_CUDA(x) TORCH_CHECK(x.is_cuda(), #x \" must be a CUDA tensor\")\n\n__global__ void mean_var_kernel(const float* __restrict__ logits, float* __restrict__ mean, float* __restrict__ invstd,\n float* __restrict__ running_mean, float* __restrict__ running_var,\n int M, int C, float eps, float momentum) {\n int c = blockIdx.x, tid = threadIdx.x;\n extern __shared__ float smem[];\n float* ssum = smem; float* ssq = smem + blockDim.x;\n float sum = 0.0f, sq = 0.0f;\n for (int i = tid; i < M; i += blockDim.x) {\n float v = logits[i * C + c]; sum += v; sq = fmaf(v, v, sq);\n }\n ssum[tid] = sum; ssq[tid] = sq; __syncthreads();\n for (int stride = blockDim.x >> 1; stride > 0; stride >>= 1) { if (tid < stride) { ssum[tid] += ssum[tid+stride]; ssq[tid] += ssq[tid+stride]; } __syncthreads(); }\n if (tid == 0) {\n float m = ssum[0] / (float)M;\n float var = fmaxf(ssq[0] / (float)M - m * m, 0.0f);\n mean[c] = m; invstd[c] = rsqrtf(var + eps);\n if (running_mean != nullptr && running_var != nullptr) {\n float unbias = var * ((float)M / (float)(M - 1));\n running_mean[c] = running_mean[c] * (1.0f - momentum) + momentum * m;\n running_var[c] = running_var[c] * (1.0f - momentum) + momentum * unbias;\n }\n }\n}\n\n__global__ void eval_stats_kernel(const float* __restrict__ running_mean, const float* __restrict__ running_var,\n float* __restrict__ mean, float* __restrict__ invstd, int C, float eps) {\n int c = blockIdx.x * blockDim.x + threadIdx.x;\n if (c < C) { mean[c] = running_mean[c]; invstd[c] = rsqrtf(running_var[c] + eps); }\n}\n\n__global__ void softmax_assign_kernel(const float* __restrict__ logits,\n const float* __restrict__ mean, const float* __restrict__ invstd,\n const float* __restrict__ gamma, const float* __restrict__ beta,\n float* __restrict__ assign, float* __restrict__ asum,\n int M, int N, int K, int C) {\n int row = blockIdx.x, tid = threadIdx.x;\n extern __shared__ float sh[];\n float v = -INFINITY;\n if (tid < C) { v = (logits[row * C + tid] - mean[tid]) * invstd[tid]; v = fmaf(v, gamma[tid], beta[tid]); }\n sh[tid] = v; __syncthreads();\n for (int stride = blockDim.x >> 1; stride > 0; stride >>= 1) { if (tid < stride) sh[tid] = fmaxf(sh[tid], sh[tid+stride]); __syncthreads(); }\n float mx = sh[0];\n float e = (tid < C) ? expf(v - mx) : 0.0f;\n sh[tid] = e; __syncthreads();\n for (int stride = blockDim.x >> 1; stride > 0; stride >>= 1) { if (tid < stride) sh[tid] += sh[tid+stride]; __syncthreads(); }\n if (tid < K) {\n float p = e / sh[0]; int b = row / N; int n = row - b * N;\n assign[(b * K + tid) * N + n] = p;\n atomicAdd(&asum[b * K + tid], p);\n }\n}\n\n__global__ void finalize_kernel(const float* __restrict__ vlad, const float* __restrict__ asum,\n const float* __restrict__ clusters2, float* __restrict__ out,\n int B, int K, int D) {\n int bk = blockIdx.x, b = bk / K, k = bk - b * K, tid = threadIdx.x;\n extern __shared__ float sm[];\n float local = 0.0f, a = asum[b * K + k];\n for (int d = tid; d < D; d += blockDim.x) { float r = vlad[(b*K+k)*D+d] - a * clusters2[d*K+k]; local = fmaf(r, r, local); }\n sm[tid] = local; __syncthreads();\n for (int stride = blockDim.x >> 1; stride > 0; stride >>= 1) { if (tid < stride) sm[tid] += sm[tid+stride]; __syncthreads(); }\n float scale = (sm[0] > 1.0e-24f ? rsqrtf(sm[0]) : 1.0e12f) * rsqrtf((float)K);\n for (int d = tid; d < D; d += blockDim.x) { float r = vlad[(b*K+k)*D+d] - a * clusters2[d*K+k]; out[b*(D*K)+d*K+k] = r * scale; }\n}\n\ntorch::Tensor netvlad_forward(torch::Tensor x, torch::Tensor clusters, torch::Tensor bn_weight, torch::Tensor bn_bias,\n torch::Tensor running_mean, torch::Tensor running_var, torch::Tensor clusters2,\n bool training, double eps, double momentum, int64_t cluster_size, int64_t ghost_clusters) {\n CHECK_CUDA(x); CHECK_CUDA(clusters); CHECK_CUDA(clusters2);\n auto xc = x.contiguous(); auto cc = clusters.contiguous(); auto g = bn_weight.contiguous(); auto be = bn_bias.contiguous(); auto c2 = clusters2.contiguous();\n int B=(int)xc.size(0), N=(int)xc.size(1), D=(int)xc.size(2), K=(int)cluster_size, C=K+(int)ghost_clusters, M=B*N;\n auto logits = at::mm(xc.view({M,D}), cc);\n auto opts = xc.options(); auto mean = torch::empty({C}, opts); auto invstd = torch::empty({C}, opts);\n cudaStream_t stream = at::cuda::getCurrentCUDAStream();\n if (training) mean_var_kernel<<<C,256,2*256*sizeof(float),stream>>>(logits.data_ptr<float>(), mean.data_ptr<float>(), invstd.data_ptr<float>(), running_mean.data_ptr<float>(), running_var.data_ptr<float>(), M, C, (float)eps, (float)momentum);\n else eval_stats_kernel<<<(C+127)/128,128,0,stream>>>(running_mean.data_ptr<float>(), running_var.data_ptr<float>(), mean.data_ptr<float>(), invstd.data_ptr<float>(), C, (float)eps);\n auto assign = torch::empty({B,K,N}, opts); auto asum = torch::zeros({B,K}, opts);\n softmax_assign_kernel<<<M,64,64*sizeof(float),stream>>>(logits.data_ptr<float>(), mean.data_ptr<float>(), invstd.data_ptr<float>(), g.data_ptr<float>(), be.data_ptr<float>(), assign.data_ptr<float>(), asum.data_ptr<float>(), M, N, K, C);\n auto vlad = at::bmm(assign, xc);\n auto out = torch::empty({B,D*K}, opts);\n finalize_kernel<<<B*K,256,256*sizeof(float),stream>>>(vlad.data_ptr<float>(), asum.data_ptr<float>(), c2.data_ptr<float>(), out.data_ptr<float>(), B, K, D);\n return out;\n}\n\"\"\"\n\n_ext = load_inline(name=\"netvlad_ext_v2\", cpp_sources=_cpp_src, cuda_sources=_cuda_src,\n functions=[\"netvlad_forward\"], extra_cuda_cflags=[\"-O3\", \"--use_fast_math\"], verbose=False)\n\nclass ModelNew(nn.Module):\n def __init__(self, cluster_size, feature_size, ghost_clusters):\n super(ModelNew, self).__init__()\n self.feature_size = feature_size; self.cluster_size = cluster_size; self.ghost_clusters = ghost_clusters\n init_sc = (1 / math.sqrt(feature_size)); clusters = cluster_size + ghost_clusters\n self.clusters = nn.Parameter(init_sc * torch.randn(feature_size, clusters))\n self.batch_norm = nn.BatchNorm1d(clusters)\n self.clusters2 = nn.Parameter(init_sc * torch.randn(1, feature_size, cluster_size))\n self.out_dim = self.cluster_size * feature_size\n def forward(self, x, mask=None):\n if self.batch_norm.training and self.batch_norm.track_running_stats:\n self.batch_norm.num_batches_tracked.add_(1)\n return _ext.netvlad_forward(x, self.clusters, self.batch_norm.weight, self.batch_norm.bias,\n self.batch_norm.running_mean, self.batch_norm.running_var, self.clusters2,\n self.batch_norm.training, self.batch_norm.eps, self.batch_norm.momentum,\n self.cluster_size, self.ghost_clusters)\n"
}