← Back to Leaderboard

The AI CUDA Engineer 👷

76_Gemm_Add_ReLU • vectorized_warp_reduction_base

Level 2 • Task 76
import torch
import torch.nn as nn
import torch.nn.functional as F


def module_fn(
    x: torch.Tensor,
    weight: torch.Tensor,
    bias: torch.Tensor,
) -> torch.Tensor:
    """
    Performs matrix multiplication, adds bias, and applies ReLU activation.

    Args:
        x (torch.Tensor): Input tensor with shape (batch_size, in_features)
        weight (torch.Tensor): Weight matrix with shape (out_features, in_features)
        bias (torch.Tensor): Bias tensor with shape (out_features,)

    Returns:
        torch.Tensor: Output tensor with shape (batch_size, out_features)
    """
    x = F.linear(x, weight)
    x = x + bias
    x = F.relu(x)
    return x


class Model(nn.Module):
    """
    Simple model that performs a matrix multiplication, adds a bias term, and applies ReLU.
    """

    def __init__(self, in_features, out_features, bias_shape):
        super(Model, self).__init__()
        gemm = nn.Linear(in_features, out_features, bias=False)
        self.weight = nn.Parameter(gemm.weight)
        self.bias = nn.Parameter(torch.randn(bias_shape) * 0.02)

    def forward(self, x, fn=module_fn):
        return fn(x, self.weight, self.bias)


batch_size = 128
in_features = 1024
out_features = 512
bias_shape = (out_features,)


def get_inputs():
    return [torch.randn(batch_size, in_features)]


def get_init_inputs():
    return [in_features, out_features, bias_shape]

import torch
import torch.nn as nn

class Model(nn.Module):
    """
    Simple model that performs a matrix multiplication, adds a bias term, and applies ReLU.
    """
    def __init__(self, in_features, out_features, bias_shape):
        super(Model, self).__init__()
        self.gemm = nn.Linear(in_features, out_features, bias=False)
        self.bias = nn.Parameter(torch.randn(bias_shape)*0.02)

    def forward(self, x):   
        """
        Args:
            x (torch.Tensor): Input tensor with shape (batch_size, in_features).
        Returns:
            torch.Tensor: Output tensor with shape (batch_size, out_features).
        """
        x = self.gemm(x)
        x = x + self.bias
        x = torch.relu(x)
        return x

batch_size = 128
in_features = 1024
out_features = 512
bias_shape = (out_features,)

def get_inputs():
    return [torch.randn(batch_size, in_features)]

def get_init_inputs():
    return [in_features, out_features, bias_shape]
Download Evaluation Download PyTorch Download CUDA Download Profiles

Kernel Information

Operation Name 76_Gemm_Add_ReLU
Level ID 2
Task ID 76
Kernel Name vectorized_warp_reduction_base
CUDA Speedup (Native) 0.671x
CUDA Speedup (Compile) 1.118x
CUDA Runtime 0.040 ms
PyTorch Runtime (Native) 0.027 ms
PyTorch Runtime (Compile) 0.045 ms
Correct True
Max Diff (vs. Reference) 0.000000
Model openrouter-deepseek-reasoner
Temperature 1.00
View Experiment Progress Details

Related Kernels (Level 2, Task 76 • 76_Gemm_Add_ReLU)

Rank Kernel Name Runtime (ms) Speedup Native Speedup Compile
🥇 shared_warp_tile_kernel_base 0.03 0.93 1.54
🥇 combined_warp_tile_base 0.03 0.93 1.54
🥉 optimized_block_size_kernel_base 0.03 0.89 1.49
4 warp_tile_ldg_base 0.03 0.87 1.44
4 even_workload_dist_base_base 0.03 0.87 1.44
4 hybrid_warp_tile_kernel_base 0.03 0.87 1.44
4 warp_tile_hybrid_base 0.03 0.87 1.44
8 warp_tile_ldg_opt_base 0.03 0.81 1.36
8 warp_reduction_optimized_base_base 0.03 0.81 1.36
10 optimized_shared_memory_base_base 0.03 0.79 1.32
10 warp_tile_base_base 0.03 0.79 1.32
12 hybrid_optimized_kernel_base 0.04 0.77 1.28
13 warp_reduction_gemm_base 0.04 0.71 1.18
13 warp_tile_aligned_base_base 0.04 0.71 1.18
15 vectorized_warp_unroll_base_base 0.04 0.69 1.15
15 vectorized_warp_unroll_base_edit_1 0.04 0.69 1.15
15 warp_reduction_unrolled_gemm_edit_1 0.04 0.69 1.15
18 unrolled_warp_gemm_edit_1 0.04 0.67 1.12
18 unrolled_warp_gemm_base 0.04 0.67 1.12
18 vectorized_warp_reduction_base 0.04 0.67 1.12 
#include <torch/extension.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <c10/cuda/CUDAGuard.h>

__global__ void linear_relu_vector_warp_kernel(const float* x, const float* weight, const float* bias, float* out,
                                               int batch_size, int in_features, int out_features) {
  // Global warp index determines which output element we compute
  int warp_id = (blockIdx.x * blockDim.x + threadIdx.x) / 32;
  int lane_id = threadIdx.x % 32;

  if (warp_id >= batch_size * out_features) return;

  int row = warp_id / out_features;
  int col = warp_id % out_features;

  float sum = 0.0f;

  const float* x_row = x + row * in_features;
  const float* w_row = weight + col * in_features;

  // Process aligned sections with vectorized loads
  int vec_end = (in_features / 4) * 4;
  for (int k = lane_id * 4; k < vec_end; k += 32 * 4) {
    float4 x_val = *reinterpret_cast<const float4*>(x_row + k);
    float4 w_val = *reinterpret_cast<const float4*>(w_row + k);
    sum += x_val.x * w_val.x + x_val.y * w_val.y + 
           x_val.z * w_val.z + x_val.w * w_val.w;
  }

  // Process remaining elements with scalar loads
  for (int k = vec_end + lane_id; k < in_features; k += 32) {
    sum += x_row[k] * w_row[k];
  }

  // Warp-level reduction
  for (int offset = 16; offset > 0; offset /= 2) {
    sum += __shfl_down_sync(0xffffffff, sum, offset);
  }

  if (lane_id == 0) {
    sum += bias[col];
    out[row * out_features + col] = fmaxf(sum, 0.0f);
  }
}

torch::Tensor linear_relu_forward(torch::Tensor x, torch::Tensor weight, torch::Tensor bias) {
  TORCH_CHECK(x.is_cuda(), "x must be CUDA tensor");
  TORCH_CHECK(weight.is_cuda(), "weight must be CUDA tensor");
  TORCH_CHECK(bias.is_cuda(), "bias must be CUDA tensor");

  const int batch_size = x.size(0);
  const int in_features = x.size(1);
  const int out_features = weight.size(0);

  auto out = torch::empty({batch_size, out_features}, x.options());

  // Each warp handles one output element
  int total_warps = batch_size * out_features;
  int total_threads = total_warps * 32;
  const int threads_per_block = 256;
  int blocks = (total_threads + threads_per_block - 1) / threads_per_block;

  cudaStream_t stream = c10::cuda::getCurrentCUDAStream();
  linear_relu_vector_warp_kernel<<<blocks, threads_per_block, 0, stream>>>(
      x.data_ptr<float>(),
      weight.data_ptr<float>(),
      bias.data_ptr<float>(),
      out.data_ptr<float>(),
      batch_size,
      in_features,
      out_features
  );

  return out;
}

PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
  m.def("forward", &linear_relu_forward, "Vectorized warp reduction GEMM+Bias+ReLU");
}
Performance Metrics
Metric Value Unit Variance Samples
Executed Ipc Active 1.564 inst/cycle 0.000 5
Executed Ipc Elapsed 1.438 inst/cycle 0.000 5
Issue Slots Busy 39.190 % 0.200 5
Issued Ipc Active 1.566 inst/cycle 0.000 5
SM Busy 39.190 % 0.200 5
Memory Throughput 67773237745.008 byte/second 122398151738564656.000 5
Mem Busy 59.074 % 0.092 5
Max Bandwidth 70.832 % 0.149 5
L1/TEX Hit Rate 48.982 % 0.013 5
L2 Hit Rate 97.894 % 0.576 5
Mem Pipes Busy 19.514 % 0.010 5
Warp Cycles Per Issued Instruction 28.202 cycle 0.020 5
Warp Cycles Per Executed Instruction 28.252 cycle 0.020 5
Avg. Active Threads Per Warp 30.070 0.000 5
Avg. Not Predicated Off Threads Per Warp 28.440 0.000 5
Max Active Clusters 0.000 cluster 0.000 5
Max Cluster Size 8.000 block 0.000 5
Overall GPU Occupancy 0.000 % 0.000 5
Cluster Occupancy 0.000 % 0.000 5
Block Limit SM 32.000 block 0.000 5
Block Limit Registers 6.000 block 0.000 5
Block Limit Shared Mem 32.000 block 0.000 5
Block Limit Warps 8.000 block 0.000 5
Theoretical Active Warps per SM 48.000 warp 0.000 5
Theoretical Occupancy 75.000 % 0.000 5
Achieved Occupancy 69.582 % 0.023 5
Achieved Active Warps Per SM 44.534 warp 0.010 5
Analysis Rules
Rule Description
INF HighPipeUtilization ALU is the highest-utilized pipeline (27.1%) based on active cycles, taking into account the rates of its different instructions. It executes integer and logic operations. It is well-utilized, but should not be a bottleneck.
INF CPIStall Check the Warp Stall Sampling (All Cycles) table for the top stall locations in your source based on sampling data. The Kernel Profiling Guide (https://docs.nvidia.com/nsight-compute/ProfilingGuide/index.html#metrics-reference) provides more details on each stall reason.
WRN Occupancy This kernel's theoretical occupancy (75.0%) is limited by the number of required registers. See the CUDA Best Practices Guide (https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#occupancy) for more details on optimizing occupancy.
Operation / Metric Value Unit
aten::to
CPU Time 482080.26 μs
Device Time 173.89 μs
Self CPU Time 62.44 μs
Self Device Time 0.00 μs
CPU Memory Usage 0 B
Device Memory Usage 0 B
Self CPU Memory Usage 0 B
Self Device Memory Usage 0 B
aten::_to_copy
CPU Time 482017.83 μs
Device Time 173.89 μs
Self CPU Time 113.62 μs
Self Device Time 0.00 μs
CPU Memory Usage 0 B
Device Memory Usage 0 B
Self CPU Memory Usage 0 B
Self Device Memory Usage 0 B
aten::empty_strided
CPU Time 481411.10 μs
Device Time 0.00 μs
Self CPU Time 122.98 μs
Self Device Time 0.00 μs
CPU Memory Usage 0 B
Device Memory Usage 0 B
Self CPU Memory Usage 0 B
Self Device Memory Usage 0 B
cudaDeviceGetStreamPriorityRange
CPU Time 469799.13 μs
Device Time 0.00 μs
Self CPU Time 469799.13 μs
Self Device Time 0.00 μs
CPU Memory Usage 0 B
Device Memory Usage 0 B
Self CPU Memory Usage 0 B
Self Device Memory Usage 0 B
cudaLaunchKernel
CPU Time 706369.62 μs
Device Time 16186.82 μs
Self CPU Time 706369.62 μs
Self Device Time 16186.82 μs
CPU Memory Usage 0 B
Device Memory Usage 0 B
Self CPU Memory Usage 0 B
Self Device Memory Usage 0 B
linear_relu_vector_warp_kernel(float const*, float const*, float const*, float*, int, int, int)
CPU Time 0.00 μs
Device Time 284559.11 μs
Self CPU Time 0.00 μs
Self Device Time 284559.11 μs
CPU Memory Usage 0 B
Device Memory Usage 0 B
Self CPU Memory Usage 0 B
Self Device Memory Usage 0 B
cudaEventRecord
CPU Time 21047.77 μs
Device Time 31436.41 μs
Self CPU Time 21047.77 μs
Self Device Time 31436.41 μs
CPU Memory Usage 0 B
Device Memory Usage 0 B
Self CPU Memory Usage 0 B
Self Device Memory Usage 0 B
aten::zero_
CPU Time 244831.75 μs
Device Time 606805.38 μs
Self CPU Time 13376.65 μs
Self Device Time 0.00 μs
CPU Memory Usage 0 B
Device Memory Usage 0 B
Self CPU Memory Usage 0 B
Self Device Memory Usage 0 B
aten::fill_
CPU Time 231456.47 μs
Device Time 606805.38 μs
Self CPU Time 17352.55 μs
Self Device Time 606805.38 μs
CPU Memory Usage 0 B
Device Memory Usage 0 B
Self CPU Memory Usage 0 B
Self Device Memory Usage 0 B
void at::native::vectorized_elementwise_kernel<4, at::native::FillFunctor<int>, at::detail::Array<char*, 1> >(int, at::native::FillFunctor<int>, at::detail::Array<char*, 1>)
CPU Time 0.00 μs
Device Time 606805.38 μs
Self CPU Time 0.00 μs
Self Device Time 606805.38 μs
CPU Memory Usage 0 B
Device Memory Usage 0 B
Self CPU Memory Usage 0 B
Self Device Memory Usage 0 B
Status: Completed 
45312 warnings generated when compiling for host.
Suppressed 45347 warnings (45300 in non-user code, 47 NOLINT).
Use -header-filter=.* to display errors from all non-system headers. Use -system-headers to display errors from system headers as well.
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_2/task_76/b3_s3_vectorized_warp_reduction/base/base.cu:6:48 bugprone-easily-swappable-parameters
6 | __global__ void linear_relu_vector_warp_kernel(const float* x, const float* weight, const float* bias, float* out,
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_2/task_76/b3_s3_vectorized_warp_reduction/base/base.cu:6:61: note: the first parameter in the range is 'x'
6 | __global__ void linear_relu_vector_warp_kernel(const float* x, const float* weight, const float* bias, float* out,
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_2/task_76/b3_s3_vectorized_warp_reduction/base/base.cu:6:98: note: the last parameter in the range is 'bias'
6 | __global__ void linear_relu_vector_warp_kernel(const float* x, const float* weight, const float* bias, float* out,
| ^~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_2/task_76/b3_s3_vectorized_warp_reduction/base/base.cu:7:48: warning: 2 adjacent parameters of 'linear_relu_vector_warp_kernel' of similar type ('int') are easily swapped by mistake [bugprone-easily-swappable-parameters]
7 | int batch_size, int in_features, int out_features) {
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_2/task_76/b3_s3_vectorized_warp_reduction/base/base.cu:7:52: note: the first parameter in the range is 'batch_size'
7 | int batch_size, int in_features, int out_features) {
| ^~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_2/task_76/b3_s3_vectorized_warp_reduction/base/base.cu:7:68: note: the last parameter in the range is 'in_features'
7 | int batch_size, int in_features, int out_features) {
| ^~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_2/task_76/b3_s3_vectorized_warp_reduction/base/base.cu:9:17: warning: narrowing conversion from 'unsigned int' to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
9 | int warp_id = (blockIdx.x * blockDim.x + threadIdx.x) / 32;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_2/task_76/b3_s3_vectorized_warp_reduction/base/base.cu:10:17: warning: narrowing conversion from 'unsigned int' to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
10 | int lane_id = threadIdx.x % 32;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_2/task_76/b3_s3_vectorized_warp_reduction/base/base.cu:19:24: warning: result of multiplication in type 'int' is used as a pointer offset after an implicit widening conversion to type 'ptrdiff_t' [bugprone-implicit-widening-of-multiplication-result]
19 | const float* x_row = x + row * in_features;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_2/task_76/b3_s3_vectorized_warp_reduction/base/base.cu:19:28: note: make conversion explicit to silence this warning
5 |
6 | __global__ void linear_relu_vector_warp_kernel(const float* x, const float* weight, const float* bias, float* out,
7 | int batch_size, int in_features, int out_features) {
8 | // Global warp index determines which output element we compute
9 | int warp_id = (blockIdx.x * blockDim.x + threadIdx.x) / 32;
10 | int lane_id = threadIdx.x % 32;
11 |
12 | if (warp_id >= batch_size * out_features) return;
13 |
14 | int row = warp_id / out_features;
15 | int col = warp_id % out_features;
16 |
17 | float sum = 0.0f;
18 |
19 | const float* x_row = x + row * in_features;
| ^~~~~~~~~~~~~~~~~
| static_cast<ptrdiff_t>( )
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_2/task_76/b3_s3_vectorized_warp_reduction/base/base.cu:19:28: note: perform multiplication in a wider type
19 | const float* x_row = x + row * in_features;
| ^~~
| static_cast<ptrdiff_t>( )
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_2/task_76/b3_s3_vectorized_warp_reduction/base/base.cu:20:24: warning: result of multiplication in type 'int' is used as a pointer offset after an implicit widening conversion to type 'ptrdiff_t' [bugprone-implicit-widening-of-multiplication-result]
20 | const float* w_row = weight + col * in_features;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_2/task_76/b3_s3_vectorized_warp_reduction/base/base.cu:20:33: note: make conversion explicit to silence this warning
20 | const float* w_row = weight + col * in_features;
| ^~~~~~~~~~~~~~~~~
| static_cast<ptrdiff_t>( )
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_2/task_76/b3_s3_vectorized_warp_reduction/base/base.cu:20:33: note: perform multiplication in a wider type
20 | const float* w_row = weight + col * in_features;
| ^~~
| static_cast<ptrdiff_t>( )
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_2/task_76/b3_s3_vectorized_warp_reduction/base/base.cu:47:49: warning: the parameter 'x' is copied for each invocation but only used as a const reference; consider making it a const reference [performance-unnecessary-value-param]
47 | torch::Tensor linear_relu_forward(torch::Tensor x, torch::Tensor weight, torch::Tensor bias) {
| ^
| const &
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_2/task_76/b3_s3_vectorized_warp_reduction/base/base.cu:47:66: warning: the parameter 'weight' is copied for each invocation but only used as a const reference; consider making it a const reference [performance-unnecessary-value-param]
47 | torch::Tensor linear_relu_forward(torch::Tensor x, torch::Tensor weight, torch::Tensor bias) {
| ^
| const &
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_2/task_76/b3_s3_vectorized_warp_reduction/base/base.cu:47:88: warning: the parameter 'bias' is copied for each invocation but only used as a const reference; consider making it a const reference [performance-unnecessary-value-param]
47 | torch::Tensor linear_relu_forward(torch::Tensor x, torch::Tensor weight, torch::Tensor bias) {
| ^
| const &
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_2/task_76/b3_s3_vectorized_warp_reduction/base/base.cu:52:26: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
52 | const int batch_size = x.size(0);
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_2/task_76/b3_s3_vectorized_warp_reduction/base/base.cu:53:27: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
53 | const int in_features = x.size(1);
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250212_optimize_b5_s4_e1_v2/level_2/task_76/b3_s3_vectorized_warp_reduction/base/base.cu:54:28: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
54 | const int out_features = weight.size(0);
| ^