turn 0210.9s turn · 34.2s LLM · 176.7s tools · tokens in=1,434 · out=2,292 · total=4,162 · FINAL
reasoning
The task is to replace the forward computation of this PyTorch model with a custom CUDA kernel. The model is a stack of planar normalizing flow layers.
I need to implement a custom CUDA kernel for the forward pass of a stack of planar normalizing flow layers in PyTorch. The model uses PlanarFlowBuilder with parameters like num_flows=8, dim=32, batch=102 variable batch size up to 1024.
assistant```python
import torch
import torch.nn as nn
import torch.nn.functional as F
from torch.utils.cpp_extension import load_inline
cuda_source = r'''
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <math.h>
__device__ __forceinline__ float softplus(float x) {
return log1pf(expf(x));
}
__global__ void planar_flow_kernel(
const float* __restrict__ z_in,
float* __restrict__ z_out,
float* __restrict__ sum_log_det,
const float* __restrict__ w_all,
const float* __restrict__ u_all,
const float* __restrict__ b_all,
int batch_size,
int dim,
int num_flows
) {
int b = blockIdx.x * blockDim.x + threadIdx.x;
if (b >= batch_size) return;
float local_z[32]; // assume dim <= 32
for (int d = 0; d < dim; d++) {
local_z[d] = z_in[b * dim + d];
}
float log_det_acc = 0.0f;
for (int i = 0; i < num_flows; i++) {
const float* w = w_all + i * dim;
const float* u = u_all + i * dim;
float b = b_all[i];
// constrain u
float wtu = 0.0f;
float w_norm2 = 0.0f;
for (int d = 0; d < dim; d++) {
wtu += w[d] * u[d];
w_norm2 += w[d] * w[d];
}
float m = -1.0f + softplus(wtu);
float scale = (m - wtu) / (w_norm2 + 1e-12f);
float linear = b;
for (int d = 0; d < dim; d++) {
linear += local_z[d] * w[d];
}
float h = tanhf(linear);
float h_prime = 1.0f - h * h;
float psi_dot_u = 0.0f;
for (int d = 0; d < dim; d++) {
float u_hat_d = u[d] + scale * w[d];
float psi_d = h_prime * w[d];
local_z[d] += u_hat_d * h;
psi_dot_u += psi_d * u_hat_d;
}
float log_det = logf(fabsf(1.0f + psi_dot_u) + 1e-8f);
log_det_acc += log_det;
}
for (int d = 0; d < dim; d++) {
z_out[b * dim + d] = local_z[d];
}
sum_log_det[b] = log_det_acc;
}
torch::Tensor planar_flow_cuda(
torch::Tensor z,
torch::Tensor w_all,
torch::Tensor u_all,
torch::Tensor b_all,
int num_flows
) {
int batch_size = z.size(0);
int dim = z.size(1);
auto z_out = torch::empty_like(z);
auto sum_log_det = torch::zeros({batch_size}, z.options());
const int threads = 256;
const int blocks = (batch_size + threads - 1) / threads;
planar_flow_kernel<<<blocks, threads>>>(
z.data_ptr<float>(),
z_out.data_ptr<float>(),
sum_log_det.data_ptr<float>(),
w_all.data_ptr<float>(),
u_all.data_ptr<float>(),
b_all.data_ptr<float>(),
batch_size,
dim,
num_flows
);
return z_out, sum_log_det;
}
'''
cpp_source = r'''
#include <torch/extension.h>
std::tuple<torch::Tensor, torch::Tensor> planar_flow_cuda(
torch::Tensor z,
torch::Tensor w_all,
torch::Tensor u_all,
torch::Tensor b_all,
int num_flows
);
'''
module = load_inline(
name='planar_flow_ext',
cpp_sources=[cpp_source],
cuda_sources=[cuda_source],
functions=['planar_flow_cuda'],
verbose=False,
extra_cflags=['-O3']
)
class ModelNew(nn.Module):
def __init__(self, dim, num_flows):
super().__init__()
self.dim = dim
self.num_flows = num_flows
self.w = nn.ParameterList([nn.Parameter(torch.randn(dim)) for _ in range(num_flows)])
self.u = nn.ParameterList([nn.Parameter(torch.randn(dim)) for _ in range(num_flows)])
self.b = nn.ParameterList([nn.Parameter(torch.zeros(1)) for _ in range(num_flows)])
def forward(self, z: torch.Tensor) -> tuple:
w_all = torch.stack([w.data for w in self.w])
u_all = torch.stack([u.data for u in self.u])
b_all = torch.stack([b.data for b in self.b]).view(-1)
return module.planar_flow_cuda(z, w_all, u_all, b_all, self.num_flows)
```
The kernel uses register tiling for dim=32, fuses the entire flow stack per thread, and avoids Python overhead. It should be numerically equivalent and significantly faster.
call → submit_kernel
callsubmit_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\ncuda_source = r'''\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <math.h>\n\n__device__ __forceinline__ float softplus(float x) {\n return log1pf(expf(x));\n}\n\n__global__ void planar_flow_kernel(\n const float* __restrict__ z_in,\n float* __restrict__ z_out,\n float* __restrict__ sum_log_det,\n const float* __restrict__ w_all,\n const float* __restrict__ u_all,\n const float* __restrict__ b_all,\n int batch_size,\n int dim,\n int num_flows\n) {\n int b = blockIdx.x * blockDim.x + threadIdx.x;\n if (b >= batch_size) return;\n\n float local_z[32]; // assume dim <= 32\n for (int d = 0; d < dim; d++) {\n local_z[d] = z_in[b * dim + d];\n }\n\n float log_det_acc = 0.0f;\n\n for (int i = 0; i < num_flows; i++) {\n const float* w = w_all + i * dim;\n const float* u = u_all + i * dim;\n float b = b_all[i];\n\n // constrain u\n float wtu = 0.0f;\n float w_norm2 = 0.0f;\n for (int d = 0; d < dim; d++) {\n wtu += w[d] * u[d];\n w_norm2 += w[d] * w[d];\n }\n float m = -1.0f + softplus(wtu);\n float scale = (m - wtu) / (w_norm2 + 1e-12f);\n\n float linear = b;\n for (int d = 0; d < dim; d++) {\n linear += local_z[d] * w[d];\n }\n float h = tanhf(linear);\n float h_prime = 1.0f - h * h;\n\n float psi_dot_u = 0.0f;\n for (int d = 0; d < dim; d++) {\n float u_hat_d = u[d] + scale * w[d];\n float psi_d = h_prime * w[d];\n local_z[d] += u_hat_d * h;\n psi_dot_u += psi_d * u_hat_d;\n }\n float log_det = logf(fabsf(1.0f + psi_dot_u) + 1e-8f);\n log_det_acc += log_det;\n }\n\n for (int d = 0; d < dim; d++) {\n z_out[b * dim + d] = local_z[d];\n }\n sum_log_det[b] = log_det_acc;\n}\n\nstd::tuple<torch::Tensor, torch::Tensor> planar_flow_cuda(\n torch::Tensor z,\n torch::Tensor w_all,\n torch::Tensor u_all,\n torch::Tensor b_all,\n int num_flows\n) {\n int batch_size = z.size(0);\n int dim = z.size(1);\n auto z_out = torch::empty_like(z);\n auto sum_log_det = torch::zeros({batch_size}, z.options());\n\n const int threads = 256;\n const int blocks = (batch_size + threads - 1) / threads;\n\n planar_flow_kernel<<<blocks, threads>>>(\n z.data_ptr<float>(),\n z_out.data_ptr<float>(),\n sum_log_det.data_ptr<float>(),\n w_all.data_ptr<float>(),\n u_all.data_ptr<float>(),\n b_all.data_ptr<float>(),\n batch_size,\n dim,\n num_flows\n );\n\n return std::make_tuple(z_out, sum_log_det);\n}\n'''\n\ncpp_source = r'''\n#include <torch/extension.h>\nstd::tuple<torch::Tensor, torch::Tensor> planar_flow_cuda(\n torch::Tensor z,\n torch::Tensor w_all,\n torch::Tensor u_all,\n torch::Tensor b_all,\n int num_flows\n);\n'''\n\nmodule = load_inline(\n name='planar_flow_ext',\n cpp_sources=[cpp_source],\n cuda_sources=[cuda_source],\n functions=['planar_flow_cuda'],\n verbose=False,\n extra_cflags=['-O3']\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, dim, num_flows):\n super().__init__()\n self.dim = dim\n self.num_flows = num_flows\n self.w = nn.ParameterList([nn.Parameter(torch.randn(dim)) for _ in range(num_flows)])\n self.u = nn.ParameterList([nn.Parameter(torch.randn(dim)) for _ in range(num_flows)])\n self.b = nn.ParameterList([nn.Parameter(torch.zeros(1)) for _ in range(num_flows)])\n\n def forward(self, z: torch.Tensor) -> tuple:\n w_all = torch.stack([w.data for w in self.w])\n u_all = torch.stack([u.data for u in self.u])\n b_all = torch.stack([b.data for b in self.b]).view(-1)\n return module.planar_flow_cuda(z, w_all, u_all, b_all, self.num_flows)\n```"
}
result ← submit_kernel (fail)
failsubmit_kernel
submit_kernel FAILED: kernel did not compile.
Syntax error in custom generated code or ModelNew not found