correct
33_VectorizedPoseTransformFused.py
started
2026-05-11T15:54:57Z
timingkernelbench.timing
kernel mean / median / std
0.01 / — / 2.70e-04 μs
ref mean / median / std
0.25 / — / 0.02 μs
numerical precisioneval_kernel_against_ref
gpu memoryextended_metrics.measure_memory
kernel launches / fusiontorch.profiler
fusion ratio (ref / kernel)
45.50x
SOL (speed-of-light)nsight_hardware_counters
roofline / occupancynsight_hardware_counters
warp stalls (top 5)nsight
reference PyTorch source (the prompt the agent sees)
import torch
import torch.nn as nn
import torch.nn.functional as F
# Default ModelNew is PyTorch-only. run_and_check: use check_kernel=False with backend=cuda
# until the candidate file contains a real CUDA kernel (see kernel_static_checker).
def _quaternion_to_rotation_matrix(q: torch.Tensor) -> torch.Tensor:
"""
Unit quaternion ``q = (w, x, y, z)`` (scalar-first) to rotation matrix ``R``,
shape ``(B, 3, 3)``, with ``p' = p @ R.T`` for row vectors ``p`` (``N×3``).
"""
q = F.normalize(q, dim=-1, eps=1e-8)
w, x, y, z = q.unbind(dim=-1)
xx = x * x
yy = y * y
zz = z * z
xy = x * y
xz = x * z
yz = y * z
wx = w * x
wy = w * y
wz = w * z
m00 = 1.0 - 2.0 * (yy + zz)
m01 = 2.0 * (xy - wz)
m02 = 2.0 * (xz + wy)
m10 = 2.0 * (xy + wz)
m11 = 1.0 - 2.0 * (xx + zz)
m12 = 2.0 * (yz - wx)
m20 = 2.0 * (xz - wy)
m21 = 2.0 * (yz + wx)
m22 = 1.0 - 2.0 * (xx + yy)
row0 = torch.stack((m00, m01, m02), dim=-1)
row1 = torch.stack((m10, m11, m12), dim=-1)
row2 = torch.stack((m20, m21, m22), dim=-1)
return torch.stack((row0, row1, row2), dim=-2)
class Model(nn.Module):
"""
**Fused vectorized pose transform**: rotation from unit quaternion ``q`` and
translation ``t`` applied to a batch of 3D points.
One forward fuses **quaternion → ``3×3``** (Hamilton, scalar-first ``w,x,y,z``)
with **batched matmul** and **broadcast translation**—typical for LiDAR / camera
extrinsics, hand-eye calibration, and world ↔ body frame changes in VLA stacks.
``points' = points @ R.T + t`` with ``points`` row-wise ``(B, N, 3)``.
"""
def __init__(self):
super().__init__()
def forward(
self,
points: torch.Tensor,
quat: torch.Tensor,
trans: torch.Tensor,
) -> torch.Tensor:
"""
Args:
points: (B, N, 3) points in the **source** frame.
quat: (B, 4) unit quaternion ``(w, x, y, z)`` mapping source → target rotation.
trans: (B, 3) translation in the **target** frame (same convention as ``R``).
Returns:
(B, N, 3) points in the **target** frame.
"""
B, N, _ = points.shape
assert quat.shape == (B, 4) and trans.shape == (B, 3)
R = _quaternion_to_rotation_matrix(quat)
return torch.matmul(points, R.transpose(-1, -2)) + trans.unsqueeze(1)
class ModelNew(nn.Module):
"""KernelBench candidate entry point; default delegates to ``Model``."""
def __init__(self):
super().__init__()
self._impl = Model()
def forward(
self,
points: torch.Tensor,
quat: torch.Tensor,
trans: torch.Tensor,
) -> torch.Tensor:
return self._impl(points, quat, trans)
batch_size = 8
num_points = 512
def get_inputs():
p = popcorn_pri
points = torch.randn(p.jitter_int(batch_size), p.jitter_int(num_points), 3)
quat = torch.randn(p.jitter_int(batch_size), 4)
trans = torch.randn(p.jitter_int(batch_size), 3)
return [points, quat, trans]
def get_init_inputs():
return []
turn 03.0s turn · 3.0s LLM · 0.0s tools · tokens in=2,185 · out=188 · total=2,373
reasoning
Let me understand the task:
1. I need to write a custom CUDA kernel that replaces the PyTorch model's forward computation
2. The model takes:
- `points`: (B, N, 3) - batch of 3D points
- `quat`: (B, 4) - unit quaternions (w, x, y, z)
- `trans`: (B, 3) - translations
3. The forward pass:
- Converts quaternion to rotation matrix R (3x3)
- Applies rotation: `points @ R.T`
- Adds translation: `+ trans.unsqueeze(1)`
Let me first get the GPU specs to understand the hardware I'm working with.
assistantI'll start by understanding the GPU hardware specs and then develop a custom CUDA kernel for this quaternion-to-rotation-matrix transform with point transformation.
call → get_gpu_specs
callget_gpu_specs
{}
result ← get_gpu_specs (ok)
okget_gpu_specs
GPU specs for NVIDIA H100 80GB HBM3:
total memory (runtime): 79.1 GB
spec entry: H100
GPU Architecture: Hopper
GPU Memory: 80GB
Memory Bandwidth: 3.35 TB/s
FP64 TFLOPS: 34
FP64 Tensor Core TFLOPS: 67
FP32 TFLOPS: 67
TF32 Tensor Core TFLOPS: 989 with sparsity
BFLOAT16 Tensore Core TFLOPS: 1979 with sparsity
FP16 Tensor Core TFLOPS: 1979 with sparsity
FP8 Tensor Core TFLOPS: 3958 with sparsity
INT8 Tensor Core TOPS: 3958 with sparsity
Register File Size: 64K 32-bit registers per SM
Maximum number of registers per thread: 255
Maximum number of thread blocks per SM: 32
Shared memory capacity per SM: 228 KB
Maximum shared memory per thread block: 227 KB
turn 1 · 1 tool call(s) · 99.7s turn · 24.6s LLM · tokens in=2,462 · out=2,462 · total=4,924 · 9,403 chars (click to expand)
turn 199.7s turn · 24.6s LLM · 75.1s tools · tokens in=2,462 · out=2,462 · total=4,924
reasoning
Now let me design the CUDA kernel. The problem involves:
1. Converting quaternions to rotation matrices (B, 4) → (B, 3, 3)
2. Transforming points: points @ R.T + trans, where points is (B, N, 3), R.T is (B, 3, 3), and trans is (B, 3)
For optimal performance, I can fuse these operations:
- Each thread can handle one point (or multiple points)
- The rotation matrix needs to be computed once per batch element
- Then each point is transformed using that rotation matrix
Strategy:
1. Compute the rotation matrix in shared memory or registers for each batch element
2. Transform all points for that batch element
Let me write a kernel that:
- Launches with (B, N) threads
- Each block handles one batch element
- Threads within a block cooperate to compute the rotation matrix (using shared memory)
- Then each thread transforms its point(s)
Actually, let me think about this more carefully:
- Batch size = 8, num_points = 512
- Total points = 8 * 512 = 4096
The rotation matrix computation is very lightweight (just a few operations per batch element).
The point transformation is just a 3x3 matrix multiplication plus translation.
For a simple approach:
- Each thread handles one point
- Threads cooperate to compute the rotation matrix using shared memory
Let me write the kernel:
```cuda
__global__ void quaternion_transform_kernel(
const float* __restrict__ points, // (B, N, 3)
const float* __restrict__ quat, // (B, 4)
const float* __restrict__ trans, // (B, 3)
float* __restrict__ output, // (B, N, 3)
int batch_size, int num_points
) {
// Shared memory for the rotation matrix and translation
__shared__ float R[9]; // 3x3 rotation matrix, row-major
__shared__ float t[3]; // translation
int b = blockIdx.x; // batch index
int tid = threadIdx.x;
// Load quaternion and compute rotation matrix (first few threads)
if (tid < 4) {
// Load quaternion
}
__syncthreads();
// Compute rotation matrix elements
// ...
// Transform points
int n = tid; // point index
if (n < num_points) {
// output[b, n, :] = points[b, n, :] @ R.T + t
}
}
```
Actually, let me think about this differently. Since N=512 and threads per block is typically up to 1024, I can use:
- Grid dimension: batch_size blocks
- Block dimension: at least num_points threads (512)
- Each block handles one batch element
The rotation matrix computation is very small - just 4 values need to be loaded and 9 values computed. The first few threads can do this and store in shared memory.
Let me write a complete kernel:
assistantNow I'll write a custom CUDA kernel for this operation. The key insight is:
1. Quaternion-to-rotation-matrix is per-batch (only 8 computations)
2. Point transformation is per-point (4096 operations)
3. I can fuse these: compute rotation matrix in shared memory per block, then transform all points
Let me write and compile the kernel:
call → compile_kernel
callcompile_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\n# CUDA source for quaternion transform\ncuda_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <math.h>\n\n__global__ void quaternion_transform_kernel(\n const float* __restrict__ points, // (B, N, 3) - flattened\n const float* __restrict__ quat, // (B, 4) - flattened\n const float* __restrict__ trans, // (B, 3) - flattened\n float* __restrict__ output, // (B, N, 3) - flattened\n int batch_size, int num_points\n) {\n // Shared memory for rotation matrix (3x3) and translation (3)\n __shared__ float R[9]; // row-major rotation matrix\n __shared__ float t[3]; // translation\n \n int b = blockIdx.x; // batch index (one block per batch element)\n int n = threadIdx.x; // point index\n \n // Load quaternion and compute rotation matrix (first 4 threads)\n if (n < 4) {\n float q_val = quat[b * 4 + n];\n // Normalize quaternion (compute squared norm first)\n // We'll do normalization after loading all 4 components\n }\n __syncthreads();\n \n // First thread computes normalized quaternion and rotation matrix\n if (n == 0) {\n float qw = quat[b * 4 + 0];\n float qx = quat[b * 4 + 1];\n float qy = quat[b * 4 + 2];\n float qz = quat[b * 4 + 3];\n \n // Normalize\n float norm = sqrtf(qw*qw + qx*qx + qy*qy + qz*qz);\n float inv_norm = 1.0f / (norm + 1e-8f);\n qw *= inv_norm;\n qx *= inv_norm;\n qy *= inv_norm;\n qz *= inv_norm;\n \n // Compute rotation matrix elements\n float xx = qx * qx;\n float yy = qy * qy;\n float zz = qz * qz;\n float xy = qx * qy;\n float xz = qx * qz;\n float yz = qy * qz;\n float wx = qw * qx;\n float wy = qw * qy;\n float wz = qw * qz;\n \n R[0] = 1.0f - 2.0f * (yy + zz); // m00\n R[1] = 2.0f * (xy - wz); // m01\n R[2] = 2.0f * (xz + wy); // m02\n R[3] = 2.0f * (xy + wz); // m10\n R[4] = 1.0f - 2.0f * (xx + zz); // m11\n R[5] = 2.0f * (yz - wx); // m12\n R[6] = 2.0f * (xz - wy); // m20\n R[7] = 2.0f * (yz + wx); // m21\n R[8] = 1.0f - 2.0f * (xx + yy); // m22\n }\n \n // Load translation (first 3 threads)\n if (n < 3) {\n t[n] = trans[b * 3 + n];\n }\n __syncthreads();\n \n // Transform points\n // output[b, n, :] = points[b, n, :] @ R.T + t\n // R.T means we use R[col * 3 + row] instead of R[row * 3 + col]\n if (n < num_points) {\n float px = points[b * num_points * 3 + n * 3 + 0];\n float py = points[b * num_points * 3 + n * 3 + 1];\n float pz = points[b * num_points * 3 + n * 3 + 2];\n \n // Matrix multiply: p @ R.T (which is p @ R^T)\n // R^T is column-major view of R, so result[k] = sum(p[j] * R[k*3 + j])\n // Actually, for p @ R.T, each output component is:\n // out[0] = p[0]*R[0] + p[1]*R[3] + p[2]*R[6] (first column of R.T = first row of R)\n // out[1] = p[0]*R[1] + p[1]*R[4] + p[2]*R[7]\n // out[2] = p[0]*R[2] + p[1]*R[5] + p[2]*R[8]\n // This is p @ R (not p @ R.T!), so let me reconsider.\n \n // For row vectors: p' = p @ R.T, where R is row-major\n // R.T transposes rows and columns, so R.T[j,k] = R[k,j]\n // p'[i] = sum_j(p[j] * R.T[j,i]) = sum_j(p[j] * R[i,j])\n // So: out[0] = p[0]*R[0] + p[1]*R[1] + p[2]*R[2]\n // out[1] = p[0]*R[3] + p[1]*R[4] + p[2]*R[5]\n // out[2] = p[0]*R[6] + p[1]*R[7] + p[2]*R[8]\n \n float out0 = px * R[0] + py * R[1] + pz * R[2] + t[0];\n float out1 = px * R[3] + py * R[4] + pz * R[5] + t[1];\n float out2 = px * R[6] + py * R[7] + pz * R[8] + t[2];\n \n output[b * num_points * 3 + n * 3 + 0] = out0;\n output[b * num_points * 3 + n * 3 + 1] = out1;\n output[b * num_points * 3 + n * 3 + 2] = out2;\n }\n}\n\ntorch::Tensor quaternion_transform_cuda(\n torch::Tensor points,\n torch::Tensor quat,\n torch::Tensor trans\n) {\n int batch_size = points.size(0);\n int num_points = points.size(1);\n \n auto output = torch::empty_like(points);\n \n int threads_per_block = num_points;\n int num_blocks = batch_size;\n \n quaternion_transform_kernel<<<num_blocks, threads_per_block>>>(\n points.data_ptr<float>(),\n quat.data_ptr<float>(),\n trans.data_ptr<float>(),\n output.data_ptr<float>(),\n batch_size,\n num_points\n );\n \n return output;\n}\n\"\"\"\n\n# Load the CUDA extension\nquaternion_transform_ext = load_inline(\n name='quaternion_transform_ext',\n cpp_sources='',\n cuda_sources=cuda_source,\n functions=['quaternion_transform_cuda'],\n extra_cuda_cflags=['-O3'],\n verbose=True\n)\n\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.ext = quaternion_transform_ext\n\n def forward(\n self,\n points: torch.Tensor,\n quat: torch.Tensor,\n trans: torch.Tensor,\n ) -> torch.Tensor:\n return self.ext.quaternion_transform_cuda(points, quat, trans)"
}
result ← compile_kernel (fail)
failcompile_kernel
compile_kernel FAILED: RuntimeError.
RuntimeError: Error building extension 'quaternion_transform_ext'
turn 285.7s turn · 10.1s LLM · 75.6s tools · tokens in=4,283 · out=1,158 · total=5,441
reasoning
The build failed with a generic error. Let me try to get more details or simplify the kernel. Let me try again with verbose=True to see the actual error message. Actually, I already had verbose=True. Let me try a cleaner approach and add error checking.
Let me rewrite with better handling and see if there's a more specific error.
assistantLet me try a simpler version with better error handling:
call → compile_kernel
callcompile_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\n# CUDA source for quaternion transform\ncuda_source = r\"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <cmath>\n\n__global__ void quaternion_transform_kernel(\n const float* __restrict__ points,\n const float* __restrict__ quat,\n const float* __restrict__ trans,\n float* __restrict__ output,\n int batch_size, int num_points\n) {\n __shared__ float R[9];\n __shared__ float t[3];\n \n int b = blockIdx.x;\n int n = threadIdx.x;\n \n // First thread computes rotation matrix\n if (n == 0) {\n float qw = quat[b * 4 + 0];\n float qx = quat[b * 4 + 1];\n float qy = quat[b * 4 + 2];\n float qz = quat[b * 4 + 3];\n \n float norm = sqrtf(qw*qw + qx*qx + qy*qy + qz*qz);\n float inv_norm = 1.0f / (norm + 1e-8f);\n qw *= inv_norm;\n qx *= inv_norm;\n qy *= inv_norm;\n qz *= inv_norm;\n \n float xx = qx * qx;\n float yy = qy * qy;\n float zz = qz * qz;\n float xy = qx * qy;\n float xz = qx * qz;\n float yz = qy * qz;\n float wx = qw * qx;\n float wy = qw * qy;\n float wz = qw * qz;\n \n R[0] = 1.0f - 2.0f * (yy + zz);\n R[1] = 2.0f * (xy - wz);\n R[2] = 2.0f * (xz + wy);\n R[3] = 2.0f * (xy + wz);\n R[4] = 1.0f - 2.0f * (xx + zz);\n R[5] = 2.0f * (yz - wx);\n R[6] = 2.0f * (xz - wy);\n R[7] = 2.0f * (yz + wx);\n R[8] = 1.0f - 2.0f * (xx + yy);\n }\n \n if (n < 3) {\n t[n] = trans[b * 3 + n];\n }\n __syncthreads();\n \n if (n < num_points) {\n float px = points[b * num_points * 3 + n * 3 + 0];\n float py = points[b * num_points * 3 + n * 3 + 1];\n float pz = points[b * num_points * 3 + n * 3 + 2];\n \n float out0 = px * R[0] + py * R[1] + pz * R[2] + t[0];\n float out1 = px * R[3] + py * R[4] + pz * R[5] + t[1];\n float out2 = px * R[6] + py * R[7] + pz * R[8] + t[2];\n \n output[b * num_points * 3 + n * 3 + 0] = out0;\n output[b * num_points * 3 + n * 3 + 1] = out1;\n output[b * num_points * 3 + n * 3 + 2] = out2;\n }\n}\n\ntorch::Tensor quaternion_transform_cuda(\n torch::Tensor points,\n torch::Tensor quat,\n torch::Tensor trans\n) {\n int batch_size = points.size(0);\n int num_points = points.size(1);\n \n auto output = torch::empty_like(points);\n \n int threads_per_block = num_points;\n int num_blocks = batch_size;\n \n quaternion_transform_kernel<<<num_blocks, threads_per_block>>>(\n points.data_ptr<float>(),\n quat.data_ptr<float>(),\n trans.data_ptr<float>(),\n output.data_ptr<float>(),\n batch_size,\n num_points\n );\n \n return output;\n}\n\"\"\"\n\n# Load the CUDA extension\nquaternion_transform_ext = load_inline(\n name='quaternion_transform_ext',\n cpp_sources='',\n cuda_sources=cuda_source,\n functions=['quaternion_transform_cuda'],\n extra_cuda_cflags=['-O3'],\n verbose=True\n)\n\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.ext = quaternion_transform_ext\n\n def forward(\n self,\n points: torch.Tensor,\n quat: torch.Tensor,\n trans: torch.Tensor,\n ) -> torch.Tensor:\n return self.ext.quaternion_transform_cuda(points, quat, trans)"
}
result ← compile_kernel (fail)
failcompile_kernel
compile_kernel FAILED: RuntimeError.
RuntimeError: Error building extension 'quaternion_transform_ext_v1'
turn 383.9s turn · 10.0s LLM · 73.9s tools · tokens in=5,395 · out=1,103 · total=6,498
reasoning
The compilation still fails. Let me try a different approach - maybe I need to specify more compiler flags or handle the torch extension differently. Let me check if there's an issue with the verbose output or the extension loading.
assistantLet me try with explicit cpp_sources to help debug:
call → compile_kernel
callcompile_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = r\"\"\"\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n\n__global__ void quaternion_transform_kernel(\n const float* points,\n const float* quat,\n const float* trans,\n float* output,\n int batch_size, int num_points\n) {\n __shared__ float R[9];\n __shared__ float t[3];\n \n int b = blockIdx.x;\n int n = threadIdx.x;\n \n if (n == 0) {\n float qw = quat[b * 4 + 0];\n float qx = quat[b * 4 + 1];\n float qy = quat[b * 4 + 2];\n float qz = quat[b * 4 + 3];\n \n float norm = sqrtf(qw*qw + qx*qx + qy*qy + qz*qz);\n float inv_norm = 1.0f / (norm + 1e-8f);\n qw *= inv_norm;\n qx *= inv_norm;\n qy *= inv_norm;\n qz *= inv_norm;\n \n float xx = qx * qx;\n float yy = qy * qy;\n float zz = qz * qz;\n float xy = qx * qy;\n float xz = qx * qz;\n float yz = qy * qz;\n float wx = qw * qx;\n float wy = qw * qy;\n float wz = qw * qz;\n \n R[0] = 1.0f - 2.0f * (yy + zz);\n R[1] = 2.0f * (xy - wz);\n R[2] = 2.0f * (xz + wy);\n R[3] = 2.0f * (xy + wz);\n R[4] = 1.0f - 2.0f * (xx + zz);\n R[5] = 2.0f * (yz - wx);\n R[6] = 2.0f * (xz - wy);\n R[7] = 2.0f * (yz + wx);\n R[8] = 1.0f - 2.0f * (xx + yy);\n }\n \n if (n < 3) {\n t[n] = trans[b * 3 + n];\n }\n __syncthreads();\n \n if (n < num_points) {\n float px = points[b * num_points * 3 + n * 3 + 0];\n float py = points[b * num_points * 3 + n * 3 + 1];\n float pz = points[b * num_points * 3 + n * 3 + 2];\n \n output[b * num_points * 3 + n * 3 + 0] = px * R[0] + py * R[1] + pz * R[2] + t[0];\n output[b * num_points * 3 + n * 3 + 1] = px * R[3] + py * R[4] + pz * R[5] + t[1];\n output[b * num_points * 3 + n * 3 + 2] = px * R[6] + py * R[7] + pz * R[8] + t[2];\n }\n}\n\ntorch::Tensor quaternion_transform_cuda(\n torch::Tensor points,\n torch::Tensor quat,\n torch::Tensor trans\n) {\n int batch_size = points.size(0);\n int num_points = points.size(1);\n \n auto output = torch::empty_like(points);\n \n int threads_per_block = num_points;\n int num_blocks = batch_size;\n \n quaternion_transform_kernel<<<num_blocks, threads_per_block>>>(\n points.data_ptr<float>(),\n quat.data_ptr<float>(),\n trans.data_ptr<float>(),\n output.data_ptr<float>(),\n batch_size,\n num_points\n );\n \n return output;\n}\n\"\"\"\n\ncpp_source = r\"\"\"\ntorch::Tensor quaternion_transform_cuda(\n torch::Tensor points,\n torch::Tensor quat,\n torch::Tensor trans\n);\n\"\"\"\n\nquaternion_transform_ext = load_inline(\n name='quaternion_transform_ext_v2',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['quaternion_transform_cuda'],\n extra_cuda_cflags=['-O3'],\n verbose=True\n)\n\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.ext = quaternion_transform_ext\n\n def forward(\n self,\n points: torch.Tensor,\n quat: torch.Tensor,\n trans: torch.Tensor,\n ) -> torch.Tensor:\n return self.ext.quaternion_transform_cuda(points, quat, trans)"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 412.5s turn · 10.1s LLM · 2.4s tools · tokens in=6,467 · out=1,069 · total=7,536
reasoning
Good, it compiled. Now let me run correctness check.
assistantNow let me run the correctness check:
call → run_correctness
callrun_correctness
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = r\"\"\"\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n\n__global__ void quaternion_transform_kernel(\n const float* points,\n const float* quat,\n const float* trans,\n float* output,\n int batch_size, int num_points\n) {\n __shared__ float R[9];\n __shared__ float t[3];\n \n int b = blockIdx.x;\n int n = threadIdx.x;\n \n if (n == 0) {\n float qw = quat[b * 4 + 0];\n float qx = quat[b * 4 + 1];\n float qy = quat[b * 4 + 2];\n float qz = quat[b * 4 + 3];\n \n float norm = sqrtf(qw*qw + qx*qx + qy*qy + qz*qz);\n float inv_norm = 1.0f / (norm + 1e-8f);\n qw *= inv_norm;\n qx *= inv_norm;\n qy *= inv_norm;\n qz *= inv_norm;\n \n float xx = qx * qx;\n float yy = qy * qy;\n float zz = qz * qz;\n float xy = qx * qy;\n float xz = qx * qz;\n float yz = qy * qz;\n float wx = qw * qx;\n float wy = qw * qy;\n float wz = qw * qz;\n \n R[0] = 1.0f - 2.0f * (yy + zz);\n R[1] = 2.0f * (xy - wz);\n R[2] = 2.0f * (xz + wy);\n R[3] = 2.0f * (xy + wz);\n R[4] = 1.0f - 2.0f * (xx + zz);\n R[5] = 2.0f * (yz - wx);\n R[6] = 2.0f * (xz - wy);\n R[7] = 2.0f * (yz + wx);\n R[8] = 1.0f - 2.0f * (xx + yy);\n }\n \n if (n < 3) {\n t[n] = trans[b * 3 + n];\n }\n __syncthreads();\n \n if (n < num_points) {\n float px = points[b * num_points * 3 + n * 3 + 0];\n float py = points[b * num_points * 3 + n * 3 + 1];\n float pz = points[b * num_points * 3 + n * 3 + 2];\n \n output[b * num_points * 3 + n * 3 + 0] = px * R[0] + py * R[1] + pz * R[2] + t[0];\n output[b * num_points * 3 + n * 3 + 1] = px * R[3] + py * R[4] + pz * R[5] + t[1];\n output[b * num_points * 3 + n * 3 + 2] = px * R[6] + py * R[7] + pz * R[8] + t[2];\n }\n}\n\ntorch::Tensor quaternion_transform_cuda(\n torch::Tensor points,\n torch::Tensor quat,\n torch::Tensor trans\n) {\n int batch_size = points.size(0);\n int num_points = points.size(1);\n \n auto output = torch::empty_like(points);\n \n int threads_per_block = num_points;\n int num_blocks = batch_size;\n \n quaternion_transform_kernel<<<num_blocks, threads_per_block>>>(\n points.data_ptr<float>(),\n quat.data_ptr<float>(),\n trans.data_ptr<float>(),\n output.data_ptr<float>(),\n batch_size,\n num_points\n );\n \n return output;\n}\n\"\"\"\n\ncpp_source = r\"\"\"\ntorch::Tensor quaternion_transform_cuda(\n torch::Tensor points,\n torch::Tensor quat,\n torch::Tensor trans\n);\n\"\"\"\n\nquaternion_transform_ext = load_inline(\n name='quaternion_transform_ext_v2',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['quaternion_transform_cuda'],\n extra_cuda_cflags=['-O3'],\n verbose=True\n)\n\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.ext = quaternion_transform_ext\n\n def forward(\n self,\n points: torch.Tensor,\n quat: torch.Tensor,\n trans: torch.Tensor,\n ) -> torch.Tensor:\n return self.ext.quaternion_transform_cuda(points, quat, trans)"
}
result ← run_correctness (ok)
okrun_correctness
run_correctness PASSED: (5 / 5) trials all matched the reference.
Numerical precision: max_abs_err=1.43e-06 mean_abs_err=9.40e-08 max_rel_err=2.62e-03
turn 59.4s turn · 9.3s LLM · 0.0s tools · tokens in=7,581 · out=1,077 · total=8,658
reasoning
The kernel is correct. Now let me run a static check and then submit it.
assistantThe kernel is correct. Let me run a static check:
call → static_check
callstatic_check
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = r\"\"\"\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n\n__global__ void quaternion_transform_kernel(\n const float* points,\n const float* quat,\n const float* trans,\n float* output,\n int batch_size, int num_points\n) {\n __shared__ float R[9];\n __shared__ float t[3];\n \n int b = blockIdx.x;\n int n = threadIdx.x;\n \n if (n == 0) {\n float qw = quat[b * 4 + 0];\n float qx = quat[b * 4 + 1];\n float qy = quat[b * 4 + 2];\n float qz = quat[b * 4 + 3];\n \n float norm = sqrtf(qw*qw + qx*qx + qy*qy + qz*qz);\n float inv_norm = 1.0f / (norm + 1e-8f);\n qw *= inv_norm;\n qx *= inv_norm;\n qy *= inv_norm;\n qz *= inv_norm;\n \n float xx = qx * qx;\n float yy = qy * qy;\n float zz = qz * qz;\n float xy = qx * qy;\n float xz = qx * qz;\n float yz = qy * qz;\n float wx = qw * qx;\n float wy = qw * qy;\n float wz = qw * qz;\n \n R[0] = 1.0f - 2.0f * (yy + zz);\n R[1] = 2.0f * (xy - wz);\n R[2] = 2.0f * (xz + wy);\n R[3] = 2.0f * (xy + wz);\n R[4] = 1.0f - 2.0f * (xx + zz);\n R[5] = 2.0f * (yz - wx);\n R[6] = 2.0f * (xz - wy);\n R[7] = 2.0f * (yz + wx);\n R[8] = 1.0f - 2.0f * (xx + yy);\n }\n \n if (n < 3) {\n t[n] = trans[b * 3 + n];\n }\n __syncthreads();\n \n if (n < num_points) {\n float px = points[b * num_points * 3 + n * 3 + 0];\n float py = points[b * num_points * 3 + n * 3 + 1];\n float pz = points[b * num_points * 3 + n * 3 + 2];\n \n output[b * num_points * 3 + n * 3 + 0] = px * R[0] + py * R[1] + pz * R[2] + t[0];\n output[b * num_points * 3 + n * 3 + 1] = px * R[3] + py * R[4] + pz * R[5] + t[1];\n output[b * num_points * 3 + n * 3 + 2] = px * R[6] + py * R[7] + pz * R[8] + t[2];\n }\n}\n\ntorch::Tensor quaternion_transform_cuda(\n torch::Tensor points,\n torch::Tensor quat,\n torch::Tensor trans\n) {\n int batch_size = points.size(0);\n int num_points = points.size(1);\n \n auto output = torch::empty_like(points);\n \n int threads_per_block = num_points;\n int num_blocks = batch_size;\n \n quaternion_transform_kernel<<<num_blocks, threads_per_block>>>(\n points.data_ptr<float>(),\n quat.data_ptr<float>(),\n trans.data_ptr<float>(),\n output.data_ptr<float>(),\n batch_size,\n num_points\n );\n \n return output;\n}\n\"\"\"\n\ncpp_source = r\"\"\"\ntorch::Tensor quaternion_transform_cuda(\n torch::Tensor points,\n torch::Tensor quat,\n torch::Tensor trans\n);\n\"\"\"\n\nquaternion_transform_ext = load_inline(\n name='quaternion_transform_ext_v2',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['quaternion_transform_cuda'],\n extra_cuda_cflags=['-O3'],\n verbose=True\n)\n\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.ext = quaternion_transform_ext\n\n def forward(\n self,\n points: torch.Tensor,\n quat: torch.Tensor,\n trans: torch.Tensor,\n ) -> torch.Tensor:\n return self.ext.quaternion_transform_cuda(points, quat, trans)"
}
result ← static_check (ok)
okstatic_check
static_check PASSED: no violations or warnings detected.
turn 658.7s turn · 9.9s LLM · 48.8s tools · tokens in=8,655 · out=1,119 · total=9,774 · FINAL
reasoning
The kernel compiles, passes correctness check, and passes static analysis. Let me submit it to get the timing. The current kernel uses one block per batch element with one thread per point, which is fine for the current problem size (8 blocks with 512 threads each).
assistantThe kernel passes all checks. Let me submit it to get the final timing:
call → submit_kernel
callsubmit_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = r\"\"\"\n#include <torch/extension.h>\n#include <cuda_runtime.h>\n\n__global__ void quaternion_transform_kernel(\n const float* points,\n const float* quat,\n const float* trans,\n float* output,\n int batch_size, int num_points\n) {\n __shared__ float R[9];\n __shared__ float t[3];\n \n int b = blockIdx.x;\n int n = threadIdx.x;\n \n if (n == 0) {\n float qw = quat[b * 4 + 0];\n float qx = quat[b * 4 + 1];\n float qy = quat[b * 4 + 2];\n float qz = quat[b * 4 + 3];\n \n float norm = sqrtf(qw*qw + qx*qx + qy*qy + qz*qz);\n float inv_norm = 1.0f / (norm + 1e-8f);\n qw *= inv_norm;\n qx *= inv_norm;\n qy *= inv_norm;\n qz *= inv_norm;\n \n float xx = qx * qx;\n float yy = qy * qy;\n float zz = qz * qz;\n float xy = qx * qy;\n float xz = qx * qz;\n float yz = qy * qz;\n float wx = qw * qx;\n float wy = qw * qy;\n float wz = qw * qz;\n \n R[0] = 1.0f - 2.0f * (yy + zz);\n R[1] = 2.0f * (xy - wz);\n R[2] = 2.0f * (xz + wy);\n R[3] = 2.0f * (xy + wz);\n R[4] = 1.0f - 2.0f * (xx + zz);\n R[5] = 2.0f * (yz - wx);\n R[6] = 2.0f * (xz - wy);\n R[7] = 2.0f * (yz + wx);\n R[8] = 1.0f - 2.0f * (xx + yy);\n }\n \n if (n < 3) {\n t[n] = trans[b * 3 + n];\n }\n __syncthreads();\n \n if (n < num_points) {\n float px = points[b * num_points * 3 + n * 3 + 0];\n float py = points[b * num_points * 3 + n * 3 + 1];\n float pz = points[b * num_points * 3 + n * 3 + 2];\n \n output[b * num_points * 3 + n * 3 + 0] = px * R[0] + py * R[1] + pz * R[2] + t[0];\n output[b * num_points * 3 + n * 3 + 1] = px * R[3] + py * R[4] + pz * R[5] + t[1];\n output[b * num_points * 3 + n * 3 + 2] = px * R[6] + py * R[7] + pz * R[8] + t[2];\n }\n}\n\ntorch::Tensor quaternion_transform_cuda(\n torch::Tensor points,\n torch::Tensor quat,\n torch::Tensor trans\n) {\n int batch_size = points.size(0);\n int num_points = points.size(1);\n \n auto output = torch::empty_like(points);\n \n int threads_per_block = num_points;\n int num_blocks = batch_size;\n \n quaternion_transform_kernel<<<num_blocks, threads_per_block>>>(\n points.data_ptr<float>(),\n quat.data_ptr<float>(),\n trans.data_ptr<float>(),\n output.data_ptr<float>(),\n batch_size,\n num_points\n );\n \n return output;\n}\n\"\"\"\n\ncpp_source = r\"\"\"\ntorch::Tensor quaternion_transform_cuda(\n torch::Tensor points,\n torch::Tensor quat,\n torch::Tensor trans\n);\n\"\"\"\n\nquaternion_transform_ext = load_inline(\n name='quaternion_transform_ext_v2',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['quaternion_transform_cuda'],\n extra_cuda_cflags=['-O3'],\n verbose=True\n)\n\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n self.ext = quaternion_transform_ext\n\n def forward(\n self,\n points: torch.Tensor,\n quat: torch.Tensor,\n trans: torch.Tensor,\n ) -> torch.Tensor:\n return self.ext.quaternion_transform_cuda(points, quat, trans)"
}
result ← submit_kernel (ok)
oksubmit_kernel
submit_kernel PASSED: (5 / 5) correctness trials all passed.
Kernel runtime: 0.01 μs
Runtime stats: mean=0.01μs median=?μs std=0.00μs
Numerical precision: max_abs_err=1.43e-06 mean_abs_err=9.40e-08 max_rel_err=2.62e-03
Memory: 32.09 MB peak (ref: 32.14 MB, ratio: 0.9985x)
Kernel launches: 2 (ref: 91, fusion ratio: 45.5)
SOL score: 0.0044
Flagged for excessive speedup — this submission may have been rejected by automated review.
final submitted kernel (127 lines, 3,312 chars) — full source, untruncated
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline
cuda_source = r"""
#include <torch/extension.h>
#include <cuda_runtime.h>
__global__ void quaternion_transform_kernel(
const float* points,
const float* quat,
const float* trans,
float* output,
int batch_size, int num_points
) {
__shared__ float R[9];
__shared__ float t[3];
int b = blockIdx.x;
int n = threadIdx.x;
if (n == 0) {
float qw = quat[b * 4 + 0];
float qx = quat[b * 4 + 1];
float qy = quat[b * 4 + 2];
float qz = quat[b * 4 + 3];
float norm = sqrtf(qw*qw + qx*qx + qy*qy + qz*qz);
float inv_norm = 1.0f / (norm + 1e-8f);
qw *= inv_norm;
qx *= inv_norm;
qy *= inv_norm;
qz *= inv_norm;
float xx = qx * qx;
float yy = qy * qy;
float zz = qz * qz;
float xy = qx * qy;
float xz = qx * qz;
float yz = qy * qz;
float wx = qw * qx;
float wy = qw * qy;
float wz = qw * qz;
R[0] = 1.0f - 2.0f * (yy + zz);
R[1] = 2.0f * (xy - wz);
R[2] = 2.0f * (xz + wy);
R[3] = 2.0f * (xy + wz);
R[4] = 1.0f - 2.0f * (xx + zz);
R[5] = 2.0f * (yz - wx);
R[6] = 2.0f * (xz - wy);
R[7] = 2.0f * (yz + wx);
R[8] = 1.0f - 2.0f * (xx + yy);
}
if (n < 3) {
t[n] = trans[b * 3 + n];
}
__syncthreads();
if (n < num_points) {
float px = points[b * num_points * 3 + n * 3 + 0];
float py = points[b * num_points * 3 + n * 3 + 1];
float pz = points[b * num_points * 3 + n * 3 + 2];
output[b * num_points * 3 + n * 3 + 0] = px * R[0] + py * R[1] + pz * R[2] + t[0];
output[b * num_points * 3 + n * 3 + 1] = px * R[3] + py * R[4] + pz * R[5] + t[1];
output[b * num_points * 3 + n * 3 + 2] = px * R[6] + py * R[7] + pz * R[8] + t[2];
}
}
torch::Tensor quaternion_transform_cuda(
torch::Tensor points,
torch::Tensor quat,
torch::Tensor trans
) {
int batch_size = points.size(0);
int num_points = points.size(1);
auto output = torch::empty_like(points);
int threads_per_block = num_points;
int num_blocks = batch_size;
quaternion_transform_kernel<<<num_blocks, threads_per_block>>>(
points.data_ptr<float>(),
quat.data_ptr<float>(),
trans.data_ptr<float>(),
output.data_ptr<float>(),
batch_size,
num_points
);
return output;
}
"""
cpp_source = r"""
torch::Tensor quaternion_transform_cuda(
torch::Tensor points,
torch::Tensor quat,
torch::Tensor trans
);
"""
quaternion_transform_ext = load_inline(
name='quaternion_transform_ext_v2',
cpp_sources=cpp_source,
cuda_sources=cuda_source,
functions=['quaternion_transform_cuda'],
extra_cuda_cflags=['-O3'],
verbose=True
)
class ModelNew(nn.Module):
def __init__(self):
super().__init__()
self.ext = quaternion_transform_ext
def forward(
self,
points: torch.Tensor,
quat: torch.Tensor,
trans: torch.Tensor,
) -> torch.Tensor:
return self.ext.quaternion_transform_cuda(points, quat, trans)