correct
4_GJKRobotEnvironmentIntersection.py
started
2026-05-02T07:13:41Z
timingkernelbench.timing
kernel mean / median / std
0.01 / — / 1.28e-04 μs
ref mean / median / std
0.96 / — / 0.01 μs
numerical precisioneval_kernel_against_ref
gpu memoryextended_metrics.measure_memory
kernel launches / fusiontorch.profiler
fusion ratio (ref / kernel)
295.00x
SOL (speed-of-light)nsight_hardware_counters
roofline / occupancynsight_hardware_counters
warp stalls (top 5)nsight
reference PyTorch source (the prompt the agent sees)
"""
Robot–environment convex collision in **xy** (ground plane).
For an **axis-aligned robot footprint** vs. a **single convex obstacle triangle**, a
**separating-axis test (SAT)** gives the same **intersection predicate** as **GJK** in
2D, with a compact fully batched reference. A fused device kernel often implements
**GJK (+ EPA for penetration depth)** for arbitrary 3D polyhedra; this problem is
the canonical 2D footprint special case.
"""
import torch
import torch.nn as nn
# Default ModelNew is PyTorch-only. run_and_check: use check_kernel=False with backend=cuda
def _rect_corners_xy(robot_min: torch.Tensor, robot_max: torch.Tensor) -> torch.Tensor:
"""(B, 4, 2) corners of axis-aligned rectangles."""
x0, x1 = robot_min[:, 0], robot_max[:, 0]
y0, y1 = robot_min[:, 1], robot_max[:, 1]
return torch.stack(
(
torch.stack((x0, y0), dim=-1),
torch.stack((x1, y0), dim=-1),
torch.stack((x1, y1), dim=-1),
torch.stack((x0, y1), dim=-1),
),
dim=1,
)
def _project_min_max(points: torch.Tensor, axis: torch.Tensor) -> tuple[torch.Tensor, torch.Tensor]:
"""points (B, K, 2), axis (B, 2) or (2,) broadcast → (B,) min and max projections."""
if axis.dim() == 1:
axis = axis.view(1, 2).expand(points.shape[0], -1)
proj = (points * axis.unsqueeze(1)).sum(dim=-1)
return proj.min(dim=-1).values, proj.max(dim=-1).values
def _batched_sat_rect_triangle(
robot_min: torch.Tensor,
robot_max: torch.Tensor,
env_triangle: torch.Tensor,
eps: float,
) -> torch.Tensor:
"""
Separating-axis test: AABB vs triangle in 2D, fully vectorized over batch.
Returns (B,) float 1.0 if overlapping else 0.0.
"""
B = robot_min.shape[0]
corners = _rect_corners_xy(robot_min, robot_max)
hit = torch.ones(B, device=robot_min.device, dtype=robot_min.dtype)
# World axes (rectangle is AABB)
for ax in (
torch.tensor([1.0, 0.0], device=robot_min.device, dtype=robot_min.dtype),
torch.tensor([0.0, 1.0], device=robot_min.device, dtype=robot_min.dtype),
):
r0, r1 = _project_min_max(corners, ax)
t0, t1 = _project_min_max(env_triangle, ax)
separated = (r1 < t0 - eps) | (t1 < r0 - eps)
hit = hit * (~separated).to(hit.dtype)
# Triangle edge normals (per batch)
for i in range(3):
e = env_triangle[:, (i + 1) % 3] - env_triangle[:, i]
n = torch.stack((-e[:, 1], e[:, 0]), dim=-1)
n = n / (n.norm(dim=-1, keepdim=True) + 1e-8)
r0, r1 = _project_min_max(corners, n)
t0, t1 = _project_min_max(env_triangle, n)
separated = (r1 < t0 - eps) | (t1 < r0 - eps)
hit = hit * (~separated).to(hit.dtype)
return hit
class Model(nn.Module):
"""
**Batched robot footprint vs. environment triangle** intersection in **xy**.
Robot: axis-aligned box ``[robot_min, robot_max]`` per batch element.
Environment: one **triangle** ``(B, 3, 2)``.
Reference: **vectorized SAT** (5 axis tests). Fused-kernel target: same
predicate via **GJK/EPA** or specialized SAT SIMD for planning loops.
"""
def __init__(self, eps: float = 1e-6):
super().__init__()
self.eps = eps
def forward(
self,
robot_min: torch.Tensor,
robot_max: torch.Tensor,
env_triangle: torch.Tensor,
) -> torch.Tensor:
"""
Args:
robot_min: (B, 2), robot_max: (B, 2) with max >= min componentwise.
env_triangle: (B, 3, 2) obstacle vertices in order (any consistent winding).
Returns:
(B, 1) float in ``{0.0, 1.0}`` — **1** if the sets intersect.
"""
B = robot_min.shape[0]
assert robot_max.shape == (B, 2) and env_triangle.shape == (B, 3, 2)
h = _batched_sat_rect_triangle(robot_min, robot_max, env_triangle, self.eps)
return h.unsqueeze(-1)
class ModelNew(nn.Module):
"""KernelBench candidate entry point; default delegates to ``Model``."""
def __init__(self, eps: float = 1e-6):
super().__init__()
self._impl = Model(eps=eps)
def forward(
self,
robot_min: torch.Tensor,
robot_max: torch.Tensor,
env_triangle: torch.Tensor,
) -> torch.Tensor:
return self._impl(robot_min, robot_max, env_triangle)
batch_size = 16
def get_inputs():
bs = popcorn_pri.jitter_int(batch_size)
robot_min = torch.empty(bs, 2).uniform_(-1.0, 0.5)
robot_max = robot_min + torch.empty(bs, 2).uniform_(0.2, 0.8)
env_triangle = torch.empty(bs, 3, 2).uniform_(-0.8, 0.8)
return [robot_min, robot_max, env_triangle]
def get_init_inputs():
return []
turn 011.7s turn · 11.7s LLM · 0.0s tools
reasoning
**Optimizing code for GPU**
I'm focusing on producing code that needs to be optimized, possibly with a batch size set to 16. It's likely that the batch size will remain fixed. I need to implement a custom CUDA load inline, and the kernel should be simple. Knowing the GPU specifications will definitely help, so I should keep that in mind. I’m planning to use various tools to assist with this optimization.
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 137.7s turn · 37.7s LLM · 0.0s tools
call → compile_kernel
callcompile_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\n_cpp_src = r'''\n#include <torch/extension.h>\n\ntorch::Tensor sat_forward_cuda(torch::Tensor robot_min, torch::Tensor robot_max, torch::Tensor env_triangle, double eps);\n\ntorch::Tensor sat_forward(torch::Tensor robot_min, torch::Tensor robot_max, torch::Tensor env_triangle, double eps) {\n return sat_forward_cuda(robot_min, robot_max, env_triangle, eps);\n}\n'''\n\n_cuda_src = r'''\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cuda_runtime.h>\n\n__device__ __forceinline__ bool sep_axis(float rmin, float rmax, float tmin, float tmax, float eps) {\n return (rmax < tmin - eps) || (tmax < rmin - eps);\n}\n\n__global__ void sat_kernel(const float* __restrict__ rminp,\n const float* __restrict__ rmaxp,\n const float* __restrict__ trip,\n float* __restrict__ out,\n int B,\n float eps) {\n int b = blockIdx.x * blockDim.x + threadIdx.x;\n if (b >= B) return;\n\n float x0 = rminp[b*2 + 0];\n float y0 = rminp[b*2 + 1];\n float x1 = rmaxp[b*2 + 0];\n float y1 = rmaxp[b*2 + 1];\n\n const float* tri = trip + b*6;\n float ax = tri[0], ay = tri[1];\n float bx = tri[2], by = tri[3];\n float cx = tri[4], cy = tri[5];\n\n float tminx = fminf(ax, fminf(bx, cx));\n float tmaxx = fmaxf(ax, fmaxf(bx, cx));\n float tminy = fminf(ay, fminf(by, cy));\n float tmaxy = fmaxf(ay, fmaxf(by, cy));\n\n bool hit = true;\n if (sep_axis(x0, x1, tminx, tmaxx, eps) || sep_axis(y0, y1, tminy, tmaxy, eps)) {\n hit = false;\n } else {\n // Triangle edge normals, normalized as in the PyTorch reference.\n #pragma unroll\n for (int i = 0; i < 3; ++i) {\n float px, py, qx, qy;\n if (i == 0) { px = ax; py = ay; qx = bx; qy = by; }\n else if (i == 1) { px = bx; py = by; qx = cx; qy = cy; }\n else { px = cx; py = cy; qx = ax; qy = ay; }\n float ex = qx - px;\n float ey = qy - py;\n float inv = 1.0f / (sqrtf(ex*ex + ey*ey) + 1.0e-8f);\n float nx = -ey * inv;\n float ny = ex * inv;\n\n float rx_a = x0 * nx;\n float rx_b = x1 * nx;\n float ry_a = y0 * ny;\n float ry_b = y1 * ny;\n float rlo = fminf(rx_a, rx_b) + fminf(ry_a, ry_b);\n float rhi = fmaxf(rx_a, rx_b) + fmaxf(ry_a, ry_b);\n\n float pa = ax*nx + ay*ny;\n float pb = bx*nx + by*ny;\n float pc = cx*nx + cy*ny;\n float tlo = fminf(pa, fminf(pb, pc));\n float thi = fmaxf(pa, fmaxf(pb, pc));\n\n if (sep_axis(rlo, rhi, tlo, thi, eps)) { hit = false; break; }\n }\n }\n out[b] = hit ? 1.0f : 0.0f;\n}\n\ntorch::Tensor sat_forward_cuda(torch::Tensor robot_min, torch::Tensor robot_max, torch::Tensor env_triangle, double eps) {\n const int B = (int)robot_min.size(0);\n auto out = torch::empty({B, 1}, robot_min.options());\n if (B > 0) {\n const int threads = 32;\n const int blocks = (B + threads - 1) / threads;\n sat_kernel<<<blocks, threads, 0, at::cuda::getCurrentCUDAStream()>>>(\n robot_min.data_ptr<float>(), robot_max.data_ptr<float>(), env_triangle.data_ptr<float>(),\n out.data_ptr<float>(), B, (float)eps);\n }\n return out;\n}\n'''\n\n_ext = load_inline(\n name='sat_rect_tri_ext_v1',\n cpp_sources=[_cpp_src],\n cuda_sources=[_cuda_src],\n functions=['sat_forward'],\n extra_cflags=['-O3'],\n extra_cuda_cflags=['-O3', '--use_fast_math'],\n verbose=False,\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, eps: float = 1e-6):\n super().__init__()\n self.eps = float(eps)\n\n def forward(self, robot_min: torch.Tensor, robot_max: torch.Tensor, env_triangle: torch.Tensor) -> torch.Tensor:\n return _ext.sat_forward(robot_min, robot_max, env_triangle, self.eps)\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 28.6s turn · 8.6s LLM · 0.0s tools
call → run_correctness
callrun_correctness
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\n_cpp_src = r'''\n#include <torch/extension.h>\n\ntorch::Tensor sat_forward_cuda(torch::Tensor robot_min, torch::Tensor robot_max, torch::Tensor env_triangle, double eps);\n\ntorch::Tensor sat_forward(torch::Tensor robot_min, torch::Tensor robot_max, torch::Tensor env_triangle, double eps) {\n return sat_forward_cuda(robot_min, robot_max, env_triangle, eps);\n}\n'''\n\n_cuda_src = r'''\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cuda_runtime.h>\n\n__device__ __forceinline__ bool sep_axis(float rmin, float rmax, float tmin, float tmax, float eps) {\n return (rmax < tmin - eps) || (tmax < rmin - eps);\n}\n\n__global__ void sat_kernel(const float* __restrict__ rminp,\n const float* __restrict__ rmaxp,\n const float* __restrict__ trip,\n float* __restrict__ out,\n int B,\n float eps) {\n int b = blockIdx.x * blockDim.x + threadIdx.x;\n if (b >= B) return;\n\n float x0 = rminp[b*2 + 0];\n float y0 = rminp[b*2 + 1];\n float x1 = rmaxp[b*2 + 0];\n float y1 = rmaxp[b*2 + 1];\n\n const float* tri = trip + b*6;\n float ax = tri[0], ay = tri[1];\n float bx = tri[2], by = tri[3];\n float cx = tri[4], cy = tri[5];\n\n float tminx = fminf(ax, fminf(bx, cx));\n float tmaxx = fmaxf(ax, fmaxf(bx, cx));\n float tminy = fminf(ay, fminf(by, cy));\n float tmaxy = fmaxf(ay, fmaxf(by, cy));\n\n bool hit = true;\n if (sep_axis(x0, x1, tminx, tmaxx, eps) || sep_axis(y0, y1, tminy, tmaxy, eps)) {\n hit = false;\n } else {\n // Triangle edge normals, normalized as in the PyTorch reference.\n #pragma unroll\n for (int i = 0; i < 3; ++i) {\n float px, py, qx, qy;\n if (i == 0) { px = ax; py = ay; qx = bx; qy = by; }\n else if (i == 1) { px = bx; py = by; qx = cx; qy = cy; }\n else { px = cx; py = cy; qx = ax; qy = ay; }\n float ex = qx - px;\n float ey = qy - py;\n float inv = 1.0f / (sqrtf(ex*ex + ey*ey) + 1.0e-8f);\n float nx = -ey * inv;\n float ny = ex * inv;\n\n float rx_a = x0 * nx;\n float rx_b = x1 * nx;\n float ry_a = y0 * ny;\n float ry_b = y1 * ny;\n float rlo = fminf(rx_a, rx_b) + fminf(ry_a, ry_b);\n float rhi = fmaxf(rx_a, rx_b) + fmaxf(ry_a, ry_b);\n\n float pa = ax*nx + ay*ny;\n float pb = bx*nx + by*ny;\n float pc = cx*nx + cy*ny;\n float tlo = fminf(pa, fminf(pb, pc));\n float thi = fmaxf(pa, fmaxf(pb, pc));\n\n if (sep_axis(rlo, rhi, tlo, thi, eps)) { hit = false; break; }\n }\n }\n out[b] = hit ? 1.0f : 0.0f;\n}\n\ntorch::Tensor sat_forward_cuda(torch::Tensor robot_min, torch::Tensor robot_max, torch::Tensor env_triangle, double eps) {\n const int B = (int)robot_min.size(0);\n auto out = torch::empty({B, 1}, robot_min.options());\n if (B > 0) {\n const int threads = 32;\n const int blocks = (B + threads - 1) / threads;\n sat_kernel<<<blocks, threads, 0, at::cuda::getCurrentCUDAStream()>>>(\n robot_min.data_ptr<float>(), robot_max.data_ptr<float>(), env_triangle.data_ptr<float>(),\n out.data_ptr<float>(), B, (float)eps);\n }\n return out;\n}\n'''\n\n_ext = load_inline(\n name='sat_rect_tri_ext_v1',\n cpp_sources=[_cpp_src],\n cuda_sources=[_cuda_src],\n functions=['sat_forward'],\n extra_cflags=['-O3'],\n extra_cuda_cflags=['-O3', '--use_fast_math'],\n verbose=False,\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, eps: float = 1e-6):\n super().__init__()\n self.eps = float(eps)\n\n def forward(self, robot_min: torch.Tensor, robot_max: torch.Tensor, env_triangle: torch.Tensor) -> torch.Tensor:\n return _ext.sat_forward(robot_min, robot_max, env_triangle, self.eps)\n"
}
result ← run_correctness (ok)
okrun_correctness
run_correctness PASSED: (5 / 5) trials all matched the reference.
Numerical precision: max_abs_err=0.00e+00 mean_abs_err=0.00e+00 max_rel_err=0.00e+00
turn 312.6s turn · 12.6s LLM · 0.0s tools
reasoning
**Optimizing performance timing**
I need to optimize the timing of the current extension call overhead, which seems to be around 10 microseconds, with PyTorch operations clocking in around 100 microseconds. I wonder if avoiding the kernel launch could help. For a batch size of 16, this overhead might be significant. I should think about using `torch.empty` and consider the CUDA kernel options, but I need to address potential aliasing concerns. Caching output buffers could reduce overhead, but there are complexities to consider as well.
**Considering output buffer reuse**
I'm thinking about reusing the output buffer across trials. If the previous output is still in use, it could cause problems. If I store outputs in a list for later comparison, that reuse would break things. Maybe it's better to avoid this method altogether. Static checks might help flag lazy approaches, but reusing output could lead to questionable outcomes. It seems like steering clear of this approach is the way to go.
call → static_check
callstatic_check
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\n_cpp_src = r'''\n#include <torch/extension.h>\n\ntorch::Tensor sat_forward_cuda(torch::Tensor robot_min, torch::Tensor robot_max, torch::Tensor env_triangle, double eps);\n\ntorch::Tensor sat_forward(torch::Tensor robot_min, torch::Tensor robot_max, torch::Tensor env_triangle, double eps) {\n return sat_forward_cuda(robot_min, robot_max, env_triangle, eps);\n}\n'''\n\n_cuda_src = r'''\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cuda_runtime.h>\n\n__device__ __forceinline__ bool sep_axis(float rmin, float rmax, float tmin, float tmax, float eps) {\n return (rmax < tmin - eps) || (tmax < rmin - eps);\n}\n\n__global__ void sat_kernel(const float* __restrict__ rminp,\n const float* __restrict__ rmaxp,\n const float* __restrict__ trip,\n float* __restrict__ out,\n int B,\n float eps) {\n int b = blockIdx.x * blockDim.x + threadIdx.x;\n if (b >= B) return;\n\n float x0 = rminp[b*2 + 0];\n float y0 = rminp[b*2 + 1];\n float x1 = rmaxp[b*2 + 0];\n float y1 = rmaxp[b*2 + 1];\n\n const float* tri = trip + b*6;\n float ax = tri[0], ay = tri[1];\n float bx = tri[2], by = tri[3];\n float cx = tri[4], cy = tri[5];\n\n float tminx = fminf(ax, fminf(bx, cx));\n float tmaxx = fmaxf(ax, fmaxf(bx, cx));\n float tminy = fminf(ay, fminf(by, cy));\n float tmaxy = fmaxf(ay, fmaxf(by, cy));\n\n bool hit = true;\n if (sep_axis(x0, x1, tminx, tmaxx, eps) || sep_axis(y0, y1, tminy, tmaxy, eps)) {\n hit = false;\n } else {\n // Triangle edge normals, normalized as in the PyTorch reference.\n #pragma unroll\n for (int i = 0; i < 3; ++i) {\n float px, py, qx, qy;\n if (i == 0) { px = ax; py = ay; qx = bx; qy = by; }\n else if (i == 1) { px = bx; py = by; qx = cx; qy = cy; }\n else { px = cx; py = cy; qx = ax; qy = ay; }\n float ex = qx - px;\n float ey = qy - py;\n float inv = 1.0f / (sqrtf(ex*ex + ey*ey) + 1.0e-8f);\n float nx = -ey * inv;\n float ny = ex * inv;\n\n float rx_a = x0 * nx;\n float rx_b = x1 * nx;\n float ry_a = y0 * ny;\n float ry_b = y1 * ny;\n float rlo = fminf(rx_a, rx_b) + fminf(ry_a, ry_b);\n float rhi = fmaxf(rx_a, rx_b) + fmaxf(ry_a, ry_b);\n\n float pa = ax*nx + ay*ny;\n float pb = bx*nx + by*ny;\n float pc = cx*nx + cy*ny;\n float tlo = fminf(pa, fminf(pb, pc));\n float thi = fmaxf(pa, fmaxf(pb, pc));\n\n if (sep_axis(rlo, rhi, tlo, thi, eps)) { hit = false; break; }\n }\n }\n out[b] = hit ? 1.0f : 0.0f;\n}\n\ntorch::Tensor sat_forward_cuda(torch::Tensor robot_min, torch::Tensor robot_max, torch::Tensor env_triangle, double eps) {\n const int B = (int)robot_min.size(0);\n auto out = torch::empty({B, 1}, robot_min.options());\n if (B > 0) {\n const int threads = 32;\n const int blocks = (B + threads - 1) / threads;\n sat_kernel<<<blocks, threads, 0, at::cuda::getCurrentCUDAStream()>>>(\n robot_min.data_ptr<float>(), robot_max.data_ptr<float>(), env_triangle.data_ptr<float>(),\n out.data_ptr<float>(), B, (float)eps);\n }\n return out;\n}\n'''\n\n_ext = load_inline(\n name='sat_rect_tri_ext_v1',\n cpp_sources=[_cpp_src],\n cuda_sources=[_cuda_src],\n functions=['sat_forward'],\n extra_cflags=['-O3'],\n extra_cuda_cflags=['-O3', '--use_fast_math'],\n verbose=False,\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, eps: float = 1e-6):\n super().__init__()\n self.eps = float(eps)\n\n def forward(self, robot_min: torch.Tensor, robot_max: torch.Tensor, env_triangle: torch.Tensor) -> torch.Tensor:\n return _ext.sat_forward(robot_min, robot_max, env_triangle, self.eps)\n"
}
result ← static_check (ok)
okstatic_check
static_check PASSED: no violations or warnings detected.
turn 4 · FINAL · 1 tool call(s) · 13.0s turn · 13.0s LLM · 6,251 chars (click to expand)
turn 413.0s turn · 13.0s LLM · 0.0s tools · FINAL
reasoning
**Evaluating submission approach**
I’m considering whether I should submit now or improve it first. I need to know the runtime only requires the final submission without references. I should also think about potential issues, like the constructor arguments being the same and handling the B variable as dtype fp32. There are concerns about name collisions across compiled attempts and how using --use_fast_math might affect boundary calculations. I wonder about avoiding sqrt normalization since it may not change the separation. Could I combine aspects with a CPU input approach?
**Finalizing the submission**
I’m considering using __launch_bounds__(32), although it may not make much difference for B>32 blocks. I could vectorize each batch by thread, but it might lead to most of the warp being idle. I should ensure there’s only one block. The launch overhead for H100 is about 190us, which seems enough. I need to finalize a complete Python file since submitting is the next step, so I’ll make sure everything is included before I call submit.
call → submit_kernel
callsubmit_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\n_cpp_src = r'''\n#include <torch/extension.h>\n\ntorch::Tensor sat_forward_cuda(torch::Tensor robot_min, torch::Tensor robot_max, torch::Tensor env_triangle, double eps);\n\ntorch::Tensor sat_forward(torch::Tensor robot_min, torch::Tensor robot_max, torch::Tensor env_triangle, double eps) {\n return sat_forward_cuda(robot_min, robot_max, env_triangle, eps);\n}\n'''\n\n_cuda_src = r'''\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cuda_runtime.h>\n\n__device__ __forceinline__ bool sep_axis(float rmin, float rmax, float tmin, float tmax, float eps) {\n return (rmax < tmin - eps) || (tmax < rmin - eps);\n}\n\n__global__ void sat_kernel(const float* __restrict__ rminp,\n const float* __restrict__ rmaxp,\n const float* __restrict__ trip,\n float* __restrict__ out,\n int B,\n float eps) {\n int b = blockIdx.x * blockDim.x + threadIdx.x;\n if (b >= B) return;\n\n float x0 = rminp[b*2 + 0];\n float y0 = rminp[b*2 + 1];\n float x1 = rmaxp[b*2 + 0];\n float y1 = rmaxp[b*2 + 1];\n\n const float* tri = trip + b*6;\n float ax = tri[0], ay = tri[1];\n float bx = tri[2], by = tri[3];\n float cx = tri[4], cy = tri[5];\n\n float tminx = fminf(ax, fminf(bx, cx));\n float tmaxx = fmaxf(ax, fmaxf(bx, cx));\n float tminy = fminf(ay, fminf(by, cy));\n float tmaxy = fmaxf(ay, fmaxf(by, cy));\n\n bool hit = true;\n if (sep_axis(x0, x1, tminx, tmaxx, eps) || sep_axis(y0, y1, tminy, tmaxy, eps)) {\n hit = false;\n } else {\n #pragma unroll\n for (int i = 0; i < 3; ++i) {\n float px, py, qx, qy;\n if (i == 0) { px = ax; py = ay; qx = bx; qy = by; }\n else if (i == 1) { px = bx; py = by; qx = cx; qy = cy; }\n else { px = cx; py = cy; qx = ax; qy = ay; }\n float ex = qx - px;\n float ey = qy - py;\n float inv = 1.0f / (sqrtf(ex*ex + ey*ey) + 1.0e-8f);\n float nx = -ey * inv;\n float ny = ex * inv;\n\n float rx_a = x0 * nx;\n float rx_b = x1 * nx;\n float ry_a = y0 * ny;\n float ry_b = y1 * ny;\n float rlo = fminf(rx_a, rx_b) + fminf(ry_a, ry_b);\n float rhi = fmaxf(rx_a, rx_b) + fmaxf(ry_a, ry_b);\n\n float pa = ax*nx + ay*ny;\n float pb = bx*nx + by*ny;\n float pc = cx*nx + cy*ny;\n float tlo = fminf(pa, fminf(pb, pc));\n float thi = fmaxf(pa, fmaxf(pb, pc));\n\n if (sep_axis(rlo, rhi, tlo, thi, eps)) { hit = false; break; }\n }\n }\n out[b] = hit ? 1.0f : 0.0f;\n}\n\ntorch::Tensor sat_forward_cuda(torch::Tensor robot_min, torch::Tensor robot_max, torch::Tensor env_triangle, double eps) {\n const int B = (int)robot_min.size(0);\n auto out = torch::empty({B, 1}, robot_min.options());\n if (B > 0) {\n const int threads = 32;\n const int blocks = (B + threads - 1) / threads;\n sat_kernel<<<blocks, threads, 0, at::cuda::getCurrentCUDAStream()>>>(\n robot_min.data_ptr<float>(), robot_max.data_ptr<float>(), env_triangle.data_ptr<float>(),\n out.data_ptr<float>(), B, (float)eps);\n }\n return out;\n}\n'''\n\n_ext = load_inline(\n name='sat_rect_tri_ext_final',\n cpp_sources=[_cpp_src],\n cuda_sources=[_cuda_src],\n functions=['sat_forward'],\n extra_cflags=['-O3'],\n extra_cuda_cflags=['-O3', '--use_fast_math'],\n verbose=False,\n)\n\nclass ModelNew(nn.Module):\n def __init__(self, eps: float = 1e-6):\n super().__init__()\n self.eps = float(eps)\n\n def forward(self, robot_min: torch.Tensor, robot_max: torch.Tensor, env_triangle: torch.Tensor) -> torch.Tensor:\n return _ext.sat_forward(robot_min, robot_max, env_triangle, self.eps)\n"
}
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=0.00e+00 mean_abs_err=0.00e+00 max_rel_err=0.00e+00
Kernel launches: 1 (ref: 295, fusion ratio: 295.0)
SOL score: 0.0005
Flagged for excessive speedup — this submission may have been rejected by automated review.
final submitted kernel (114 lines, 3,926 chars) — full source, untruncated
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline
_cpp_src = r'''
#include <torch/extension.h>
torch::Tensor sat_forward_cuda(torch::Tensor robot_min, torch::Tensor robot_max, torch::Tensor env_triangle, double eps);
torch::Tensor sat_forward(torch::Tensor robot_min, torch::Tensor robot_max, torch::Tensor env_triangle, double eps) {
return sat_forward_cuda(robot_min, robot_max, env_triangle, eps);
}
'''
_cuda_src = r'''
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>
__device__ __forceinline__ bool sep_axis(float rmin, float rmax, float tmin, float tmax, float eps) {
return (rmax < tmin - eps) || (tmax < rmin - eps);
}
__global__ void sat_kernel(const float* __restrict__ rminp,
const float* __restrict__ rmaxp,
const float* __restrict__ trip,
float* __restrict__ out,
int B,
float eps) {
int b = blockIdx.x * blockDim.x + threadIdx.x;
if (b >= B) return;
float x0 = rminp[b*2 + 0];
float y0 = rminp[b*2 + 1];
float x1 = rmaxp[b*2 + 0];
float y1 = rmaxp[b*2 + 1];
const float* tri = trip + b*6;
float ax = tri[0], ay = tri[1];
float bx = tri[2], by = tri[3];
float cx = tri[4], cy = tri[5];
float tminx = fminf(ax, fminf(bx, cx));
float tmaxx = fmaxf(ax, fmaxf(bx, cx));
float tminy = fminf(ay, fminf(by, cy));
float tmaxy = fmaxf(ay, fmaxf(by, cy));
bool hit = true;
if (sep_axis(x0, x1, tminx, tmaxx, eps) || sep_axis(y0, y1, tminy, tmaxy, eps)) {
hit = false;
} else {
#pragma unroll
for (int i = 0; i < 3; ++i) {
float px, py, qx, qy;
if (i == 0) { px = ax; py = ay; qx = bx; qy = by; }
else if (i == 1) { px = bx; py = by; qx = cx; qy = cy; }
else { px = cx; py = cy; qx = ax; qy = ay; }
float ex = qx - px;
float ey = qy - py;
float inv = 1.0f / (sqrtf(ex*ex + ey*ey) + 1.0e-8f);
float nx = -ey * inv;
float ny = ex * inv;
float rx_a = x0 * nx;
float rx_b = x1 * nx;
float ry_a = y0 * ny;
float ry_b = y1 * ny;
float rlo = fminf(rx_a, rx_b) + fminf(ry_a, ry_b);
float rhi = fmaxf(rx_a, rx_b) + fmaxf(ry_a, ry_b);
float pa = ax*nx + ay*ny;
float pb = bx*nx + by*ny;
float pc = cx*nx + cy*ny;
float tlo = fminf(pa, fminf(pb, pc));
float thi = fmaxf(pa, fmaxf(pb, pc));
if (sep_axis(rlo, rhi, tlo, thi, eps)) { hit = false; break; }
}
}
out[b] = hit ? 1.0f : 0.0f;
}
torch::Tensor sat_forward_cuda(torch::Tensor robot_min, torch::Tensor robot_max, torch::Tensor env_triangle, double eps) {
const int B = (int)robot_min.size(0);
auto out = torch::empty({B, 1}, robot_min.options());
if (B > 0) {
const int threads = 32;
const int blocks = (B + threads - 1) / threads;
sat_kernel<<<blocks, threads, 0, at::cuda::getCurrentCUDAStream()>>>(
robot_min.data_ptr<float>(), robot_max.data_ptr<float>(), env_triangle.data_ptr<float>(),
out.data_ptr<float>(), B, (float)eps);
}
return out;
}
'''
_ext = load_inline(
name='sat_rect_tri_ext_final',
cpp_sources=[_cpp_src],
cuda_sources=[_cuda_src],
functions=['sat_forward'],
extra_cflags=['-O3'],
extra_cuda_cflags=['-O3', '--use_fast_math'],
verbose=False,
)
class ModelNew(nn.Module):
def __init__(self, eps: float = 1e-6):
super().__init__()
self.eps = float(eps)
def forward(self, robot_min: torch.Tensor, robot_max: torch.Tensor, env_triangle: torch.Tensor) -> torch.Tensor:
return _ext.sat_forward(robot_min, robot_max, env_triangle, self.eps)