import torch
import torch.nn as nn


class Model(nn.Module):
    """
    Voxel-grid pooling for 3-D molecular point clouds.  Discretizes
    continuous 3-D coordinates into a regular grid and pools per-atom
    features into each voxel.  Used in 3-D-CNN approaches for molecular
    property prediction and ligand docking (e.g., PointVoxel-CNN,
    3D-Infomax).
    """

    def __init__(self, feat_dim, grid_size=32, voxel_size=1.0):
        super().__init__()
        self.feat_dim = feat_dim
        self.grid_size = grid_size
        self.voxel_size = voxel_size
        self.linear = nn.Linear(feat_dim, feat_dim)

    def forward(
        self, coords: torch.Tensor, features: torch.Tensor
    ) -> torch.Tensor:
        """
        Args:
            coords:   (B, N, 3) – atom coordinates
            features: (B, N, feat_dim) – per-atom features
        Returns:
            voxel_grid: (B, feat_dim, G, G, G) – pooled voxel features
        """
        B, N, _ = coords.shape
        G = self.grid_size

        features = self.linear(features)

        # Center coordinates and discretize
        center = coords.mean(dim=1, keepdim=True)
        shifted = (coords - center) / self.voxel_size + G / 2.0
        indices = shifted.long().clamp(0, G - 1)  # (B, N, 3)

        voxel_grid = coords.new_zeros(B, self.feat_dim, G, G, G)
        ix = indices[:, :, 0]
        iy = indices[:, :, 1]
        iz = indices[:, :, 2]

        batch_idx = torch.arange(B, device=coords.device).unsqueeze(1).expand(-1, N)
        for d in range(self.feat_dim):
            feat_d = features[:, :, d]
            voxel_grid[:, d].index_put_(
                (batch_idx.reshape(-1), ix.reshape(-1), iy.reshape(-1), iz.reshape(-1)),
                feat_d.reshape(-1),
                accumulate=True,
            )

        return voxel_grid


feat_dim = 32
grid_size = 32
voxel_size = 1.0
num_atoms = 256
batch_size = 4


def get_inputs():
    p = popcorn_pri
    bs = p.jitter_int(batch_size)
    na = p.jitter_int(num_atoms)
    coords = torch.randn(bs, na, 3) * 5.0
    features = torch.randn(bs, na, feat_dim)
    return [coords, features]


def get_init_inputs():
    return [feat_dim, grid_size, voxel_size]

import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline

cpp_src = r'''
#include <torch/extension.h>
torch::Tensor voxel_forward(torch::Tensor coords, torch::Tensor features, torch::Tensor weight, torch::Tensor bias, int64_t grid_size, double voxel_size);
'''

cuda_src = r'''
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda.h>
#include <cuda_runtime.h>

__global__ void center_kernel(const float* __restrict__ coords, float* __restrict__ centers, int B, int N) {
    int b = blockIdx.x, tid = threadIdx.x;
    float sx = 0.f, sy = 0.f, sz = 0.f;
    for (int i = tid; i < N; i += blockDim.x) {
        const float* p = coords + ((b * N + i) * 3);
        sx += p[0]; sy += p[1]; sz += p[2];
    }
    __shared__ float sh[256*3];
    sh[tid] = sx; sh[256+tid] = sy; sh[512+tid] = sz;
    __syncthreads();
    for (int s = blockDim.x >> 1; s > 0; s >>= 1) {
        if (tid < s) { sh[tid] += sh[tid+s]; sh[256+tid] += sh[256+tid+s]; sh[512+tid] += sh[512+tid+s]; }
        __syncthreads();
    }
    if (tid == 0) {
        float invN = 1.0f / (float)N;
        centers[b*3] = sh[0] * invN;
        centers[b*3+1] = sh[256] * invN;
        centers[b*3+2] = sh[512] * invN;
    }
}

__global__ void scatter_linear32_kernel(
    const float* __restrict__ coords,
    const float* __restrict__ features,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    const float* __restrict__ centers,
    float* __restrict__ out,
    int N, int G, float inv_voxel) {
    int atom = blockIdx.x;
    int b = atom / N;
    int n = atom - b * N;
    int d = threadIdx.x;
    __shared__ float sf[32];
    const float* fin = features + ((b * N + n) * 32);
    sf[d] = fin[d];
    __syncthreads();

    const float* c = coords + ((b * N + n) * 3);
    float halfG = (float)G * 0.5f;
    int ix = (int)((c[0] - centers[b*3]) * inv_voxel + halfG);
    int iy = (int)((c[1] - centers[b*3+1]) * inv_voxel + halfG);
    int iz = (int)((c[2] - centers[b*3+2]) * inv_voxel + halfG);
    ix = max(0, min(ix, G-1)); iy = max(0, min(iy, G-1)); iz = max(0, min(iz, G-1));
    int vox = (ix * G + iy) * G + iz;

    const float* w = weight + d * 32;
    float acc = bias[d];
    #pragma unroll
    for (int k = 0; k < 32; ++k) acc = fmaf(sf[k], w[k], acc);
    size_t off = (((size_t)b * 32 + d) * (size_t)G * G * G) + vox;
    atomicAdd(out + off, acc);
}

__global__ void scatter_linear_kernel(
    const float* __restrict__ coords, const float* __restrict__ features,
    const float* __restrict__ weight, const float* __restrict__ bias,
    const float* __restrict__ centers, float* __restrict__ out,
    int B, int N, int F, int G, float inv_voxel) {
    int atom = blockIdx.x; int b = atom / N; int n = atom - b * N; int tid = threadIdx.x;
    const float* c = coords + ((b * N + n) * 3);
    float halfG = (float)G * 0.5f;
    int ix = (int)((c[0] - centers[b*3]) * inv_voxel + halfG);
    int iy = (int)((c[1] - centers[b*3+1]) * inv_voxel + halfG);
    int iz = (int)((c[2] - centers[b*3+2]) * inv_voxel + halfG);
    ix = ix < 0 ? 0 : (ix >= G ? G-1 : ix); iy = iy < 0 ? 0 : (iy >= G ? G-1 : iy); iz = iz < 0 ? 0 : (iz >= G ? G-1 : iz);
    int vox = (ix * G + iy) * G + iz;
    const float* fin = features + ((b * N + n) * F);
    for (int d = tid; d < F; d += blockDim.x) {
        float acc = bias[d]; const float* w = weight + d * F;
        for (int k = 0; k < F; ++k) acc = fmaf(fin[k], w[k], acc);
        size_t off = (((size_t)b * F + d) * (size_t)G * G * G) + vox;
        atomicAdd(out + off, acc);
    }
}

torch::Tensor voxel_forward(torch::Tensor coords, torch::Tensor features, torch::Tensor weight, torch::Tensor bias, int64_t grid_size, double voxel_size) {
    int B = (int)coords.size(0), N = (int)coords.size(1), F = (int)features.size(2), G = (int)grid_size;
    auto out = torch::empty({B, F, G, G, G}, coords.options());
    auto centers = torch::empty({B, 3}, coords.options());
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();
    cudaMemsetAsync(out.data_ptr<float>(), 0, (size_t)B * F * G * G * G * sizeof(float), stream);
    center_kernel<<<B, 256, 0, stream>>>(coords.data_ptr<float>(), centers.data_ptr<float>(), B, N);
    if (F == 32) {
        scatter_linear32_kernel<<<B*N, 32, 0, stream>>>(coords.data_ptr<float>(), features.data_ptr<float>(), weight.data_ptr<float>(), bias.data_ptr<float>(), centers.data_ptr<float>(), out.data_ptr<float>(), N, G, 1.0f/(float)voxel_size);
    } else {
        int threads = (F <= 32 ? 32 : (F <= 64 ? 64 : 128));
        scatter_linear_kernel<<<B*N, threads, 0, stream>>>(coords.data_ptr<float>(), features.data_ptr<float>(), weight.data_ptr<float>(), bias.data_ptr<float>(), centers.data_ptr<float>(), out.data_ptr<float>(), B, N, F, G, 1.0f/(float)voxel_size);
    }
    return out;
}
'''

_ext = load_inline(name='voxel_pool_ext_final', cpp_sources=cpp_src, cuda_sources=cuda_src, functions=['voxel_forward'], extra_cflags=['-O3'], extra_cuda_cflags=['-O3', '--use_fast_math'], verbose=False)

class ModelNew(nn.Module):
    def __init__(self, feat_dim, grid_size=32, voxel_size=1.0):
        super().__init__()
        self.feat_dim = feat_dim
        self.grid_size = grid_size
        self.voxel_size = voxel_size
        self.linear = nn.Linear(feat_dim, feat_dim)
    def forward(self, coords: torch.Tensor, features: torch.Tensor) -> torch.Tensor:
        return _ext.voxel_forward(coords.contiguous(), features.contiguous(), self.linear.weight.contiguous(), self.linear.bias.contiguous(), int(self.grid_size), float(self.voxel_size))