correct
19_PairwiseDistanceMatrix.py
started
2026-05-11T19:17:58Z
timingkernelbench.timing
kernel mean / median / std
0.01 / — / 4.45e-04 μs
ref mean / median / std
0.08 / — / 4.64e-04 μs
numerical precisioneval_kernel_against_ref
gpu memoryextended_metrics.measure_memory
kernel launches / fusiontorch.profiler
fusion ratio (ref / kernel)
12.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)
import torch
import torch.nn as nn
class Model(nn.Module):
"""
Compute all-pairs Euclidean distance matrix from a 3-D point cloud.
A core primitive in molecular dynamics simulations, docking, and
distance-geometry-based structure prediction.
"""
def __init__(self):
super().__init__()
def forward(self, coords: torch.Tensor) -> torch.Tensor:
"""
Args:
coords: (B, N, 3) – 3-D coordinates of N atoms / residues
Returns:
dist: (B, N, N) – pairwise Euclidean distance matrix
"""
diff = coords.unsqueeze(2) - coords.unsqueeze(1) # (B, N, N, 3)
dist = torch.sqrt((diff ** 2).sum(dim=-1) + 1e-8)
return dist
num_atoms = 512
batch_size = 8
def get_inputs():
p = popcorn_pri
return [torch.randn(p.jitter_int(batch_size), p.jitter_int(num_atoms), 3)]
def get_init_inputs():
return []
turn 0307.8s turn · 11.7s LLM · 296.0s tools · tokens in=2,031 · out=969 · total=3,000
reasoning
Let me understand the task:
1. I need to replace the forward computation of a PyTorch model that computes all-pairs Euclidean distance matrix from a 3-D point cloud
2. The input is `(B, N, 3)` where B is batch size, N is number of atoms/residues
3. The output is `(B, N, N)` - pairwise Euclidean distance matrix
4. The computation is:
- `diff = coords.unsqueeze(2) - coords.unsqueeze(1)` gives `(B, N, N, 3)`
- `dist = sqrt((diff ** 2).sum(dim=-1) + 1e-8)`
This is a classic all-pairs distance computation. Let me think about how to optimize this with a CUDA kernel.
For the kernel:
- Each thread could compute one element of the output distance matrix
- We need to compute the squared differences for each of the 3 coordinates
- Sum them up and take the square root
Let me first get the GPU specs to understand what I'm working with.
The naive approach would be:
- Grid: (B * N * N) threads (one per output element)
- Each thread computes sqrt((x_i - x_j)^2 + (y_i - y_j)^2 + (z_i - z_j)^2)
But this might be memory-bound since we're accessing coords multiple times. Let me think about optimizations:
1. **Shared memory tiling**: Load tiles of coordinates into shared memory to reduce global memory access
2. **Block structure**: Have each block handle a tile of the output matrix
3. **Coalesced reads**: Make sure threads read coordinates in a coalesced manner
Let me start with a straightforward implementation and then profile it.
Actually, for distance matrix computation, a common optimization is:
- Use shared memory to load a tile of points
- Each block computes a tile of the output matrix
- Threads within a block cooperate to load points into shared memory
Let me write the kernel:
```python
import torch
from torch.utils.cpp_extension import load_inline
cuda_source = """
__global__ void distance_matrix_kernel(const float* __restrict__ coords,
float* __restrict__ dist,
int B, int N) {
// Each thread computes one element of the output
int b = blockIdx.z;
int i = blockIdx.y * blockDim.y + threadIdx.y;
int j = blockIdx.x * blockDim.x + threadIdx.x;
if (b < B && i < N && j < N) {
float dx = coords[b * N * 3 + i * 3 + 0] - coords[b * N * 3 + j * 3 + 0];
float dy = coords[b * N * 3 + i * 3 + 1] - coords[b * N * 3 + j * 3 + 1];
float dz = coords[b * N * 3 + i * 3 + 2] - coords[b * N * 3 + j * 3 + 2];
dist[b * N * N + i * N + j] = sqrtf(dx * dx + dy * dy + dz * dz + 1e-8f);
}
}
torch::Tensor distance_matrix_cuda(torch::Tensor coords) {
int B = coords.size(0);
int N = coords.size(1);
auto dist = torch::empty({B, N, N}, coords.options());
dim3 threads(16, 16);
dim3 blocks((N + threads.x - 1) / threads.x,
(N + threads.y - 1) / threads.y,
B);
distance_matrix_kernel<<<blocks, threads>>>(
coords.data_ptr<float>(),
dist.data_ptr<float>(),
B, N
);
return dist;
}
"""
cpp_source = "torch::Tensor distance_matrix_cuda(torch::Tensor coords);"
module = load_inline(
name='distance_matrix',
cpp_sources=cpp_source,
cuda_sources=cuda_source,
functions=['distance_matrix_cuda'],
extra_cuda_cflags=['-O3'],
verbose=True
)
class ModelNew(nn.Module):
def __init__(self):
super().__init__()
def forward(self, coords: torch.Tensor) -> torch.Tensor:
return module.distance_matrix_cuda(coords)
```
This is a basic approach with each thread handling one output element. The real bottleneck is repeated coordinate access, so I should use shared memory tiling for better performance. Let me get the GPU specs first.
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) · 291.0s turn · 10.5s LLM · tokens in=2,279 · out=993 · total=3,272 · 7,574 chars (click to expand)
turn 1291.0s turn · 10.5s LLM · 280.5s tools · tokens in=2,279 · out=993 · total=3,272
reasoning
Now let me write my first kernel implementation. This computes pairwise Euclidean distances for a 3D point cloud. The output is (B, N, N) where each element is the distance between two points.
Key observations:
1. The computation is symmetric (dist[i,j] == dist[j,i])
2. We can compute distances directly: sqrt((x1-x2)^2 + (y1-y2)^2 + (z1-z2)^2)
3. Output size is B*N*N = 8*512*512 = 2M elements
Let me write a kernel that:
- Uses shared memory to cache coords for faster access
- Has each thread compute one distance (or multiple if needed for occupancy)
- Avoids redundant computation by only computing half the matrix and mirroring
Actually, for simplicity, let me first write a straightforward kernel that computes all pairs, then optimize.
assistantLet me start by writing a basic CUDA kernel for pairwise Euclidean distances:
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 = \"\"\"\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <cmath>\n\n// Basic pairwise distance kernel\n__global__ void pairwise_distance_kernel(\n const float* __restrict__ coords,\n float* __restrict__ dist,\n const int B,\n const int N\n) {\n // coords: (B, N, 3), dist: (B, N, N)\n int b = blockIdx.z;\n int i = blockIdx.y * blockDim.y + threadIdx.y;\n int j = blockIdx.x * blockDim.x + threadIdx.x;\n \n if (b < B && i < N && j < N) {\n // Load coordinates\n float x1 = coords[b * N * 3 + i * 3 + 0];\n float y1 = coords[b * N * 3 + i * 3 + 1];\n float z1 = coords[b * N * 3 + i * 3 + 2];\n \n float x2 = coords[b * N * 3 + j * 3 + 0];\n float y2 = coords[b * N * 3 + j * 3 + 1];\n float z2 = coords[b * N * 3 + j * 3 + 2];\n \n // Compute squared distance\n float dx = x1 - x2;\n float dy = y1 - y2;\n float dz = z1 - z2;\n float dist_sq = dx * dx + dy * dy + dz * dz;\n \n // Store distance (with epsilon to match PyTorch)\n dist[b * N * N + i * N + j] = sqrtf(dist_sq + 1e-8f);\n }\n}\n\nvoid launch_pairwise_distance(const float* coords, float* dist, int B, int N, cudaStream_t stream) {\n dim3 threadsPerBlock(16, 16);\n dim3 numBlocks((N + threadsPerBlock.x - 1) / threadsPerBlock.x,\n (N + threadsPerBlock.y - 1) / threadsPerBlock.y,\n B);\n pairwise_distance_kernel<<<numBlocks, threadsPerBlock, 0, stream>>>(coords, dist, B, N);\n}\n\"\"\"\n\ncpp_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\nvoid launch_pairwise_distance(const float* coords, float* dist, int B, int N, cudaStream_t stream);\n\ntorch::Tensor pairwise_distance_cuda(torch::Tensor coords) {\n TORCH_CHECK(coords.dim() == 3, \"coords must be 3D\");\n TORCH_CHECK(coords.size(2) == 3, \"coords last dimension must be 3\");\n \n int B = coords.size(0);\n int N = coords.size(1);\n \n auto dist = torch::empty({B, N, N}, coords.options());\n \n const float* coords_ptr = coords.data_ptr<float>();\n float* dist_ptr = dist.data_ptr<float>();\n \n cudaStream_t stream = at::cuda::getCurrentCUDAStream();\n launch_pairwise_distance(coords_ptr, dist_ptr, B, N, stream);\n \n return dist;\n}\n\"\"\"\n\nmodule = load_inline(\n name='pairwise_distance_ext',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['pairwise_distance_cuda'],\n extra_cuda_cflags=['-O3'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n\n def forward(self, coords: torch.Tensor) -> torch.Tensor:\n return module.pairwise_distance_cuda(coords.contiguous())"
}
result ← compile_kernel (fail)
failcompile_kernel
compile_kernel FAILED: RuntimeError.
RuntimeError: Error building extension 'pairwise_distance_ext': [1/3] /opt/conda/bin/x86_64-conda-linux-gnu-c++ -MMD -MF main.o.d -DTORCH_EXTENSION_NAME=pairwise_distance_ext -DTORCH_API_INCLUDE_EXTENSION_H -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /usr/local/cuda/include -isystem /opt/conda/include/python3.10 -fPIC -std=c++17 -c /scratch/abaner/KernelBench/runs/pop_l123_all_gpt/popcorn/FW-GLM-5-1/level_1_problem_19_cache/k_b9fcecf41bde/pairwise_distance_ext/main.cpp -o main.o
FAILED: [code=1] main.o
/opt/conda/bin/x86_64-conda-linux-gnu-c++ -MMD -MF main.o.d -DTORCH_EXTENSION_NAME=pairwise_distance_ext -DTORCH_API_INCLUDE_EXTENSION_H -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /usr/local/cuda/include -isystem /opt/conda/include/python3.10 -fPIC -std=c++17 -c /scratch/abaner/KernelBench/runs/pop_l123_all_gpt/popcorn/FW-GLM-5-1/level_1_problem_19_cache/k_b9fcecf41bde/pairwise_distance_ext/main.cpp -o main.o
/scratch/abaner/KernelBench/runs/pop_l123_all_gpt/popcorn/FW-GLM-5-1/level_1_problem_19_cache/k_b9fcecf41bde/pairwise_distance_ext/main.cpp: In function 'at::Tensor pairwise_distance_cuda(at::Tensor)':
/scratch/abaner/KernelBench/runs/pop_l123_all_gpt/popcorn/FW-GLM-5-1/level_1_problem_19_cache/k_b9fcecf41bde/pairwise_distance_ext/main.cpp:21:37: error: 'getCurrentCUDAStream' is not a member of 'at::cuda'
21 | cudaStream_t stream = at::cuda::getCurrentCUDAStream();
| ^~~~~~~~~~~~~~~~~~~~
[2/3] /usr/local/cuda/bin/nvcc -MD -MF cuda.cuda.o.d -ccbin /opt/conda/bin/x86_64-conda-linux-gnu-cc -DTORCH_EXTENSION_NAME=pairwise_distance_ext -DTORCH_API_INCLUDE_EXTENSION_H -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /usr/local/cuda/include -isystem /opt/conda/include/python3.10 -D__CUDA_NO_HALF_OPERATORS__ -D__CUDA_NO_HALF_CONVERSIONS__ -D__CUDA_NO_BFLOAT16_CONVERSIONS__ -D__CUDA_NO_HALF2_OPERATORS__ --expt-relaxed-constexpr -gencode=arch=compute_90,code=compute_90 -gencode=arch=compute_90,code=sm_90 --compiler-options '-fPIC' -O3 -std=c++17 -c /scratch/abaner/KernelBench/runs/pop_l123_all_gpt/popcorn/FW-GLM-5-1/level_1_problem_19_cache/k_b9fcecf41bde/pairwise_distance_ext/cuda.cu -o cuda.cuda.o
nvcc warning : incompatible redefinition for option 'compiler-bindir', the last value of this option was used
ninja: build stopped: subcommand failed.
turn 2 · 1 tool call(s) · 267.4s turn · 7.6s LLM · tokens in=3,980 · out=830 · total=4,810 · 6,908 chars (click to expand)
turn 2267.4s turn · 7.6s LLM · 259.9s tools · tokens in=3,980 · out=830 · total=4,810
reasoning
The error is about `at::cuda::getCurrentCUDAStream()` not being a member. Let me fix this using the correct API:
assistantLet me fix the CUDA stream API:
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 = \"\"\"\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <cmath>\n\n// Basic pairwise distance kernel\n__global__ void pairwise_distance_kernel(\n const float* __restrict__ coords,\n float* __restrict__ dist,\n const int B,\n const int N\n) {\n // coords: (B, N, 3), dist: (B, N, N)\n int b = blockIdx.z;\n int i = blockIdx.y * blockDim.y + threadIdx.y;\n int j = blockIdx.x * blockDim.x + threadIdx.x;\n \n if (b < B && i < N && j < N) {\n // Load coordinates\n float x1 = coords[b * N * 3 + i * 3 + 0];\n float y1 = coords[b * N * 3 + i * 3 + 1];\n float z1 = coords[b * N * 3 + i * 3 + 2];\n \n float x2 = coords[b * N * 3 + j * 3 + 0];\n float y2 = coords[b * N * 3 + j * 3 + 1];\n float z2 = coords[b * N * 3 + j * 3 + 2];\n \n // Compute squared distance\n float dx = x1 - x2;\n float dy = y1 - y2;\n float dz = z1 - z2;\n float dist_sq = dx * dx + dy * dy + dz * dz;\n \n // Store distance (with epsilon to match PyTorch)\n dist[b * N * N + i * N + j] = sqrtf(dist_sq + 1e-8f);\n }\n}\n\nvoid launch_pairwise_distance(const float* coords, float* dist, int B, int N, cudaStream_t stream) {\n dim3 threadsPerBlock(16, 16);\n dim3 numBlocks((N + threadsPerBlock.x - 1) / threadsPerBlock.x,\n (N + threadsPerBlock.y - 1) / threadsPerBlock.y,\n B);\n pairwise_distance_kernel<<<numBlocks, threadsPerBlock, 0, stream>>>(coords, dist, B, N);\n}\n\"\"\"\n\ncpp_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\nvoid launch_pairwise_distance(const float* coords, float* dist, int B, int N, cudaStream_t stream);\n\ntorch::Tensor pairwise_distance_cuda(torch::Tensor coords) {\n TORCH_CHECK(coords.dim() == 3, \"coords must be 3D\");\n TORCH_CHECK(coords.size(2) == 3, \"coords last dimension must be 3\");\n \n int B = coords.size(0);\n int N = coords.size(1);\n \n auto dist = torch::empty({B, N, N}, coords.options());\n \n const float* coords_ptr = coords.data_ptr<float>();\n float* dist_ptr = dist.data_ptr<float>();\n \n cudaStream_t stream = c10::cuda::getCurrentCUDAStream();\n launch_pairwise_distance(coords_ptr, dist_ptr, B, N, stream);\n \n return dist;\n}\n\"\"\"\n\nmodule = load_inline(\n name='pairwise_distance_ext',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['pairwise_distance_cuda'],\n extra_cuda_cflags=['-O3'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n\n def forward(self, coords: torch.Tensor) -> torch.Tensor:\n return module.pairwise_distance_cuda(coords.contiguous())"
}
result ← compile_kernel (fail)
failcompile_kernel
compile_kernel FAILED: RuntimeError.
RuntimeError: Error building extension 'pairwise_distance_ext_v1': [1/3] /opt/conda/bin/x86_64-conda-linux-gnu-c++ -MMD -MF main.o.d -DTORCH_EXTENSION_NAME=pairwise_distance_ext_v1 -DTORCH_API_INCLUDE_EXTENSION_H -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /usr/local/cuda/include -isystem /opt/conda/include/python3.10 -fPIC -std=c++17 -c /scratch/abaner/KernelBench/runs/pop_l123_all_gpt/popcorn/FW-GLM-5-1/level_1_problem_19_cache/k_af96e3cb9878/pairwise_distance_ext/main.cpp -o main.o
FAILED: [code=1] main.o
/opt/conda/bin/x86_64-conda-linux-gnu-c++ -MMD -MF main.o.d -DTORCH_EXTENSION_NAME=pairwise_distance_ext_v1 -DTORCH_API_INCLUDE_EXTENSION_H -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /usr/local/cuda/include -isystem /opt/conda/include/python3.10 -fPIC -std=c++17 -c /scratch/abaner/KernelBench/runs/pop_l123_all_gpt/popcorn/FW-GLM-5-1/level_1_problem_19_cache/k_af96e3cb9878/pairwise_distance_ext/main.cpp -o main.o
/scratch/abaner/KernelBench/runs/pop_l123_all_gpt/popcorn/FW-GLM-5-1/level_1_problem_19_cache/k_af96e3cb9878/pairwise_distance_ext/main.cpp: In function 'at::Tensor pairwise_distance_cuda(at::Tensor)':
/scratch/abaner/KernelBench/runs/pop_l123_all_gpt/popcorn/FW-GLM-5-1/level_1_problem_19_cache/k_af96e3cb9878/pairwise_distance_ext/main.cpp:21:38: error: 'getCurrentCUDAStream' is not a member of 'c10::cuda'
21 | cudaStream_t stream = c10::cuda::getCurrentCUDAStream();
| ^~~~~~~~~~~~~~~~~~~~
[2/3] /usr/local/cuda/bin/nvcc -MD -MF cuda.cuda.o.d -ccbin /opt/conda/bin/x86_64-conda-linux-gnu-cc -DTORCH_EXTENSION_NAME=pairwise_distance_ext_v1 -DTORCH_API_INCLUDE_EXTENSION_H -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include -isystem /scratch/abaner/KernelBench/.venv/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /usr/local/cuda/include -isystem /opt/conda/include/python3.10 -D__CUDA_NO_HALF_OPERATORS__ -D__CUDA_NO_HALF_CONVERSIONS__ -D__CUDA_NO_BFLOAT16_CONVERSIONS__ -D__CUDA_NO_HALF2_OPERATORS__ --expt-relaxed-constexpr -gencode=arch=compute_90,code=compute_90 -gencode=arch=compute_90,code=sm_90 --compiler-options '-fPIC' -O3 -std=c++17 -c /scratch/abaner/KernelBench/runs/pop_l123_all_gpt/popcorn/FW-GLM-5-1/level_1_problem_19_cache/k_af96e3cb9878/pairwise_distance_ext/cuda.cu -o cuda.cuda.o
nvcc warning : incompatible redefinition for option 'compiler-bindir', the last value of this option was used
ninja: build stopped: subcommand failed.
turn 3278.0s turn · 8.4s LLM · 269.6s tools · tokens in=5,690 · out=845 · total=6,535
reasoning
Let me try using the at::cuda::getCurrentCUDAStream() via at::cuda::getCurrentCUDAStream().stream() or using at::cuda::getReservedStream() or just not using a stream at all, or use 0 (default stream).
Let me try a simpler approach without specifying the stream explicitly, or using the getDefaultCUDAStream() method.
assistantLet me try using the correct CUDA stream API:
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 = \"\"\"\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <cmath>\n\n// Basic pairwise distance kernel\n__global__ void pairwise_distance_kernel(\n const float* __restrict__ coords,\n float* __restrict__ dist,\n const int B,\n const int N\n) {\n // coords: (B, N, 3), dist: (B, N, N)\n int b = blockIdx.z;\n int i = blockIdx.y * blockDim.y + threadIdx.y;\n int j = blockIdx.x * blockDim.x + threadIdx.x;\n \n if (b < B && i < N && j < N) {\n // Load coordinates\n float x1 = coords[b * N * 3 + i * 3 + 0];\n float y1 = coords[b * N * 3 + i * 3 + 1];\n float z1 = coords[b * N * 3 + i * 3 + 2];\n \n float x2 = coords[b * N * 3 + j * 3 + 0];\n float y2 = coords[b * N * 3 + j * 3 + 1];\n float z2 = coords[b * N * 3 + j * 3 + 2];\n \n // Compute squared distance\n float dx = x1 - x2;\n float dy = y1 - y2;\n float dz = z1 - z2;\n float dist_sq = dx * dx + dy * dy + dz * dz;\n \n // Store distance (with epsilon to match PyTorch)\n dist[b * N * N + i * N + j] = sqrtf(dist_sq + 1e-8f);\n }\n}\n\nvoid launch_pairwise_distance(const float* coords, float* dist, int B, int N) {\n dim3 threadsPerBlock(16, 16);\n dim3 numBlocks((N + threadsPerBlock.x - 1) / threadsPerBlock.x,\n (N + threadsPerBlock.y - 1) / threadsPerBlock.y,\n B);\n pairwise_distance_kernel<<<numBlocks, threadsPerBlock>>>(coords, dist, B, N);\n}\n\"\"\"\n\ncpp_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\nvoid launch_pairwise_distance(const float* coords, float* dist, int B, int N);\n\ntorch::Tensor pairwise_distance_cuda(torch::Tensor coords) {\n TORCH_CHECK(coords.dim() == 3, \"coords must be 3D\");\n TORCH_CHECK(coords.size(2) == 3, \"coords last dimension must be 3\");\n \n int B = coords.size(0);\n int N = coords.size(1);\n \n auto dist = torch::empty({B, N, N}, coords.options());\n \n const float* coords_ptr = coords.data_ptr<float>();\n float* dist_ptr = dist.data_ptr<float>();\n \n launch_pairwise_distance(coords_ptr, dist_ptr, B, N);\n \n return dist;\n}\n\"\"\"\n\nmodule = load_inline(\n name='pairwise_distance_ext',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['pairwise_distance_cuda'],\n extra_cuda_cflags=['-O3'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n\n def forward(self, coords: torch.Tensor) -> torch.Tensor:\n return module.pairwise_distance_cuda(coords.contiguous())"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 4320.1s turn · 6.7s LLM · 313.4s tools · tokens in=6,476 · out=770 · total=7,246
reasoning
Now let me check correctness.
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 = \"\"\"\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <cmath>\n\n// Basic pairwise distance kernel\n__global__ void pairwise_distance_kernel(\n const float* __restrict__ coords,\n float* __restrict__ dist,\n const int B,\n const int N\n) {\n // coords: (B, N, 3), dist: (B, N, N)\n int b = blockIdx.z;\n int i = blockIdx.y * blockDim.y + threadIdx.y;\n int j = blockIdx.x * blockDim.x + threadIdx.x;\n \n if (b < B && i < N && j < N) {\n // Load coordinates\n float x1 = coords[b * N * 3 + i * 3 + 0];\n float y1 = coords[b * N * 3 + i * 3 + 1];\n float z1 = coords[b * N * 3 + i * 3 + 2];\n \n float x2 = coords[b * N * 3 + j * 3 + 0];\n float y2 = coords[b * N * 3 + j * 3 + 1];\n float z2 = coords[b * N * 3 + j * 3 + 2];\n \n // Compute squared distance\n float dx = x1 - x2;\n float dy = y1 - y2;\n float dz = z1 - z2;\n float dist_sq = dx * dx + dy * dy + dz * dz;\n \n // Store distance (with epsilon to match PyTorch)\n dist[b * N * N + i * N + j] = sqrtf(dist_sq + 1e-8f);\n }\n}\n\nvoid launch_pairwise_distance(const float* coords, float* dist, int B, int N) {\n dim3 threadsPerBlock(16, 16);\n dim3 numBlocks((N + threadsPerBlock.x - 1) / threadsPerBlock.x,\n (N + threadsPerBlock.y - 1) / threadsPerBlock.y,\n B);\n pairwise_distance_kernel<<<numBlocks, threadsPerBlock>>>(coords, dist, B, N);\n}\n\"\"\"\n\ncpp_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\nvoid launch_pairwise_distance(const float* coords, float* dist, int B, int N);\n\ntorch::Tensor pairwise_distance_cuda(torch::Tensor coords) {\n TORCH_CHECK(coords.dim() == 3, \"coords must be 3D\");\n TORCH_CHECK(coords.size(2) == 3, \"coords last dimension must be 3\");\n \n int B = coords.size(0);\n int N = coords.size(1);\n \n auto dist = torch::empty({B, N, N}, coords.options());\n \n const float* coords_ptr = coords.data_ptr<float>();\n float* dist_ptr = dist.data_ptr<float>();\n \n launch_pairwise_distance(coords_ptr, dist_ptr, B, N);\n \n return dist;\n}\n\"\"\"\n\nmodule = load_inline(\n name='pairwise_distance_ext',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['pairwise_distance_cuda'],\n extra_cuda_cflags=['-O3'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n\n def forward(self, coords: torch.Tensor) -> torch.Tensor:\n return module.pairwise_distance_cuda(coords.contiguous())"
}
result ← run_correctness (ok)
okrun_correctness
run_correctness PASSED: (3 / 3) trials all matched the reference.
Numerical precision: max_abs_err=4.77e-07 mean_abs_err=2.82e-08 max_rel_err=1.65e-07
turn 5292.2s turn · 7.2s LLM · 285.0s tools · tokens in=7,297 · out=797 · total=8,094
reasoning
The kernel is correct. Now let me profile it to see how it performs and what optimizations I can make.
assistantNow let me profile the kernel to understand its performance characteristics:
call → profile_kernel
callprofile_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = \"\"\"\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <cmath>\n\n// Basic pairwise distance kernel\n__global__ void pairwise_distance_kernel(\n const float* __restrict__ coords,\n float* __restrict__ dist,\n const int B,\n const int N\n) {\n // coords: (B, N, 3), dist: (B, N, N)\n int b = blockIdx.z;\n int i = blockIdx.y * blockDim.y + threadIdx.y;\n int j = blockIdx.x * blockDim.x + threadIdx.x;\n \n if (b < B && i < N && j < N) {\n // Load coordinates\n float x1 = coords[b * N * 3 + i * 3 + 0];\n float y1 = coords[b * N * 3 + i * 3 + 1];\n float z1 = coords[b * N * 3 + i * 3 + 2];\n \n float x2 = coords[b * N * 3 + j * 3 + 0];\n float y2 = coords[b * N * 3 + j * 3 + 1];\n float z2 = coords[b * N * 3 + j * 3 + 2];\n \n // Compute squared distance\n float dx = x1 - x2;\n float dy = y1 - y2;\n float dz = z1 - z2;\n float dist_sq = dx * dx + dy * dy + dz * dz;\n \n // Store distance (with epsilon to match PyTorch)\n dist[b * N * N + i * N + j] = sqrtf(dist_sq + 1e-8f);\n }\n}\n\nvoid launch_pairwise_distance(const float* coords, float* dist, int B, int N) {\n dim3 threadsPerBlock(16, 16);\n dim3 numBlocks((N + threadsPerBlock.x - 1) / threadsPerBlock.x,\n (N + threadsPerBlock.y - 1) / threadsPerBlock.y,\n B);\n pairwise_distance_kernel<<<numBlocks, threadsPerBlock>>>(coords, dist, B, N);\n}\n\"\"\"\n\ncpp_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\nvoid launch_pairwise_distance(const float* coords, float* dist, int B, int N);\n\ntorch::Tensor pairwise_distance_cuda(torch::Tensor coords) {\n TORCH_CHECK(coords.dim() == 3, \"coords must be 3D\");\n TORCH_CHECK(coords.size(2) == 3, \"coords last dimension must be 3\");\n \n int B = coords.size(0);\n int N = coords.size(1);\n \n auto dist = torch::empty({B, N, N}, coords.options());\n \n const float* coords_ptr = coords.data_ptr<float>();\n float* dist_ptr = dist.data_ptr<float>();\n \n launch_pairwise_distance(coords_ptr, dist_ptr, B, N);\n \n return dist;\n}\n\"\"\"\n\nmodule = load_inline(\n name='pairwise_distance_ext',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['pairwise_distance_cuda'],\n extra_cuda_cflags=['-O3'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n\n def forward(self, coords: torch.Tensor) -> torch.Tensor:\n return module.pairwise_distance_cuda(coords.contiguous())"
}
result ← profile_kernel (ok)
okprofile_kernel
profile_kernel PASSED: profiling complete.
=== Kernel Profile Summary ===
--- Kernel Breakdown (by GPU time) ---
#1: aten::empty 2731.0μs (100%) calls=1
--- Timing ---
GPU time: 8.58 μs
--- Memory Bandwidth ---
DRAM BW: 6.2/3350 GB/s (0.2%)
read=6.2 write=0.0 GB/s
Hit rates: L2=97.9%
--- Compute Throughput ---
Dominant pipe: FP32 util=21.9%
Pipe util: FMA=21.9%, Tensor=0.0%
--- Roofline Analysis ---
Ridge point (FP32): 20.00 FLOPs/byte
Classification: COMPUTE-BOUND
--- Occupancy ---
Achieved: 71.0%
Eligible warps/cycle: 1.28
Launch: regs/thread=16, smem/blk=1024B, block=256, grid=8192
--- Memory Access Quality ---
Global ld coalescing: 3.67 sectors/req (3.7x wasted BW)
--- Warp Stall Reasons (% of issue cycles) ---
long_scoreboard 874.6% ####################
short_scoreboard 292.1% ####################
mio_throttle 241.8% ####################
--- Optimization Hints ---
-> 875% stalls: global memory latency (waiting for DRAM/L2). + sectors/req=3.7 -> uncoalesced global loads, restructure access pattern or use __ldg Try shared-memory tiling, prefetching, or reducing working set.
-> Global load coalescing=3.7x (ideal=1.0) -> 4x wasted DRAM BW. Make per-warp accesses contiguous.
-> Compute-bound but NOT using tensor cores. If applicable, use FP16/BF16 with torch.matmul or WMMA intrinsics.
turn 6356.5s turn · 12.7s LLM · 343.8s tools · tokens in=8,515 · out=1,201 · total=9,716
reasoning
The profiling shows:
1. Poor memory coalescing (3.67 sectors/req, should be 1.0)
2. High warp stalls due to memory latency (long_scoreboard and short_scoreboard)
3. Low occupancy (71%)
4. The kernel is compute-bound but not using tensor cores
The main issue is uncoalesced memory access. The current kernel accesses coords[b*N*3 + i*3 + k] and coords[b*N*3 + j*3 + k], which isn't coalesced well because threads in a warp are accessing different i values but reading from scattered j values.
Let me design a tiled approach using shared memory:
- Each block processes one batch and one tile of i points
- Load the i tile into shared memory
- Iterate over j in tiles and compute all (i, j) pairs
- This will improve coalescing and reuse
But since the output is symmetric, I could also just compute the upper triangle and mirror it. But let me focus on a tiled approach first.
Actually, let me think about a simpler optimization first - using a 2D shared memory tile for the coordinates. For each block:
- Process one batch element
- Each block handles a 16x16 tile of the output matrix
- Load 16 points for "row" indices into shared memory
- Load 16 points for "column" indices from global memory directly (coalesced if done right)
Actually, let me try a different approach - have each thread process one row of the output. This way:
- Each thread loads its own point (xi, yi, zi) once
- Each thread iterates over all j points, loading them in a coalesced manner
- This is more memory efficient because we reuse the i coordinate
Let me implement this:
assistantThe profiling shows poor memory coalescing (3.67 sectors/req) causing memory stalls. Let me optimize with a better access pattern - having each thread process a row of the output matrix, allowing coalesced reads of the j-coordinates:
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 = \"\"\"\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <cmath>\n\n// Optimized kernel: each thread processes one row of the distance matrix\n// This allows coalesced reads for the second coordinate\n__global__ void pairwise_distance_kernel(\n const float* __restrict__ coords,\n float* __restrict__ dist,\n const int B,\n const int N\n) {\n // Each thread handles one row (i) of the output\n int b = blockIdx.y;\n int i = blockIdx.x * blockDim.x + threadIdx.x;\n \n if (b < B && i < N) {\n // Load coordinates for point i (reused N times)\n float x1 = coords[b * N * 3 + i * 3 + 0];\n float y1 = coords[b * N * 3 + i * 3 + 1];\n float z1 = coords[b * N * 3 + i * 3 + 2];\n \n // Iterate over all j points\n for (int j = 0; j < N; j++) {\n // Coalesced reads for point j (threads read consecutive j values)\n float x2 = coords[b * N * 3 + j * 3 + 0];\n float y1_ = coords[b * N * 3 + j * 3 + 1]; // rename to avoid conflict\n float z2 = coords[b * N * 3 + j * 3 + 2];\n \n float dx = x1 - x2;\n float dy = y1 - y1_;\n float dz = z1 - z2;\n float dist_sq = dx * dx + dy * dy + dz * dz;\n \n dist[b * N * N + i * N + j] = sqrtf(dist_sq + 1e-8f);\n }\n }\n}\n\nvoid launch_pairwise_distance(const float* coords, float* dist, int B, int N) {\n int threadsPerBlock = 256;\n dim3 numBlocks((N + threadsPerBlock - 1) / threadsPerBlock, B);\n pairwise_distance_kernel<<<numBlocks, threadsPerBlock>>>(coords, dist, B, N);\n}\n\"\"\"\n\ncpp_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\nvoid launch_pairwise_distance(const float* coords, float* dist, int B, int N);\n\ntorch::Tensor pairwise_distance_cuda(torch::Tensor coords) {\n TORCH_CHECK(coords.dim() == 3, \"coords must be 3D\");\n TORCH_CHECK(coords.size(2) == 3, \"coords last dimension must be 3\");\n \n int B = coords.size(0);\n int N = coords.size(1);\n \n auto dist = torch::empty({B, N, N}, coords.options());\n \n const float* coords_ptr = coords.data_ptr<float>();\n float* dist_ptr = dist.data_ptr<float>();\n \n launch_pairwise_distance(coords_ptr, dist_ptr, B, N);\n \n return dist;\n}\n\"\"\"\n\nmodule = load_inline(\n name='pairwise_distance_ext',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['pairwise_distance_cuda'],\n extra_cuda_cflags=['-O3'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n\n def forward(self, coords: torch.Tensor) -> torch.Tensor:\n return module.pairwise_distance_cuda(coords.contiguous())"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 7 · 1 tool call(s) · 337.9s turn · 17.3s LLM · tokens in=9,354 · out=1,681 · total=11,035 · 7,375 chars (click to expand)
turn 7337.9s turn · 17.3s LLM · 320.6s tools · tokens in=9,354 · out=1,681 · total=11,035
reasoning
Wait, I realized there's a problem with my approach. When each thread iterates over all j values, the reads of j coordinates are NOT coalesced across threads - each thread reads the entire column of j coordinates independently. The coalescing happens when threads in the same warp access consecutive memory locations at the same time.
Let me think about this more carefully:
- With my new kernel, each thread processes row i
- Within a warp, threads have consecutive i values
- When all threads read coords[b*N*3 + j*3 + 0], they're accessing memory with stride N*3 floats between threads
- This is NOT coalesced!
I need to think differently. For coalesced reads:
- Threads in a warp should access consecutive memory addresses
- For coords[B, N, 3], the last dimension (3) is contiguous
Actually, I realize the problem: the way the coordinate data is laid out, we can't easily get coalesced reads when threads work on different i values. Let me reconsider.
A better approach might be:
1. Use shared memory to cache the coordinates
2. Have threads in a warp work on consecutive j values (coalesced reads of coords for j)
3. Process the distance matrix in tiles
Actually, let me reconsider the original approach. With the 16x16 block:
- Each thread computes dist[i,j]
- Threads in a warp have consecutive threadIdx.x values
- When reading coords for j, threads access coords[b*N*3 + (j)*3 + 0] where j varies by threadIdx.x
- Since threadIdx.x varies by 1, j varies by 1, so threads access coords[b*N*3 + j*3 + 0] for consecutive j
- The addresses would be: base + j*3, base + (j+1)*3, base + (j+2)*3...
- These are NOT contiguous (stride of 3 floats = 12 bytes between threads)
- But x, y, z for the same j ARE contiguous (at offsets 0, 1, 2 from j*3)
Hmm, stride of 3 floats is actually decently coalesced - each thread reads 3 consecutive floats (x, y, z for one j). Let me check what the issue was.
Actually wait - in the original kernel, the memory access pattern was:
- Thread (ty, tx) reads coords[b, i, :] and coords[b, j, :] where i = by*blockDim.y + ty, j = bx*blockDim.x + tx
- Within a warp, tx varies (typically first 16 threads have the same ty, consecutive tx)
- So threads read coords[b, j*3+0] for consecutive j values
- Addresses: base + j*3, base + (j+1)*3, ... (stride 3)
Actually this should be decent for coalescing. But wait - each thread reads 6 floats total (both coords). And the i coordinate varies with ty, which is slow-varying in a warp.
Let me think again. With a 16x16 block:
- Warp size is 32
- A warp spans 2 rows of the thread block
- For the first warp: threads 0-15 have (ty=0, tx=0-15), threads 16-31 have (ty=1, tx=0-15)
- For j-coordinate reads: threads read coords[b, j*3+0] where j = tx for first 16 threads, then j = tx for next 16
- The addresses are NOT contiguous because of the ty difference
Actually the j values repeat within a warp (first 16 threads read j=0-15, next 16 also read j=0-15), so there's no coalescing benefit.
Let me redesign with better memory access. I'll use a tile-based approach with shared memory.
Actually, the real issue is that each thread needs two coordinates: i and j. The i coordinate causes the non-coalesced reads.
A better approach:
1. Each thread block handles one row of the output matrix
2. Threads in the block collaborate to load coords[j] into shared memory (coalesced)
3. Each thread computes dist[i, j] for multiple j values
Let me write this:
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 = \"\"\"\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <cmath>\n\n// Optimized kernel: each thread processes one row of the distance matrix\n// This allows coalesced reads for the second coordinate\n__global__ void pairwise_distance_kernel(\n const float* __restrict__ coords,\n float* __restrict__ dist,\n const int B,\n const int N\n) {\n // Each thread handles one row (i) of the output\n int b = blockIdx.y;\n int i = blockIdx.x * blockDim.x + threadIdx.x;\n \n if (b < B && i < N) {\n // Load coordinates for point i (reused N times)\n float x1 = coords[b * N * 3 + i * 3 + 0];\n float y1 = coords[b * N * 3 + i * 3 + 1];\n float z1 = coords[b * N * 3 + i * 3 + 2];\n \n // Iterate over all j points\n for (int j = 0; j < N; j++) {\n // Coalesced reads for point j (threads read consecutive j values)\n float x2 = coords[b * N * 3 + j * 3 + 0];\n float y1_ = coords[b * N * 3 + j * 3 + 1]; // rename to avoid conflict\n float z2 = coords[b * N * 3 + j * 3 + 2];\n \n float dx = x1 - x2;\n float dy = y1 - y1_;\n float dz = z1 - z2;\n float dist_sq = dx * dx + dy * dy + dz * dz;\n \n dist[b * N * N + i * N + j] = sqrtf(dist_sq + 1e-8f);\n }\n }\n}\n\nvoid launch_pairwise_distance(const float* coords, float* dist, int B, int N) {\n int threadsPerBlock = 256;\n dim3 numBlocks((N + threadsPerBlock - 1) / threadsPerBlock, B);\n pairwise_distance_kernel<<<numBlocks, threadsPerBlock>>>(coords, dist, B, N);\n}\n\"\"\"\n\ncpp_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\nvoid launch_pairwise_distance(const float* coords, float* dist, int B, int N);\n\ntorch::Tensor pairwise_distance_cuda(torch::Tensor coords) {\n TORCH_CHECK(coords.dim() == 3, \"coords must be 3D\");\n TORCH_CHECK(coords.size(2) == 3, \"coords last dimension must be 3\");\n \n int B = coords.size(0);\n int N = coords.size(1);\n \n auto dist = torch::empty({B, N, N}, coords.options());\n \n const float* coords_ptr = coords.data_ptr<float>();\n float* dist_ptr = dist.data_ptr<float>();\n \n launch_pairwise_distance(coords_ptr, dist_ptr, B, N);\n \n return dist;\n}\n\"\"\"\n\nmodule = load_inline(\n name='pairwise_distance_ext',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['pairwise_distance_cuda'],\n extra_cuda_cflags=['-O3'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n\n def forward(self, coords: torch.Tensor) -> torch.Tensor:\n return module.pairwise_distance_cuda(coords.contiguous())"
}
result ← run_correctness (ok)
okrun_correctness
run_correctness PASSED: (3 / 3) trials all matched the reference.
Numerical precision: max_abs_err=4.77e-07 mean_abs_err=2.82e-08 max_rel_err=1.65e-07
turn 8318.2s turn · 7.1s LLM · 311.1s tools · tokens in=10,187 · out=786 · total=10,973
reasoning
Good, correctness is maintained. Let me profile this version.
call → profile_kernel
callprofile_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = \"\"\"\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <cmath>\n\n// Optimized kernel: each thread processes one row of the distance matrix\n// This allows coalesced reads for the second coordinate\n__global__ void pairwise_distance_kernel(\n const float* __restrict__ coords,\n float* __restrict__ dist,\n const int B,\n const int N\n) {\n // Each thread handles one row (i) of the output\n int b = blockIdx.y;\n int i = blockIdx.x * blockDim.x + threadIdx.x;\n \n if (b < B && i < N) {\n // Load coordinates for point i (reused N times)\n float x1 = coords[b * N * 3 + i * 3 + 0];\n float y1 = coords[b * N * 3 + i * 3 + 1];\n float z1 = coords[b * N * 3 + i * 3 + 2];\n \n // Iterate over all j points\n for (int j = 0; j < N; j++) {\n // Coalesced reads for point j (threads read consecutive j values)\n float x2 = coords[b * N * 3 + j * 3 + 0];\n float y1_ = coords[b * N * 3 + j * 3 + 1]; // rename to avoid conflict\n float z2 = coords[b * N * 3 + j * 3 + 2];\n \n float dx = x1 - x2;\n float dy = y1 - y1_;\n float dz = z1 - z2;\n float dist_sq = dx * dx + dy * dy + dz * dz;\n \n dist[b * N * N + i * N + j] = sqrtf(dist_sq + 1e-8f);\n }\n }\n}\n\nvoid launch_pairwise_distance(const float* coords, float* dist, int B, int N) {\n int threadsPerBlock = 256;\n dim3 numBlocks((N + threadsPerBlock - 1) / threadsPerBlock, B);\n pairwise_distance_kernel<<<numBlocks, threadsPerBlock>>>(coords, dist, B, N);\n}\n\"\"\"\n\ncpp_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\nvoid launch_pairwise_distance(const float* coords, float* dist, int B, int N);\n\ntorch::Tensor pairwise_distance_cuda(torch::Tensor coords) {\n TORCH_CHECK(coords.dim() == 3, \"coords must be 3D\");\n TORCH_CHECK(coords.size(2) == 3, \"coords last dimension must be 3\");\n \n int B = coords.size(0);\n int N = coords.size(1);\n auto dist = torch::empty({B, N, N}, coords.options());\n \n const float* coords_ptr = coords.data_ptr<float>();\n float* dist_ptr = dist.data_ptr<float>();\n \n launch_pairwise_distance(coords_ptr, dist_ptr, B, N);\n \n return dist;\n}\n\"\"\"\n\nmodule = load_inline(\n name='pairwise_distance_ext',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['pairwise_distance_cuda'],\n extra_cuda_cflags=['-O3'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n\n def forward(self, coords: torch.Tensor) -> torch.Tensor:\n return module.pairwise_distance_cuda(coords.contiguous())"
}
result ← profile_kernel (ok)
okprofile_kernel
profile_kernel PASSED: profiling complete.
=== Kernel Profile Summary ===
--- Kernel Breakdown (by GPU time) ---
#1: aten::empty 1622.0μs (100%) calls=1
--- Timing ---
GPU time: 142.05 μs [↑133.5μs REGRESSED]
--- Memory Bandwidth ---
DRAM BW: 0.4/3350 GB/s (0.0%) [↓0.2% REGRESSED]
read=0.4 write=0.0 GB/s
Hit rates: L2=100.2%
--- Compute Throughput ---
Pipe util: FMA=4.3%, Tensor=0.0%
--- Roofline Analysis ---
Ridge point (FP32): 20.00 FLOPs/byte
Classification: UNKNOWN
--- Occupancy ---
Achieved: 12.5% [↓58.5% REGRESSED]
Eligible warps/cycle: 0.09
WARNING: <1 eligible warp/cycle = latency-bound, needs more parallelism
Launch: regs/thread=24, smem/blk=1024B, block=256, grid=16
--- Memory Access Quality ---
Global ld coalescing: 1.02 sectors/req (perfect)
--- Warp Stall Reasons (% of issue cycles) ---
long_scoreboard 1796.2% ####################
short_scoreboard 2.8%
--- Optimization Hints ---
-> 1796% stalls: global memory latency (waiting for DRAM/L2). Try shared-memory tiling, prefetching, or reducing working set.
-> Low occupancy (12%). Consider __launch_bounds__ or reducing register pressure.
-> Low eligible warps + low BW% + low compute% = LATENCY-BOUND. Need more parallelism (larger grid, more ILP, prefetching).
--- Delta vs Previous Iteration ---
GPU time: 8.58→142.05μs (↑133.47) REGRESSED
DRAM BW%: 0.18→0.01% (↓0.17) REGRESSED
Occupancy: 71.03→12.49% (↓58.54) REGRESSED
Eligible warps/cyc: 1.28→0.09 (↓1.19) REGRESSED
Ld sectors/req: 3.67→1.02 (↓2.65) ok
L2 hit%: 97.91→100.23% (↑2.32) ok
turn 9339.7s turn · 12.9s LLM · 326.8s tools · tokens in=11,519 · out=1,287 · total=12,806
reasoning
The kernel regressed! The issue is that with only 16 blocks total (2 blocks × 8 batches), we have very low occupancy. The first kernel had 8192 blocks which gave much better parallelism.
The problem is that with N=512 atoms and 256 threads per block, we only get 2 blocks per batch. With 8 batches, that's only 16 blocks total.
I need to go back to the approach with more blocks, but with better memory access patterns. Let me try a different approach where each thread block processes multiple rows but coordinates the memory access better.
Actually, let me think about this differently. The original kernel had:
- 32x32 blocks per batch = 1024 blocks per batch
- With 8 batches = 8192 total blocks
The issue was that each thread loaded both i and j coordinates, leading to non-coalesced reads.
A better approach is:
1. Have each thread compute one output element (i, j)
2. But use shared memory to cache coordinates
Let me try a shared memory tiling approach where each thread block loads a tile of coordinates into shared memory, then all threads in the block can compute distances using that cached data.
assistantThe kernel regressed significantly due to low occupancy (only 16 blocks total). Let me try a better approach using shared memory tiling to improve data reuse while maintaining good parallelism:
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 = \"\"\"\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <cmath>\n\n// Shared memory tiled kernel for better data reuse\n// Each block computes a TILE x TILE block of the output\ntemplate <int TILE_SIZE>\n__global__ void pairwise_distance_kernel(\n const float* __restrict__ coords,\n float* __restrict__ dist,\n const int B,\n const int N\n) {\n int b = blockIdx.z;\n int i_start = blockIdx.y * TILE_SIZE;\n int j_start = blockIdx.x * TILE_SIZE;\n \n int i_local = threadIdx.y;\n int j_local = threadIdx.x;\n \n int i = i_start + i_local;\n int j = j_start + j_local;\n \n // Shared memory for tile of coordinates\n __shared__ float s_i_coords[TILE_SIZE][3]; // coordinates for i points\n __shared__ float s_j_coords[TILE_SIZE][3]; // coordinates for j points\n \n // Load coordinates for i points (collaborative loading)\n if (i < N) {\n s_i_coords[i_local][0] = coords[b * N * 3 + i * 3 + 0];\n s_i_coords[i_local][1] = coords[b * N * 3 + i * 3 + 1];\n s_i_coords[i_local][2] = coords[b * N * 3 + i * 3 + 2];\n }\n \n // Load coordinates for j points (collaborative loading)\n if (j < N) {\n s_j_coords[j_local][0] = coords[b * N * 3 + j * 3 + 0];\n s_j_coords[j_local][1] = coords[b * N * 3 + j * 3 + 1];\n s_j_coords[j_local][2] = coords[b * N * 3 + j * 3 + 2];\n }\n \n __syncthreads();\n \n // Compute distance for this output element\n if (i < N && j < N) {\n float x1 = s_i_coords[i_local][0];\n float y1 = s_i_coords[i_local][1];\n float z1 = s_i_coords[i_local][2];\n \n float x2 = s_j_coords[j_local][0];\n float y2 = s_j_coords[j_local][1];\n float z2 = s_j_coords[j_local][2];\n \n float dx = x1 - x2;\n float dy = y1 - y2;\n float dz = z1 - z2;\n float dist_sq = dx * dx + dy * dy + dz * dz;\n \n dist[b * N * N + i * N + j] = sqrtf(dist_sq + 1e-8f);\n }\n}\n\nvoid launch_pairwise_distance(const float* coords, float* dist, int B, int N) {\n const int TILE_SIZE = 16;\n dim3 threadsPerBlock(TILE_SIZE, TILE_SIZE);\n dim3 numBlocks((N + TILE_SIZE - 1) / TILE_SIZE,\n (N + TILE_SIZE - 1) / TILE_SIZE,\n B);\n pairwise_distance_kernel<TILE_SIZE><<<numBlocks, threadsPerBlock>>>(coords, dist, B, N);\n}\n\"\"\"\n\ncpp_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\nvoid launch_pairwise_distance(const float* coords, float* dist, int B, int N);\n\ntorch::Tensor pairwise_distance_cuda(torch::Tensor coords) {\n TORCH_CHECK(coords.dim() == 3, \"coords must be 3D\");\n TORCH_CHECK(coords.size(2) == 3, \"coords last dimension must be 3\");\n \n int B = coords.size(0);\n int N = coords.size(1);\n auto dist = torch::empty({B, N, N}, coords.options());\n \n const float* coords_ptr = coords.data_ptr<float>();\n float* dist_ptr = dist.data_ptr<float>();\n \n launch_pairwise_distance(coords_ptr, dist_ptr, B, N);\n \n return dist;\n}\n\"\"\"\n\nmodule = load_inline(\n name='pairwise_distance_ext',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['pairwise_distance_cuda'],\n extra_cuda_cflags=['-O3'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n\n def forward(self, coords: torch.Tensor) -> torch.Tensor:\n return module.pairwise_distance_cuda(coords.contiguous())"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 10344.2s turn · 9.2s LLM · 335.0s tools · tokens in=12,568 · out=1,007 · total=13,575
reasoning
Good, let me test correctness.
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 = \"\"\"\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <cmath>\n\n// Shared memory tiled kernel for better data reuse\n// Each block computes a TILE x TILE block of the output\ntemplate <int TILE_SIZE>\n__global__ void pairwise_distance_kernel(\n const float* __restrict__ coords,\n float* __restrict__ dist,\n const int B,\n const int N\n) {\n int b = blockIdx.z;\n int i_start = blockIdx.y * TILE_SIZE;\n int j_start = blockIdx.x * TILE_SIZE;\n \n int i_local = threadIdx.y;\n int j_local = threadIdx.x;\n \n int i = i_start + i_local;\n int j = j_start + j_local;\n \n // Shared memory for tile of coordinates\n __shared__ float s_i_coords[TILE_SIZE][3]; // coordinates for i points\n __shared__ float s_j_coords[TILE_SIZE][3]; // coordinates for j points\n \n // Load coordinates for i points (collaborative loading)\n if (i < N) {\n s_i_coords[i_local][0] = coords[b * N * 3 + i * 3 + 0];\n s_i_coords[i_local][1] = coords[b * N * 3 + i * 3 + 1];\n s_i_coords[i_local][2] = coords[b * N * 3 + i * 3 + 2];\n }\n \n // Load coordinates for j points (collaborative loading)\n if (j < N) {\n s_j_coords[j_local][0] = coords[b * N * 3 + j * 3 + 0];\n s_j_coords[j_local][1] = coords[b * N * 3 + j * 3 + 1];\n s_j_coords[j_local][2] = coords[b * N * 3 + j * 3 + 2];\n }\n \n __syncthreads();\n \n // Compute distance for this output element\n if (i < N && j < N) {\n float x1 = s_i_coords[i_local][0];\n float y1 = s_i_coords[i_local][1];\n float z1 = s_i_coords[i_local][2];\n \n float x2 = s_j_coords[j_local][0];\n float y2 = s_j_coords[j_local][1];\n float z2 = s_j_coords[j_local][2];\n \n float dx = x1 - x2;\n float dy = y1 - y2;\n float dz = z1 - z2;\n float dist_sq = dx * dx + dy * dy + dz * dz;\n \n dist[b * N * N + i * N + j] = sqrtf(dist_sq + 1e-8f);\n }\n}\n\nvoid launch_pairwise_distance(const float* coords, float* dist, int B, int N) {\n const int TILE_SIZE = 16;\n dim3 threadsPerBlock(TILE_SIZE, TILE_SIZE);\n dim3 numBlocks((N + TILE_SIZE - 1) / TILE_SIZE,\n (N + TILE_SIZE - 1) / TILE_SIZE,\n B);\n pairwise_distance_kernel<TILE_SIZE><<<numBlocks, threadsPerBlock>>>(coords, dist, B, N);\n}\n\"\"\"\n\ncpp_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\nvoid launch_pairwise_distance(const float* coords, float* dist, int B, int N);\n\ntorch::Tensor pairwise_distance_cuda(torch::Tensor coords) {\n TORCH_CHECK(coords.dim() == 3, \"coords must be 3D\");\n TORCH_CHECK(coords.size(2) == 3, \"coords last dimension must be 3\");\n \n int B = coords.size(0);\n int N = coords.size(1);\n auto dist = torch::empty({B, N, N}, coords.options());\n \n const float* coords_ptr = coords.data_ptr<float>();\n float* dist_ptr = dist.data_ptr<float>();\n \n launch_pairwise_distance(coords_ptr, dist_ptr, B, N);\n \n return dist;\n}\n\"\"\"\n\nmodule = load_inline(\n name='pairwise_distance_ext',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['pairwise_distance_cuda'],\n extra_cuda_cflags=['-O3'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n\n def forward(self, coords: torch.Tensor) -> torch.Tensor:\n return module.pairwise_distance_cuda(coords.contiguous())"
}
result ← run_correctness (ok)
okrun_correctness
run_correctness PASSED: (3 / 3) trials all matched the reference.
Numerical precision: max_abs_err=4.77e-07 mean_abs_err=2.82e-08 max_rel_err=1.65e-07
turn 11 · 1 tool call(s) · 355.0s turn · 9.0s LLM · tokens in=13,625 · out=1,005 · total=14,630 · 6,203 chars (click to expand)
turn 11355.0s turn · 9.0s LLM · 346.0s tools · tokens in=13,625 · out=1,005 · total=14,630
reasoning
Let me profile this version.
call → profile_kernel
callprofile_kernel
{
"kernel_code": "import torch\nimport torch.nn as nn\nfrom torch.utils.cpp_extension import load_inline\n\ncuda_source = \"\"\"\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <cmath>\n\n// Shared memory tiled kernel for better data reuse\n// Each block computes a TILE x TILE block of the output\ntemplate <int TILE_SIZE>\n__global__ void pairwise_distance_kernel(\n const float* __restrict__ coords,\n float* __restrict__ dist,\n const int B,\n const int N\n) {\n int b = blockIdx.z;\n int i_start = blockIdx.y * TILE_SIZE;\n int j_start = blockIdx.x * TILE_SIZE;\n \n int i_local = threadIdx.y;\n int j_local = threadIdx.x;\n \n int i = i_start + i_local;\n int j = j_start + j_local;\n \n // Shared memory for tile of coordinates\n __shared__ float s_i_coords[TILE_SIZE][3]; // coordinates for i points\n __shared__ float s_j_coords[TILE_SIZE][3]; // coordinates for j points\n \n // Load coordinates for i points (collaborative loading)\n if (i < N) {\n s_i_coords[i_local][0] = coords[b * N * 3 + i * 3 + 0];\n s_i_coords[i_local][1] = coords[b * N * 3 + i * 3 + 1];\n s_i_coords[i_local][2] = coords[b * N * 3 + i * 3 + 2];\n }\n \n // Load coordinates for j points (collaborative loading)\n if (j < N) {\n s_j_coords[j_local][0] = coords[b * N * 3 + j * 3 + 0];\n s_j_coords[j_local][1] = coords[b * N * 3 + j * 3 + 1];\n s_j_coords[j_local][2] = coords[b * N * 3 + j * 3 + 2];\n }\n \n __syncthreads();\n \n // Compute distance for this output element\n if (i < N && j < N) {\n float x1 = s_i_coords[i_local][0];\n float y1 = s_i_coords[i_local][1];\n float z1 = s_i_coords[i_local][2];\n \n float x2 = s_j_coords[j_local][0];\n float y2 = s_j_coords[j_local][1];\n float z2 = s_j_coords[j_local][2];\n \n float dx = x1 - x2;\n float dy = y1 - y2;\n float dz = z1 - z2;\n float dist_sq = dx * dx + dy * dy + dz * dz;\n \n dist[b * N * N + i * N + j] = sqrtf(dist_sq + 1e-8f);\n }\n}\n\nvoid launch_pairwise_distance(const float* coords, float* dist, int B, int N) {\n const int TILE_SIZE = 16;\n dim3 threadsPerBlock(TILE_SIZE, TILE_SIZE);\n dim3 numBlocks((N + TILE_SIZE - 1) / TILE_SIZE,\n (N + TILE_SIZE - 1) / TILE_SIZE,\n B);\n pairwise_distance_kernel<TILE_SIZE><<<numBlocks, threadsPerBlock>>>(coords, dist, B, N);\n}\n\"\"\"\n\ncpp_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\nvoid launch_pairwise_distance(const float* coords, float* dist, int B, int N);\n\ntorch::Tensor pairwise_distance_cuda(torch::Tensor coords) {\n TORCH_CHECK(coords.dim() == 3, \"coords must be 3D\");\n TORCH_CHECK(coords.size(2) == 3, \"coords last dimension must be 3\");\n \n int B = coords.size(0);\n int N = coords.size(1);\n auto dist = torch::empty({B, N, N}, coords.options());\n \n const float* coords_ptr = coords.data_ptr<float>();\n float* dist_ptr = dist.data_ptr<float>();\n \n launch_pairwise_distance(coords_ptr, dist_ptr, B, N);\n \n return dist;\n}\n\"\"\"\n\nmodule = load_inline(\n name='pairwise_distance_ext',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['pairwise_distance_cuda'],\n extra_cuda_cflags=['-O3'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n\n def forward(self, coords: torch.Tensor) -> torch.Tensor:\n return module.pairwise_distance_cuda(coords.contiguous())"
}
result ← profile_kernel (ok)
okprofile_kernel
profile_kernel PASSED: profiling complete.
=== Kernel Profile Summary ===
--- Kernel Breakdown (by GPU time) ---
#1: aten::empty 1759.0μs (100%) calls=1
--- Timing ---
GPU time: 10.50 μs [↓131.6μs ok]
--- Memory Bandwidth ---
DRAM BW: 5.1/3350 GB/s (0.2%) [↑0.1% ok]
read=5.1 write=0.0 GB/s
Hit rates: L2=97.8%
--- Compute Throughput ---
Dominant pipe: FP32 util=17.4%
Pipe util: FMA=17.4%, Tensor=0.0%
--- Roofline Analysis ---
Ridge point (FP32): 20.00 FLOPs/byte
Classification: COMPUTE-BOUND
--- Occupancy ---
Achieved: 83.6% [↑71.1% ok]
Eligible warps/cycle: 2.28
Launch: regs/thread=25, smem/blk=1408B, block=256, grid=8192
--- Memory Access Quality ---
Global ld coalescing: 3.67 sectors/req (3.7x wasted BW)
--- Warp Stall Reasons (% of issue cycles) ---
mio_throttle 697.8% ####################
long_scoreboard 336.3% ####################
short_scoreboard 279.7% ####################
--- Optimization Hints ---
-> 698% stalls: memory-instruction throttle. Too many outstanding memory ops; batch loads, reduce address divergence.
-> Global load coalescing=3.7x (ideal=1.0) -> 4x wasted DRAM BW. Make per-warp accesses contiguous.
-> Compute-bound but NOT using tensor cores. If applicable, use FP16/BF16 with torch.matmul or WMMA intrinsics.
--- Delta vs Previous Iteration ---
GPU time: 142.05→10.50μs (↓131.55) ok
DRAM BW%: 0.01→0.15% (↑0.14) ok
Occupancy: 12.49→83.59% (↑71.10) ok
Eligible warps/cyc: 0.09→2.28 (↑2.19) ok
Ld sectors/req: 1.02→3.67 (↑2.65) REGRESSED
L2 hit%: 100.23→97.81% (↓2.42) REGRESSED
turn 12 · 1 tool call(s) · 248.0s turn · 16.2s LLM · tokens in=15,236 · out=1,672 · total=16,908 · 7,675 chars (click to expand)
turn 12248.0s turn · 16.2s LLM · 231.9s tools · tokens in=15,236 · out=1,672 · total=16,908
reasoning
The shared memory tiled kernel is now at 10.50 μs which is better than the previous version. However, the global load coalescing is still poor (3.67 sectors/req). Let me try to improve the memory access pattern.
Actually, let me look at the issue more carefully. The problem is that the shared memory loading pattern for `s_i_coords` and `s_j_coords` is not coalesced. Each thread loads coordinates for its i/j point, but the coordinates are stored as [x, y, z] for each point, so threads in the same row access scattered memory locations.
A better approach would be to have threads in a warp load consecutive points, so they can read consecutive floats from the coords array. Let me redesign the kernel with better memory access patterns.
Actually, looking at the results again - 10.50 μs with 83.6% occupancy is decent. The issue is still the memory coalescing. Let me try a different approach where I use a 1D thread block and have each thread compute multiple output elements to reduce memory pressure.
Actually, let me think about this differently. The coordinates are stored as (B, N, 3). When threads in a warp load `s_i_coords`, thread with threadIdx.y=0 loads coords for point i_start+0, thread with threadIdx.y=1 loads coords for point i_start+1, etc. But threadIdx.x varies from 0 to 15, and they all load the same i point coordinates! This is inefficient because multiple threads are loading the same data.
Let me reconsider. In the tiled approach:
- Thread (ty, tx) is responsible for output element (i_start+ty, j_start+tx)
- We load TILE_SIZE i-points and TILE_SIZE j-points into shared memory
- Each i-point is loaded by all threads in a row (same threadIdx.y)
- Each j-point is loaded by all threads in a column (same threadIdx.x)
The issue is that for loading i-points, threads with different threadIdx.x but same threadIdx.y all try to load the same coordinates. This is redundant and not coalesced.
A better approach would be:
1. Have threads in a warp cooperatively load the coordinates
2. Use a 1D block of 256 threads, each responsible for one output element
3. Use shared memory to cache all coordinates for a batch
Actually, let me try a different approach: use vectorized loads (float4) to improve memory throughput, and ensure coalesced access by reorganizing how coordinates are loaded.
Let me try a simpler optimization first - adjusting the tile loading to use coalesced accesses.
assistantThe shared memory kernel shows good occupancy (83.6%) but still has poor load coalescing. Let me try to improve the memory access pattern by reorganizing how coordinates are loaded into shared memory:
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 = \"\"\"\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <cmath>\n\n// Optimized shared memory kernel with better coalescing\n// Each thread block handles TILE x TILE output elements\n// Use 1D indexing for loading to ensure coalesced accesses\ntemplate <int TILE_SIZE>\n__global__ void pairwise_distance_kernel(\n const float* __restrict__ coords,\n float* __restrict__ dist,\n const int B,\n const int N\n) {\n int b = blockIdx.z;\n int i_start = blockIdx.y * TILE_SIZE;\n int j_start = blockIdx.x * TILE_SIZE;\n \n int i_local = threadIdx.y;\n int j_local = threadIdx.x;\n \n int i = i_start + i_local;\n int j = j_start + j_local;\n \n // Shared memory for tiles of coordinates - flattened for better bank access\n __shared__ float s_i_x[TILE_SIZE];\n __shared__ float s_i_y[TILE_SIZE];\n __shared__ float s_i_z[TILE_SIZE];\n __shared__ float s_j_x[TILE_SIZE];\n __shared__ float s_j_y[TILE_SIZE];\n __shared__ float s_j_z[TILE_SIZE];\n \n // Cooperative loading: each thread loads one coordinate\n // This ensures coalesced access when threads in a warp load consecutive data\n if (i_local == 0) {\n // Row 0 threads load i_coords for all rows\n for (int row = 0; row < TILE_SIZE; row++) {\n int idx = i_start + row;\n if (idx < N) {\n s_i_x[row] = coords[b * N * 3 + idx * 3 + 0];\n s_i_y[row] = coords[b * N * 3 + idx * 3 + 1];\n s_i_z[row] = coords[b * N * 3 + idx * 3 + 2];\n }\n }\n }\n if (j_local == 0) {\n // Column 0 threads load j_coords for all columns\n for (int col = 0; col < TILE_SIZE; col++) {\n int idx = j_start + col;\n if (idx < N) {\n s_j_x[col] = coords[b * N * 3 + idx * 3 + 0];\n s_j_y[col] = coords[b * N * 3 + idx * 3 + 1];\n s_j_z[col] = coords[b * N * 3 + idx * 3 + 2];\n }\n }\n }\n \n __syncthreads();\n \n // Compute distance for this output element\n if (i < N && j < N) {\n float dx = s_i_x[i_local] - s_j_x[j_local];\n float dy = s_i_y[i_local] - s_j_y[j_local];\n float dz = s_i_z[i_local] - s_j_z[j_local];\n float dist_sq = dx * dx + dy * dy + dz * dz;\n \n dist[b * N * N + i * N + j] = sqrtf(dist_sq + 1e-8f);\n }\n}\n\nvoid launch_pairwise_distance(const float* coords, float* dist, int B, int N) {\n const int TILE_SIZE = 16;\n dim3 threadsPerBlock(TILE_SIZE, TILE_SIZE);\n dim3 numBlocks((N + TILE_SIZE - 1) / TILE_SIZE,\n (N + TILE_SIZE - 1) / TILE_SIZE,\n B);\n pairwise_distance_kernel<TILE_SIZE><<<numBlocks, threadsPerBlock>>>(coords, dist, B, N);\n}\n\"\"\"\n\ncpp_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\nvoid launch_pairwise_distance(const float* coords, float* dist, int B, int N);\n\ntorch::Tensor pairwise_distance_cuda(torch::Tensor coords) {\n TORCH_CHECK(coords.dim() == 3, \"coords must be 3D\");\n TORCH_CHECK(coords.size(2) == 3, \"coords last dimension must be 3\");\n \n int B = coords.size(0);\n int N = coords.size(1);\n auto dist = torch::empty({B, N, N}, coords.options());\n \n const float* coords_ptr = coords.data_ptr<float>();\n float* dist_ptr = dist.data_ptr<float>();\n \n launch_pairwise_distance(coords_ptr, dist_ptr, B, N);\n \n return dist;\n}\n\"\"\"\n\nmodule = load_inline(\n name='pairwise_distance_ext',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['pairwise_distance_cuda'],\n extra_cuda_cflags=['-O3'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n\n def forward(self, coords: torch.Tensor) -> torch.Tensor:\n return module.pairwise_distance_cuda(coords.contiguous())"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 13195.5s turn · 9.5s LLM · 186.0s tools · tokens in=16,404 · out=1,001 · total=17,405
reasoning
The current shared memory kernel shows good occupancy (83.6%) and reasonable performance. The main issue is memory coalescing, but the shared memory approach is helping reduce global memory traffic. Let me run correctness and then consider submitting if it passes.
Actually, let me try one more optimization - using vectorized loads and improving the tiling strategy. But first let me check correctness of my latest kernel.
assistantLet me verify correctness and submit this kernel:
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 = \"\"\"\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <cmath>\n\n// Shared memory tiled kernel for pairwise distances\ntemplate <int TILE_SIZE>\n__global__ void pairwise_distance_kernel(\n const float* __restrict__ coords,\n float* __restrict__ dist,\n const int B,\n const int N\n) {\n int b = blockIdx.z;\n int i_start = blockIdx.y * TILE_SIZE;\n int j_start = blockIdx.x * TILE_SIZE;\n \n int i_local = threadIdx.y;\n int j_local = threadIdx.x;\n \n int i = i_start + i_local;\n int j = j_start + j_local;\n \n // Shared memory for tiles of coordinates\n __shared__ float s_i_coords[TILE_SIZE][3];\n __shared__ float s_j_coords[TILE_SIZE][3];\n \n // Load coordinates for i points\n if (i < N) {\n s_i_coords[i_local][0] = coords[b * N * 3 + i * 3 + 0];\n s_i_coords[i_local][1] = coords[b * N * 3 + i * 3 + 1];\n s_i_coords[i_local][2] = coords[b * N * 3 + i * 3 + 2];\n }\n \n // Load coordinates for j points\n if (j < N) {\n s_j_coords[j_local][0] = coords[b * N * 3 + j * 3 + 0];\n s_j_coords[j_local][1] = coords[b * N * 3 + j * 3 + 1];\n s_j_coords[j_local][2] = coords[b * N * 3 + j * 3 + 2];\n }\n \n __syncthreads();\n \n // Compute distance\n if (i < N && j < N) {\n float dx = s_i_coords[i_local][0] - s_j_coords[j_local][0];\n float dy = s_i_coords[i_local][1] - s_j_coords[j_local][1];\n float dz = s_i_coords[i_local][2] - s_j_coords[j_local][2];\n float dist_sq = dx * dx + dy * dy + dz * dz;\n dist[b * N * N + i * N + j] = sqrtf(dist_sq + 1e-8f);\n }\n}\n\nvoid launch_pairwise_distance(const float* coords, float* dist, int B, int N) {\n const int TILE_SIZE = 16;\n dim3 threadsPerBlock(TILE_SIZE, TILE_SIZE);\n dim3 numBlocks((N + TILE_SIZE - 1) / TILE_SIZE,\n (N + TILE_SIZE - 1) / TILE_SIZE,\n B);\n pairwise_distance_kernel<TILE_SIZE><<<numBlocks, threadsPerBlock>>>(coords, dist, B, N);\n}\n\"\"\"\n\ncpp_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\nvoid launch_pairwise_distance(const float* coords, float* dist, int B, int N);\n\ntorch::Tensor pairwise_distance_cuda(torch::Tensor coords) {\n TORCH_CHECK(coords.dim() == 3, \"coords must be 3D\");\n TORCH_CHECK(coords.size(2) == 3, \"coords last dimension must be 3\");\n \n int B = coords.size(0);\n int N = coords.size(1);\n auto dist = torch::empty({B, N, N}, coords.options());\n \n const float* coords_ptr = coords.data_ptr<float>();\n float* dist_ptr = dist.data_ptr<float>();\n \n launch_pairwise_distance(coords_ptr, dist_ptr, B, N);\n \n return dist;\n}\n\"\"\"\n\nmodule = load_inline(\n name='pairwise_distance_ext',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['pairwise_distance_cuda'],\n extra_cuda_cflags=['-O3'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n\n def forward(self, coords: torch.Tensor) -> torch.Tensor:\n return module.pairwise_distance_cuda(coords.contiguous())"
}
result ← run_correctness (ok)
okrun_correctness
run_correctness PASSED: (3 / 3) trials all matched the reference.
Numerical precision: max_abs_err=4.77e-07 mean_abs_err=2.82e-08 max_rel_err=1.65e-07
turn 14261.9s turn · 8.4s LLM · 253.5s tools · tokens in=17,470 · out=952 · total=18,422 · FINAL
reasoning
I have a correct kernel with good occupancy (83.6%) that passed correctness testing. The kernel uses shared memory tiling for data reuse and is showing compute-bound characteristics. I should submit this kernel now.
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 = \"\"\"\n#include <cuda.h>\n#include <cuda_runtime.h>\n#include <cmath>\n\n// Shared memory tiled kernel for pairwise distances\ntemplate <int TILE_SIZE>\n__global__ void pairwise_distance_kernel(\n const float* __restrict__ coords,\n float* __restrict__ dist,\n const int B,\n const int N\n) {\n int b = blockIdx.z;\n int i_start = blockIdx.y * TILE_SIZE;\n int j_start = blockIdx.x * TILE_SIZE;\n \n int i_local = threadIdx.y;\n int j_local = threadIdx.x;\n \n int i = i_start + i_local;\n int j = j_start + j_local;\n \n // Shared memory for tiles of coordinates\n __shared__ float s_i_coords[TILE_SIZE][3];\n __shared__ float s_j_coords[TILE_SIZE][3];\n \n // Load coordinates for i points\n if (i < N) {\n s_i_coords[i_local][0] = coords[b * N * 3 + i * 3 + 0];\n s_i_coords[i_local][1] = coords[b * N * 3 + i * 3 + 1];\n s_i_coords[i_local][2] = coords[b * N * 3 + i * 3 + 2];\n }\n \n // Load coordinates for j points\n if (j < N) {\n s_j_coords[j_local][0] = coords[b * N * 3 + j * 3 + 0];\n s_j_coords[j_local][1] = coords[b * N * 3 + j * 3 + 1];\n s_j_coords[j_local][2] = coords[b * N * 3 + j * 3 + 2];\n }\n \n __syncthreads();\n \n // Compute distance\n if (i < N && j < N) {\n float dx = s_i_coords[i_local][0] - s_j_coords[j_local][0];\n float dy = s_i_coords[i_local][1] - s_j_coords[j_local][1];\n float dz = s_i_coords[i_local][2] - s_j_coords[j_local][2];\n float dist_sq = dx * dx + dy * dy + dz * dz;\n dist[b * N * N + i * N + j] = sqrtf(dist_sq + 1e-8f);\n }\n}\n\nvoid launch_pairwise_distance(const float* coords, float* dist, int B, int N) {\n const int TILE_SIZE = 16;\n dim3 threadsPerBlock(TILE_SIZE, TILE_SIZE);\n dim3 numBlocks((N + TILE_SIZE - 1) / TILE_SIZE,\n (N + TILE_SIZE - 1) / TILE_SIZE,\n B);\n pairwise_distance_kernel<TILE_SIZE><<<numBlocks, threadsPerBlock>>>(coords, dist, B, N);\n}\n\"\"\"\n\ncpp_source = \"\"\"\n#include <torch/extension.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\nvoid launch_pairwise_distance(const float* coords, float* dist, int B, int N);\n\ntorch::Tensor pairwise_distance_cuda(torch::Tensor coords) {\n TORCH_CHECK(coords.dim() == 3, \"coords must be 3D\");\n TORCH_CHECK(coords.size(2) == 3, \"coords last dimension must be 3\");\n \n int B = coords.size(0);\n int N = coords.size(1);\n auto dist = torch::empty({B, N, N}, coords.options());\n \n const float* coords_ptr = coords.data_ptr<float>();\n float* dist_ptr = dist.data_ptr<float>();\n \n launch_pairwise_distance(coords_ptr, dist_ptr, B, N);\n \n return dist;\n}\n\"\"\"\n\nmodule = load_inline(\n name='pairwise_distance_ext',\n cpp_sources=cpp_source,\n cuda_sources=cuda_source,\n functions=['pairwise_distance_cuda'],\n extra_cuda_cflags=['-O3'],\n verbose=False\n)\n\nclass ModelNew(nn.Module):\n def __init__(self):\n super().__init__()\n\n def forward(self, coords: torch.Tensor) -> torch.Tensor:\n return module.pairwise_distance_cuda(coords.contiguous())"
}
result ← submit_kernel (ok)
oksubmit_kernel
submit_kernel PASSED: (3 / 3) 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=4.77e-07 mean_abs_err=2.82e-08 max_rel_err=1.65e-07
Memory: 41.05 MB peak (ref: 89.05 MB, ratio: 0.461x)
Kernel launches: 1 (ref: 12, fusion ratio: 12.0)
SOL score: 0.1725
final submitted kernel (106 lines, 3,180 chars) — full source, untruncated
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline
cuda_source = """
#include <cuda.h>
#include <cuda_runtime.h>
#include <cmath>
// Shared memory tiled kernel for pairwise distances
template <int TILE_SIZE>
__global__ void pairwise_distance_kernel(
const float* __restrict__ coords,
float* __restrict__ dist,
const int B,
const int N
) {
int b = blockIdx.z;
int i_start = blockIdx.y * TILE_SIZE;
int j_start = blockIdx.x * TILE_SIZE;
int i_local = threadIdx.y;
int j_local = threadIdx.x;
int i = i_start + i_local;
int j = j_start + j_local;
// Shared memory for tiles of coordinates
__shared__ float s_i_coords[TILE_SIZE][3];
__shared__ float s_j_coords[TILE_SIZE][3];
// Load coordinates for i points
if (i < N) {
s_i_coords[i_local][0] = coords[b * N * 3 + i * 3 + 0];
s_i_coords[i_local][1] = coords[b * N * 3 + i * 3 + 1];
s_i_coords[i_local][2] = coords[b * N * 3 + i * 3 + 2];
}
// Load coordinates for j points
if (j < N) {
s_j_coords[j_local][0] = coords[b * N * 3 + j * 3 + 0];
s_j_coords[j_local][1] = coords[b * N * 3 + j * 3 + 1];
s_j_coords[j_local][2] = coords[b * N * 3 + j * 3 + 2];
}
__syncthreads();
// Compute distance
if (i < N && j < N) {
float dx = s_i_coords[i_local][0] - s_j_coords[j_local][0];
float dy = s_i_coords[i_local][1] - s_j_coords[j_local][1];
float dz = s_i_coords[i_local][2] - s_j_coords[j_local][2];
float dist_sq = dx * dx + dy * dy + dz * dz;
dist[b * N * N + i * N + j] = sqrtf(dist_sq + 1e-8f);
}
}
void launch_pairwise_distance(const float* coords, float* dist, int B, int N) {
const int TILE_SIZE = 16;
dim3 threadsPerBlock(TILE_SIZE, TILE_SIZE);
dim3 numBlocks((N + TILE_SIZE - 1) / TILE_SIZE,
(N + TILE_SIZE - 1) / TILE_SIZE,
B);
pairwise_distance_kernel<TILE_SIZE><<<numBlocks, threadsPerBlock>>>(coords, dist, B, N);
}
"""
cpp_source = """
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
void launch_pairwise_distance(const float* coords, float* dist, int B, int N);
torch::Tensor pairwise_distance_cuda(torch::Tensor coords) {
TORCH_CHECK(coords.dim() == 3, "coords must be 3D");
TORCH_CHECK(coords.size(2) == 3, "coords last dimension must be 3");
int B = coords.size(0);
int N = coords.size(1);
auto dist = torch::empty({B, N, N}, coords.options());
const float* coords_ptr = coords.data_ptr<float>();
float* dist_ptr = dist.data_ptr<float>();
launch_pairwise_distance(coords_ptr, dist_ptr, B, N);
return dist;
}
"""
module = load_inline(
name='pairwise_distance_ext',
cpp_sources=cpp_source,
cuda_sources=cuda_source,
functions=['pairwise_distance_cuda'],
extra_cuda_cflags=['-O3'],
verbose=False
)
class ModelNew(nn.Module):
def __init__(self):
super().__init__()
def forward(self, coords: torch.Tensor) -> torch.Tensor:
return module.pairwise_distance_cuda(coords.contiguous())