correct
44_Average_Pooling_1D.py
started
2026-05-02T06:40:58Z
timingkernelbench.timing
kernel mean / median / std
1.76 / — / 0.15 μs
ref mean / median / std
5.46 / — / 0.06 μs
numerical precisioneval_kernel_against_ref
gpu memoryextended_metrics.measure_memory
kernel launches / fusiontorch.profiler
fusion ratio (ref / kernel)
6.00x
energyNVML (pynvml)
ratio (ref / kernel)
3.00x
SOL (speed-of-light)nsight_hardware_counters
roofline / occupancynsight_hardware_counters
dram bandwidth
2506.79 GB/s
warp stalls (top 5)nsight
reference PyTorch source (the prompt the agent sees)
import torch
import torch.nn as nn
class Model(nn.Module):
"""
Simple model that performs 1D Average Pooling.
"""
def __init__(self, kernel_size: int, stride: int = 1, padding: int = 0):
"""
Initializes the 1D Average Pooling layer.
Args:
kernel_size (int): Size of the pooling window.
stride (int, optional): Stride of the pooling operation. Defaults to 1.
padding (int, optional): Padding applied to the input tensor. Defaults to 0.
"""
super(Model, self).__init__()
self.avg_pool = nn.AvgPool1d(kernel_size=kernel_size, stride=stride, padding=padding)
def forward(self, x: torch.Tensor) -> torch.Tensor:
"""
Applies 1D Average Pooling to the input tensor.
Args:
x (torch.Tensor): Input tensor of shape (batch_size, in_channels, input_length).
Returns:
torch.Tensor: Output tensor with 1D Average Pooling applied, shape (batch_size, in_channels, output_length).
"""
return self.avg_pool(x)
batch_size = 64
in_channels = 128
input_length = 65536
kernel_size = 8
stride = 1
padding = 4
def get_inputs():
x = torch.rand(batch_size, in_channels, input_length)
return [x]
def get_init_inputs():
return [kernel_size, stride, padding]turn 07.0s turn · 7.0s LLM · 0.0s tools
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) · 31.4s turn · 31.4s LLM · 7,147 chars (click to expand)
turn 131.4s turn · 31.4s LLM · 0.0s tools
reasoning
**Evaluating output optimization**
I need to compute the sum for local output j over 8 values and think about optimizing that with vectorization. Using shared memory seems feasible, allowing for efficient loads. The kernel runtime could be around 1-2ms, but I should be cautious about a potential issue with partial tiles when outputting. Implementing rolling sums within threads could help reduce shared memory loads, improving efficiency in storing and processing outputs. Each thread can compute outputs in a structured way to maximize load efficiency.
**Analyzing thread storage**
I’m thinking about how thread four handles storing data. The compiler seems to emit each store offset as t*4+k, which, for k=0, results in addresses 0, 4, and 8. This leads to non-coalesced memory access with a stride of 16 bytes. That could slow things down a bit. I wonder if there's a more efficient way to manage these offsets to improve performance during data retrieval. Exploring coalescing might help optimize these thread interactions.
**Planning C++ and CUDA Code**
I'm preparing to write some code using C++ and CUDA for average pooling. I need to define a function for `avgpool1d_forward` and load in the corresponding CUDA files. I might not need PYBIND and will include necessary headers for torch.
I should consider padding and stride parameters, potentially using `torch.empty` for my output. Maybe I can utilize the `forward` function without copying if the tensor is already contiguous. I'll also need to compile with specific flags for optimization.
I’ll focus on ensuring correctness and performance in my final implementation.
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 avgpool1d_forward(torch::Tensor x, int64_t kernel_size, int64_t stride, int64_t padding);\n\"\"\"\n\n_cuda_src = r\"\"\"\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\nstatic constexpr int BLOCK = 256;\nstatic constexpr int TILE = 1024;\nstatic constexpr int K = 8;\nstatic constexpr int P = 4;\n\n__global__ __launch_bounds__(BLOCK, 2)\nvoid avgpool1d_k8s1p4_kernel(const float* __restrict__ x, float* __restrict__ y,\n int rows, int L, int Lout) {\n int tile = blockIdx.x;\n int row = blockIdx.y;\n int tid = threadIdx.x;\n int out0 = tile * TILE;\n int tile_count = Lout - out0;\n if (tile_count > TILE) tile_count = TILE;\n if (tile_count <= 0) return;\n\n extern __shared__ float sh[];\n // Need tile_count + K - 1 input/pad values starting at input index out0 - P.\n int load_count = tile_count + K - 1;\n const float* rowx = x + ((long long)row) * L;\n for (int i = tid; i < load_count; i += BLOCK) {\n int xi = out0 + i - P;\n float v = 0.0f;\n if ((unsigned)xi < (unsigned)L) v = rowx[xi];\n sh[i] = v;\n }\n __syncthreads();\n\n float* rowy = y + ((long long)row) * Lout;\n #pragma unroll\n for (int q = 0; q < 4; ++q) {\n int j = q * BLOCK + tid;\n if (j < tile_count) {\n float s = sh[j] + sh[j+1] + sh[j+2] + sh[j+3]\n + sh[j+4] + sh[j+5] + sh[j+6] + sh[j+7];\n rowy[out0 + j] = s * 0.125f;\n }\n }\n}\n\n// Generic, correctness-oriented fallback CUDA kernel for unusual constructor args.\n__global__ void avgpool1d_generic_kernel(const float* __restrict__ x, float* __restrict__ y,\n int rows, int L, int Lout, int kernel_size,\n int stride, int padding) {\n long long total = (long long)rows * Lout;\n long long idx = (long long)blockIdx.x * blockDim.x + threadIdx.x;\n long long step = (long long)blockDim.x * gridDim.x;\n for (; idx < total; idx += step) {\n int o = (int)(idx % Lout);\n int row = (int)(idx / Lout);\n int start = o * stride - padding;\n float sum = 0.0f;\n const float* rowx = x + ((long long)row) * L;\n for (int k = 0; k < kernel_size; ++k) {\n int xi = start + k;\n if ((unsigned)xi < (unsigned)L) sum += rowx[xi];\n }\n y[idx] = sum / (float)kernel_size;\n }\n}\n\ntorch::Tensor avgpool1d_forward(torch::Tensor x, int64_t kernel_size, int64_t stride, int64_t padding) {\n const int64_t N = x.size(0);\n const int64_t C = x.size(1);\n const int64_t L64 = x.size(2);\n const int64_t Lout64 = (L64 + 2 * padding - kernel_size) / stride + 1;\n auto y = torch::empty({N, C, Lout64}, x.options());\n int rows = (int)(N * C);\n int L = (int)L64;\n int Lout = (int)Lout64;\n cudaStream_t stream = at::cuda::getCurrentCUDAStream();\n if (x.scalar_type() == at::kFloat && kernel_size == 8 && stride == 1 && padding == 4) {\n dim3 block(BLOCK);\n dim3 grid((Lout + TILE - 1) / TILE, rows);\n size_t smem = (TILE + K - 1) * sizeof(float);\n avgpool1d_k8s1p4_kernel<<<grid, block, smem, stream>>>(\n x.data_ptr<float>(), y.data_ptr<float>(), rows, L, Lout);\n } else {\n int threads = 256;\n long long total = (long long)rows * Lout;\n int blocks = (int)((total + threads - 1) / threads);\n if (blocks > 65535) blocks = 65535;\n avgpool1d_generic_kernel<<<blocks, threads, 0, stream>>>(\n x.data_ptr<float>(), y.data_ptr<float>(), rows, L, Lout, (int)kernel_size, (int)stride, (int)padding);\n }\n return y;\n}\n\"\"\"\n\n_ext = load_inline(\n name=\"avgpool1d_ext_v1\",\n cpp_sources=_cpp_src,\n cuda_sources=_cuda_src,\n functions=[\"avgpool1d_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, kernel_size: int, stride: int = 1, padding: int = 0):\n super(ModelNew, self).__init__()\n self.kernel_size = int(kernel_size)\n self.stride = int(stride)\n self.padding = int(padding)\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n return _ext.avgpool1d_forward(x, self.kernel_size, self.stride, self.padding)\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 210.3s turn · 10.3s 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 avgpool1d_forward(torch::Tensor x, int64_t kernel_size, int64_t stride, int64_t padding);\n\"\"\"\n\n_cuda_src = r\"\"\"\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\nstatic constexpr int BLOCK = 256;\nstatic constexpr int TILE = 1024;\nstatic constexpr int K = 8;\nstatic constexpr int P = 4;\n\n__global__ __launch_bounds__(BLOCK, 2)\nvoid avgpool1d_k8s1p4_kernel(const float* __restrict__ x, float* __restrict__ y,\n int rows, int L, int Lout) {\n int tile = blockIdx.x;\n int row = blockIdx.y;\n int tid = threadIdx.x;\n int out0 = tile * TILE;\n int tile_count = Lout - out0;\n if (tile_count > TILE) tile_count = TILE;\n if (tile_count <= 0) return;\n\n extern __shared__ float sh[];\n // Need tile_count + K - 1 input/pad values starting at input index out0 - P.\n int load_count = tile_count + K - 1;\n const float* rowx = x + ((long long)row) * L;\n for (int i = tid; i < load_count; i += BLOCK) {\n int xi = out0 + i - P;\n float v = 0.0f;\n if ((unsigned)xi < (unsigned)L) v = rowx[xi];\n sh[i] = v;\n }\n __syncthreads();\n\n float* rowy = y + ((long long)row) * Lout;\n #pragma unroll\n for (int q = 0; q < 4; ++q) {\n int j = q * BLOCK + tid;\n if (j < tile_count) {\n float s = sh[j] + sh[j+1] + sh[j+2] + sh[j+3]\n + sh[j+4] + sh[j+5] + sh[j+6] + sh[j+7];\n rowy[out0 + j] = s * 0.125f;\n }\n }\n}\n\n// Generic, correctness-oriented fallback CUDA kernel for unusual constructor args.\n__global__ void avgpool1d_generic_kernel(const float* __restrict__ x, float* __restrict__ y,\n int rows, int L, int Lout, int kernel_size,\n int stride, int padding) {\n long long total = (long long)rows * Lout;\n long long idx = (long long)blockIdx.x * blockDim.x + threadIdx.x;\n long long step = (long long)blockDim.x * gridDim.x;\n for (; idx < total; idx += step) {\n int o = (int)(idx % Lout);\n int row = (int)(idx / Lout);\n int start = o * stride - padding;\n float sum = 0.0f;\n const float* rowx = x + ((long long)row) * L;\n for (int k = 0; k < kernel_size; ++k) {\n int xi = start + k;\n if ((unsigned)xi < (unsigned)L) sum += rowx[xi];\n }\n y[idx] = sum / (float)kernel_size;\n }\n}\n\ntorch::Tensor avgpool1d_forward(torch::Tensor x, int64_t kernel_size, int64_t stride, int64_t padding) {\n const int64_t N = x.size(0);\n const int64_t C = x.size(1);\n const int64_t L64 = x.size(2);\n const int64_t Lout64 = (L64 + 2 * padding - kernel_size) / stride + 1;\n auto y = torch::empty({N, C, Lout64}, x.options());\n int rows = (int)(N * C);\n int L = (int)L64;\n int Lout = (int)Lout64;\n cudaStream_t stream = at::cuda::getCurrentCUDAStream();\n if (x.scalar_type() == at::kFloat && kernel_size == 8 && stride == 1 && padding == 4) {\n dim3 block(BLOCK);\n dim3 grid((Lout + TILE - 1) / TILE, rows);\n size_t smem = (TILE + K - 1) * sizeof(float);\n avgpool1d_k8s1p4_kernel<<<grid, block, smem, stream>>>(\n x.data_ptr<float>(), y.data_ptr<float>(), rows, L, Lout);\n } else {\n int threads = 256;\n long long total = (long long)rows * Lout;\n int blocks = (int)((total + threads - 1) / threads);\n if (blocks > 65535) blocks = 65535;\n avgpool1d_generic_kernel<<<blocks, threads, 0, stream>>>(\n x.data_ptr<float>(), y.data_ptr<float>(), rows, L, Lout, (int)kernel_size, (int)stride, (int)padding);\n }\n return y;\n}\n\"\"\"\n\n_ext = load_inline(\n name=\"avgpool1d_ext_v1\",\n cpp_sources=_cpp_src,\n cuda_sources=_cuda_src,\n functions=[\"avgpool1d_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, kernel_size: int, stride: int = 1, padding: int = 0):\n super(ModelNew, self).__init__()\n self.kernel_size = int(kernel_size)\n self.stride = int(stride)\n self.padding = int(padding)\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n return _ext.avgpool1d_forward(x, self.kernel_size, self.stride, self.padding)\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 3 · 1 tool call(s) · 65.1s turn · 65.1s LLM · 6,610 chars (click to expand)
turn 365.1s turn · 65.1s LLM · 0.0s tools
reasoning
**Evaluating computation strategy**
I’m considering optimizing computation by using shared rolling with vectorization, especially with a focus on block sizes. It looks like I can utilize 256 threads, producing four outputs, which gives a tile size of 1024. I think the approach would involve loading shared memory and calculating based on each thread's ID. First, I'd compute the sum of the first eight elements, then perform three rolling calculations, storing the results as float4. That could be an effective method!
**Evaluating performance optimizations**
I’m thinking about whether to use half precision or maintain exact tolerance with fp32. The idea of considering potential boundary separation for the kernel is interesting, especially since I might be able to create an interior kernel without branching for most tiles. However, there's the issue of overhead on all those loads, especially with 549 million loads. I'm curious if that could become significant. It's a tricky balancing act!
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>\ntorch::Tensor avgpool1d_forward(torch::Tensor x, int64_t kernel_size, int64_t stride, int64_t padding);\n\"\"\"\n\n_cuda_src = r\"\"\"\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\nstatic constexpr int BLOCK = 256;\nstatic constexpr int TILE = 8192;\nstatic constexpr int K = 8;\nstatic constexpr int P = 4;\nstatic constexpr int Q = TILE / BLOCK;\n\n__global__ __launch_bounds__(BLOCK, 4)\nvoid avgpool1d_k8s1p4_kernel(const float* __restrict__ x, float* __restrict__ y,\n int rows, int L, int Lout) {\n int tile = blockIdx.x;\n int row = blockIdx.y;\n int tid = threadIdx.x;\n int out0 = tile * TILE;\n int tile_count = Lout - out0;\n if (tile_count > TILE) tile_count = TILE;\n if (tile_count <= 0) return;\n\n extern __shared__ float sh[];\n int load_count = tile_count + K - 1;\n int start = out0 - P;\n const float* rowx = x + ((long long)row) * L;\n if (start >= 0 && start + load_count <= L) {\n for (int i = tid; i < load_count; i += BLOCK) {\n sh[i] = rowx[start + i];\n }\n } else {\n for (int i = tid; i < load_count; i += BLOCK) {\n int xi = start + i;\n float v = 0.0f;\n if ((unsigned)xi < (unsigned)L) v = rowx[xi];\n sh[i] = v;\n }\n }\n __syncthreads();\n\n float* rowy = y + ((long long)row) * Lout;\n for (int q = 0; q < Q; ++q) {\n int j = q * BLOCK + tid;\n if (j < tile_count) {\n float s = sh[j] + sh[j+1] + sh[j+2] + sh[j+3]\n + sh[j+4] + sh[j+5] + sh[j+6] + sh[j+7];\n rowy[out0 + j] = s * 0.125f;\n }\n }\n}\n\n__global__ void avgpool1d_generic_kernel(const float* __restrict__ x, float* __restrict__ y,\n int rows, int L, int Lout, int kernel_size,\n int stride, int padding) {\n long long total = (long long)rows * Lout;\n long long idx = (long long)blockIdx.x * blockDim.x + threadIdx.x;\n long long step = (long long)blockDim.x * gridDim.x;\n for (; idx < total; idx += step) {\n int o = (int)(idx % Lout);\n int row = (int)(idx / Lout);\n int start = o * stride - padding;\n float sum = 0.0f;\n const float* rowx = x + ((long long)row) * L;\n for (int k = 0; k < kernel_size; ++k) {\n int xi = start + k;\n if ((unsigned)xi < (unsigned)L) sum += rowx[xi];\n }\n y[idx] = sum / (float)kernel_size;\n }\n}\n\ntorch::Tensor avgpool1d_forward(torch::Tensor x, int64_t kernel_size, int64_t stride, int64_t padding) {\n const int64_t N = x.size(0);\n const int64_t C = x.size(1);\n const int64_t L64 = x.size(2);\n const int64_t Lout64 = (L64 + 2 * padding - kernel_size) / stride + 1;\n auto y = torch::empty({N, C, Lout64}, x.options());\n int rows = (int)(N * C);\n int L = (int)L64;\n int Lout = (int)Lout64;\n cudaStream_t stream = at::cuda::getCurrentCUDAStream();\n if (x.scalar_type() == at::kFloat && kernel_size == 8 && stride == 1 && padding == 4) {\n dim3 block(BLOCK);\n dim3 grid((Lout + TILE - 1) / TILE, rows);\n size_t smem = (TILE + K - 1) * sizeof(float);\n avgpool1d_k8s1p4_kernel<<<grid, block, smem, stream>>>(\n x.data_ptr<float>(), y.data_ptr<float>(), rows, L, Lout);\n } else {\n int threads = 256;\n long long total = (long long)rows * Lout;\n int blocks = (int)((total + threads - 1) / threads);\n if (blocks > 65535) blocks = 65535;\n avgpool1d_generic_kernel<<<blocks, threads, 0, stream>>>(\n x.data_ptr<float>(), y.data_ptr<float>(), rows, L, Lout, (int)kernel_size, (int)stride, (int)padding);\n }\n return y;\n}\n\"\"\"\n\n_ext = load_inline(\n name=\"avgpool1d_ext_v2\",\n cpp_sources=_cpp_src,\n cuda_sources=_cuda_src,\n functions=[\"avgpool1d_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, kernel_size: int, stride: int = 1, padding: int = 0):\n super(ModelNew, self).__init__()\n self.kernel_size = int(kernel_size)\n self.stride = int(stride)\n self.padding = int(padding)\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n return _ext.avgpool1d_forward(x, self.kernel_size, self.stride, self.padding)\n"
}
result ← compile_kernel (ok)
okcompile_kernel
compile_kernel PASSED: kernel compiled without errors.
turn 49.1s turn · 9.1s 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>\ntorch::Tensor avgpool1d_forward(torch::Tensor x, int64_t kernel_size, int64_t stride, int64_t padding);\n\"\"\"\n\n_cuda_src = r\"\"\"\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\nstatic constexpr int BLOCK = 256;\nstatic constexpr int TILE = 8192;\nstatic constexpr int K = 8;\nstatic constexpr int P = 4;\nstatic constexpr int Q = TILE / BLOCK;\n\n__global__ __launch_bounds__(BLOCK, 4)\nvoid avgpool1d_k8s1p4_kernel(const float* __restrict__ x, float* __restrict__ y,\n int rows, int L, int Lout) {\n int tile = blockIdx.x;\n int row = blockIdx.y;\n int tid = threadIdx.x;\n int out0 = tile * TILE;\n int tile_count = Lout - out0;\n if (tile_count > TILE) tile_count = TILE;\n if (tile_count <= 0) return;\n\n extern __shared__ float sh[];\n int load_count = tile_count + K - 1;\n int start = out0 - P;\n const float* rowx = x + ((long long)row) * L;\n if (start >= 0 && start + load_count <= L) {\n for (int i = tid; i < load_count; i += BLOCK) {\n sh[i] = rowx[start + i];\n }\n } else {\n for (int i = tid; i < load_count; i += BLOCK) {\n int xi = start + i;\n float v = 0.0f;\n if ((unsigned)xi < (unsigned)L) v = rowx[xi];\n sh[i] = v;\n }\n }\n __syncthreads();\n\n float* rowy = y + ((long long)row) * Lout;\n for (int q = 0; q < Q; ++q) {\n int j = q * BLOCK + tid;\n if (j < tile_count) {\n float s = sh[j] + sh[j+1] + sh[j+2] + sh[j+3]\n + sh[j+4] + sh[j+5] + sh[j+6] + sh[j+7];\n rowy[out0 + j] = s * 0.125f;\n }\n }\n}\n\n__global__ void avgpool1d_generic_kernel(const float* __restrict__ x, float* __restrict__ y,\n int rows, int L, int Lout, int kernel_size,\n int stride, int padding) {\n long long total = (long long)rows * Lout;\n long long idx = (long long)blockIdx.x * blockDim.x + threadIdx.x;\n long long step = (long long)blockDim.x * gridDim.x;\n for (; idx < total; idx += step) {\n int o = (int)(idx % Lout);\n int row = (int)(idx / Lout);\n int start = o * stride - padding;\n float sum = 0.0f;\n const float* rowx = x + ((long long)row) * L;\n for (int k = 0; k < kernel_size; ++k) {\n int xi = start + k;\n if ((unsigned)xi < (unsigned)L) sum += rowx[xi];\n }\n y[idx] = sum / (float)kernel_size;\n }\n}\n\ntorch::Tensor avgpool1d_forward(torch::Tensor x, int64_t kernel_size, int64_t stride, int64_t padding) {\n const int64_t N = x.size(0);\n const int64_t C = x.size(1);\n const int64_t L64 = x.size(2);\n const int64_t Lout64 = (L64 + 2 * padding - kernel_size) / stride + 1;\n auto y = torch::empty({N, C, Lout64}, x.options());\n int rows = (int)(N * C);\n int L = (int)L64;\n int Lout = (int)Lout64;\n cudaStream_t stream = at::cuda::getCurrentCUDAStream();\n if (x.scalar_type() == at::kFloat && kernel_size == 8 && stride == 1 && padding == 4) {\n dim3 block(BLOCK);\n dim3 grid((Lout + TILE - 1) / TILE, rows);\n size_t smem = (TILE + K - 1) * sizeof(float);\n avgpool1d_k8s1p4_kernel<<<grid, block, smem, stream>>>(\n x.data_ptr<float>(), y.data_ptr<float>(), rows, L, Lout);\n } else {\n int threads = 256;\n long long total = (long long)rows * Lout;\n int blocks = (int)((total + threads - 1) / threads);\n if (blocks > 65535) blocks = 65535;\n avgpool1d_generic_kernel<<<blocks, threads, 0, stream>>>(\n x.data_ptr<float>(), y.data_ptr<float>(), rows, L, Lout, (int)kernel_size, (int)stride, (int)padding);\n }\n return y;\n}\n\"\"\"\n\n_ext = load_inline(\n name=\"avgpool1d_ext_v2\",\n cpp_sources=_cpp_src,\n cuda_sources=_cuda_src,\n functions=[\"avgpool1d_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, kernel_size: int, stride: int = 1, padding: int = 0):\n super(ModelNew, self).__init__()\n self.kernel_size = int(kernel_size)\n self.stride = int(stride)\n self.padding = int(padding)\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n return _ext.avgpool1d_forward(x, self.kernel_size, self.stride, self.padding)\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 59.6s turn · 9.6s LLM · 0.0s tools
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>\ntorch::Tensor avgpool1d_forward(torch::Tensor x, int64_t kernel_size, int64_t stride, int64_t padding);\n\"\"\"\n\n_cuda_src = r\"\"\"\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\nstatic constexpr int BLOCK = 256;\nstatic constexpr int TILE = 8192;\nstatic constexpr int K = 8;\nstatic constexpr int P = 4;\nstatic constexpr int Q = TILE / BLOCK;\n\n__global__ __launch_bounds__(BLOCK, 4)\nvoid avgpool1d_k8s1p4_kernel(const float* __restrict__ x, float* __restrict__ y,\n int rows, int L, int Lout) {\n int tile = blockIdx.x;\n int row = blockIdx.y;\n int tid = threadIdx.x;\n int out0 = tile * TILE;\n int tile_count = Lout - out0;\n if (tile_count > TILE) tile_count = TILE;\n if (tile_count <= 0) return;\n\n extern __shared__ float sh[];\n int load_count = tile_count + K - 1;\n int start = out0 - P;\n const float* rowx = x + ((long long)row) * L;\n if (start >= 0 && start + load_count <= L) {\n for (int i = tid; i < load_count; i += BLOCK) {\n sh[i] = rowx[start + i];\n }\n } else {\n for (int i = tid; i < load_count; i += BLOCK) {\n int xi = start + i;\n float v = 0.0f;\n if ((unsigned)xi < (unsigned)L) v = rowx[xi];\n sh[i] = v;\n }\n }\n __syncthreads();\n\n float* rowy = y + ((long long)row) * Lout;\n for (int q = 0; q < Q; ++q) {\n int j = q * BLOCK + tid;\n if (j < tile_count) {\n float s = sh[j] + sh[j+1] + sh[j+2] + sh[j+3]\n + sh[j+4] + sh[j+5] + sh[j+6] + sh[j+7];\n rowy[out0 + j] = s * 0.125f;\n }\n }\n}\n\n__global__ void avgpool1d_generic_kernel(const float* __restrict__ x, float* __restrict__ y,\n int rows, int L, int Lout, int kernel_size,\n int stride, int padding) {\n long long total = (long long)rows * Lout;\n long long idx = (long long)blockIdx.x * blockDim.x + threadIdx.x;\n long long step = (long long)blockDim.x * gridDim.x;\n for (; idx < total; idx += step) {\n int o = (int)(idx % Lout);\n int row = (int)(idx / Lout);\n int start = o * stride - padding;\n float sum = 0.0f;\n const float* rowx = x + ((long long)row) * L;\n for (int k = 0; k < kernel_size; ++k) {\n int xi = start + k;\n if ((unsigned)xi < (unsigned)L) sum += rowx[xi];\n }\n y[idx] = sum / (float)kernel_size;\n }\n}\n\ntorch::Tensor avgpool1d_forward(torch::Tensor x, int64_t kernel_size, int64_t stride, int64_t padding) {\n const int64_t N = x.size(0);\n const int64_t C = x.size(1);\n const int64_t L64 = x.size(2);\n const int64_t Lout64 = (L64 + 2 * padding - kernel_size) / stride + 1;\n auto y = torch::empty({N, C, Lout64}, x.options());\n int rows = (int)(N * C);\n int L = (int)L64;\n int Lout = (int)Lout64;\n cudaStream_t stream = at::cuda::getCurrentCUDAStream();\n if (x.scalar_type() == at::kFloat && kernel_size == 8 && stride == 1 && padding == 4) {\n dim3 block(BLOCK);\n dim3 grid((Lout + TILE - 1) / TILE, rows);\n size_t smem = (TILE + K - 1) * sizeof(float);\n avgpool1d_k8s1p4_kernel<<<grid, block, smem, stream>>>(\n x.data_ptr<float>(), y.data_ptr<float>(), rows, L, Lout);\n } else {\n int threads = 256;\n long long total = (long long)rows * Lout;\n int blocks = (int)((total + threads - 1) / threads);\n if (blocks > 65535) blocks = 65535;\n avgpool1d_generic_kernel<<<blocks, threads, 0, stream>>>(\n x.data_ptr<float>(), y.data_ptr<float>(), rows, L, Lout, (int)kernel_size, (int)stride, (int)padding);\n }\n return y;\n}\n\"\"\"\n\n_ext = load_inline(\n name=\"avgpool1d_ext_v2\",\n cpp_sources=_cpp_src,\n cuda_sources=_cuda_src,\n functions=[\"avgpool1d_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, kernel_size: int, stride: int = 1, padding: int = 0):\n super(ModelNew, self).__init__()\n self.kernel_size = int(kernel_size)\n self.stride = int(stride)\n self.padding = int(padding)\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n return _ext.avgpool1d_forward(x, self.kernel_size, self.stride, self.padding)\n"
}
result ← static_check (ok)
okstatic_check
static_check PASSED: no violations or warnings detected.
turn 614.9s turn · 14.9s LLM · 0.0s tools · FINAL
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>\ntorch::Tensor avgpool1d_forward(torch::Tensor x, int64_t kernel_size, int64_t stride, int64_t padding);\n\"\"\"\n\n_cuda_src = r\"\"\"\n#include <torch/extension.h>\n#include <ATen/cuda/CUDAContext.h>\n#include <cuda.h>\n#include <cuda_runtime.h>\n\nstatic constexpr int BLOCK = 256;\nstatic constexpr int TILE = 8192;\nstatic constexpr int K = 8;\nstatic constexpr int P = 4;\nstatic constexpr int Q = TILE / BLOCK;\n\n__global__ __launch_bounds__(BLOCK, 4)\nvoid avgpool1d_k8s1p4_kernel(const float* __restrict__ x, float* __restrict__ y,\n int rows, int L, int Lout) {\n int tile = blockIdx.x;\n int row = blockIdx.y;\n int tid = threadIdx.x;\n int out0 = tile * TILE;\n int tile_count = Lout - out0;\n if (tile_count > TILE) tile_count = TILE;\n if (tile_count <= 0) return;\n\n extern __shared__ float sh[];\n int load_count = tile_count + K - 1;\n int start = out0 - P;\n const float* rowx = x + ((long long)row) * L;\n if (start >= 0 && start + load_count <= L) {\n for (int i = tid; i < load_count; i += BLOCK) {\n sh[i] = rowx[start + i];\n }\n } else {\n for (int i = tid; i < load_count; i += BLOCK) {\n int xi = start + i;\n float v = 0.0f;\n if ((unsigned)xi < (unsigned)L) v = rowx[xi];\n sh[i] = v;\n }\n }\n __syncthreads();\n\n float* rowy = y + ((long long)row) * Lout;\n for (int q = 0; q < Q; ++q) {\n int j = q * BLOCK + tid;\n if (j < tile_count) {\n float s = sh[j] + sh[j+1] + sh[j+2] + sh[j+3]\n + sh[j+4] + sh[j+5] + sh[j+6] + sh[j+7];\n rowy[out0 + j] = s * 0.125f;\n }\n }\n}\n\n__global__ void avgpool1d_generic_kernel(const float* __restrict__ x, float* __restrict__ y,\n int rows, int L, int Lout, int kernel_size,\n int stride, int padding) {\n long long total = (long long)rows * Lout;\n long long idx = (long long)blockIdx.x * blockDim.x + threadIdx.x;\n long long step = (long long)blockDim.x * gridDim.x;\n for (; idx < total; idx += step) {\n int o = (int)(idx % Lout);\n int row = (int)(idx / Lout);\n int start = o * stride - padding;\n float sum = 0.0f;\n const float* rowx = x + ((long long)row) * L;\n for (int k = 0; k < kernel_size; ++k) {\n int xi = start + k;\n if ((unsigned)xi < (unsigned)L) sum += rowx[xi];\n }\n y[idx] = sum / (float)kernel_size;\n }\n}\n\ntorch::Tensor avgpool1d_forward(torch::Tensor x, int64_t kernel_size, int64_t stride, int64_t padding) {\n const int64_t N = x.size(0);\n const int64_t C = x.size(1);\n const int64_t L64 = x.size(2);\n const int64_t Lout64 = (L64 + 2 * padding - kernel_size) / stride + 1;\n auto y = torch::empty({N, C, Lout64}, x.options());\n int rows = (int)(N * C);\n int L = (int)L64;\n int Lout = (int)Lout64;\n cudaStream_t stream = at::cuda::getCurrentCUDAStream();\n if (x.scalar_type() == at::kFloat && kernel_size == 8 && stride == 1 && padding == 4) {\n dim3 block(BLOCK);\n dim3 grid((Lout + TILE - 1) / TILE, rows);\n size_t smem = (TILE + K - 1) * sizeof(float);\n avgpool1d_k8s1p4_kernel<<<grid, block, smem, stream>>>(\n x.data_ptr<float>(), y.data_ptr<float>(), rows, L, Lout);\n } else {\n int threads = 256;\n long long total = (long long)rows * Lout;\n int blocks = (int)((total + threads - 1) / threads);\n if (blocks > 65535) blocks = 65535;\n avgpool1d_generic_kernel<<<blocks, threads, 0, stream>>>(\n x.data_ptr<float>(), y.data_ptr<float>(), rows, L, Lout, (int)kernel_size, (int)stride, (int)padding);\n }\n return y;\n}\n\"\"\"\n\n_ext = load_inline(\n name=\"avgpool1d_ext_final\",\n cpp_sources=_cpp_src,\n cuda_sources=_cuda_src,\n functions=[\"avgpool1d_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, kernel_size: int, stride: int = 1, padding: int = 0):\n super(ModelNew, self).__init__()\n self.kernel_size = int(kernel_size)\n self.stride = int(stride)\n self.padding = int(padding)\n\n def forward(self, x: torch.Tensor) -> torch.Tensor:\n return _ext.avgpool1d_forward(x, self.kernel_size, self.stride, self.padding)\n"
}
result ← submit_kernel (ok)
oksubmit_kernel
submit_kernel PASSED: (5 / 5) correctness trials all passed.
Kernel runtime: 1.76 μs
Runtime stats: mean=1.76μs median=?μs std=0.15μs
Numerical precision: max_abs_err=0.00e+00 mean_abs_err=0.00e+00 max_rel_err=0.00e+00
Memory: 4129.03 MB peak (ref: 4129.03 MB, ratio: 1.0x)
Kernel launches: 1 (ref: 6, fusion ratio: 6.0)
Energy: 377.06 mJ/run (ref: 1132.16 mJ, ratio: 3.0026x)
SOL score: 0.7483
final submitted kernel (129 lines, 4,539 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 avgpool1d_forward(torch::Tensor x, int64_t kernel_size, int64_t stride, int64_t padding);
"""
_cuda_src = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda.h>
#include <cuda_runtime.h>
static constexpr int BLOCK = 256;
static constexpr int TILE = 8192;
static constexpr int K = 8;
static constexpr int P = 4;
static constexpr int Q = TILE / BLOCK;
__global__ __launch_bounds__(BLOCK, 4)
void avgpool1d_k8s1p4_kernel(const float* __restrict__ x, float* __restrict__ y,
int rows, int L, int Lout) {
int tile = blockIdx.x;
int row = blockIdx.y;
int tid = threadIdx.x;
int out0 = tile * TILE;
int tile_count = Lout - out0;
if (tile_count > TILE) tile_count = TILE;
if (tile_count <= 0) return;
extern __shared__ float sh[];
int load_count = tile_count + K - 1;
int start = out0 - P;
const float* rowx = x + ((long long)row) * L;
if (start >= 0 && start + load_count <= L) {
for (int i = tid; i < load_count; i += BLOCK) {
sh[i] = rowx[start + i];
}
} else {
for (int i = tid; i < load_count; i += BLOCK) {
int xi = start + i;
float v = 0.0f;
if ((unsigned)xi < (unsigned)L) v = rowx[xi];
sh[i] = v;
}
}
__syncthreads();
float* rowy = y + ((long long)row) * Lout;
for (int q = 0; q < Q; ++q) {
int j = q * BLOCK + tid;
if (j < tile_count) {
float s = sh[j] + sh[j+1] + sh[j+2] + sh[j+3]
+ sh[j+4] + sh[j+5] + sh[j+6] + sh[j+7];
rowy[out0 + j] = s * 0.125f;
}
}
}
__global__ void avgpool1d_generic_kernel(const float* __restrict__ x, float* __restrict__ y,
int rows, int L, int Lout, int kernel_size,
int stride, int padding) {
long long total = (long long)rows * Lout;
long long idx = (long long)blockIdx.x * blockDim.x + threadIdx.x;
long long step = (long long)blockDim.x * gridDim.x;
for (; idx < total; idx += step) {
int o = (int)(idx % Lout);
int row = (int)(idx / Lout);
int start = o * stride - padding;
float sum = 0.0f;
const float* rowx = x + ((long long)row) * L;
for (int k = 0; k < kernel_size; ++k) {
int xi = start + k;
if ((unsigned)xi < (unsigned)L) sum += rowx[xi];
}
y[idx] = sum / (float)kernel_size;
}
}
torch::Tensor avgpool1d_forward(torch::Tensor x, int64_t kernel_size, int64_t stride, int64_t padding) {
const int64_t N = x.size(0);
const int64_t C = x.size(1);
const int64_t L64 = x.size(2);
const int64_t Lout64 = (L64 + 2 * padding - kernel_size) / stride + 1;
auto y = torch::empty({N, C, Lout64}, x.options());
int rows = (int)(N * C);
int L = (int)L64;
int Lout = (int)Lout64;
cudaStream_t stream = at::cuda::getCurrentCUDAStream();
if (x.scalar_type() == at::kFloat && kernel_size == 8 && stride == 1 && padding == 4) {
dim3 block(BLOCK);
dim3 grid((Lout + TILE - 1) / TILE, rows);
size_t smem = (TILE + K - 1) * sizeof(float);
avgpool1d_k8s1p4_kernel<<<grid, block, smem, stream>>>(
x.data_ptr<float>(), y.data_ptr<float>(), rows, L, Lout);
} else {
int threads = 256;
long long total = (long long)rows * Lout;
int blocks = (int)((total + threads - 1) / threads);
if (blocks > 65535) blocks = 65535;
avgpool1d_generic_kernel<<<blocks, threads, 0, stream>>>(
x.data_ptr<float>(), y.data_ptr<float>(), rows, L, Lout, (int)kernel_size, (int)stride, (int)padding);
}
return y;
}
"""
_ext = load_inline(
name="avgpool1d_ext_final",
cpp_sources=_cpp_src,
cuda_sources=_cuda_src,
functions=["avgpool1d_forward"],
extra_cflags=["-O3"],
extra_cuda_cflags=["-O3", "--use_fast_math"],
verbose=False,
)
class ModelNew(nn.Module):
def __init__(self, kernel_size: int, stride: int = 1, padding: int = 0):
super(ModelNew, self).__init__()
self.kernel_size = int(kernel_size)
self.stride = int(stride)
self.padding = int(padding)
def forward(self, x: torch.Tensor) -> torch.Tensor:
return _ext.avgpool1d_forward(x, self.kernel_size, self.stride, self.padding)