← Back to Leaderboard

The AI CUDA Engineer 👷

76_Gemm_Add_ReLUeven_workload_dist_base_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]

Kernel Information

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>

#define WARP_SIZE 32
#define TILE_SIZE 4  // Each warp processes TILE_SIZE output elements

// Kernel to distribute workloads evenly, ensuring balanced computation
// across warps by dynamically mapping workloads to threads based on output size
__global__ void even_workload_distribution_kernel(const float* __restrict__ x,
                                                    const float* __restrict__ weight,
                                                    const float* __restrict__ bias,
                                                    float* __restrict__ out,
                                                    int in_features,
                                                    int out_features) {
    // Each block processes a batch sample (blockIdx.x) and a set of outputs (blockIdx.y)
    int batch_idx = blockIdx.x;

    // Warp and lane indices for work distribution
    int warp_id = threadIdx.x / WARP_SIZE;
    int lane_id = threadIdx.x % WARP_SIZE;

    // Total warps in current block, dynamic to support varied block sizes
    int warps_per_block = blockDim.x / WARP_SIZE;

    // Output base derived from block, warp-level offset yields load balancing
    int base_out = blockIdx.y * warps_per_block * TILE_SIZE + warp_id * TILE_SIZE;
    if (base_out >= out_features) return;

    // Pointer to the batch's input row
    const float* x_row = x + batch_idx * in_features;

    // Pre-allocate accumulator for each tile output
    float sums[TILE_SIZE] = {0.0f, 0.0f, 0.0f, 0.0f};

    // Vectorized load setup
    int nvec = in_features / 4;  // 128-bit alignment
    int rem = in_features % 4;  // Remainder handling

    // Processing each output in the warp's designated tile
    #pragma unroll
    for (int tile = 0; tile < TILE_SIZE; tile++) {
        int current_out = base_out + tile;
        if (current_out < out_features) {
            const float* w_row = weight + current_out * in_features;
            
            // Vectorized memory access: uses float4 for coalesced GPU loads
            const float4* x_vec = reinterpret_cast<const float4*>(x_row);
            const float4* w_vec = reinterpret_cast<const float4*>(w_row);

            float sum = 0.0f;

            // Vectorized computation loop
            for (int k = lane_id; k < nvec; k += WARP_SIZE) {
                float4 xv = __ldg(&x_vec[k]);
                float4 wv = __ldg(&w_vec[k]);
                sum += xv.x * wv.x + xv.y * wv.y + xv.z * wv.z + xv.w * wv.w;
            }

            // Handling remainder elements with unrolling
            int remainder_start = nvec * 4;
            for (int r = remainder_start + lane_id; r < in_features; r += WARP_SIZE) {
                sum += __ldg(&x_row[r]) * __ldg(&w_row[r]);
            }

            sums[tile] = sum;
        }
    }

    // Warp-level reduction for output sum
    #pragma unroll
    for (int tile = 0; tile < TILE_SIZE; tile++) {
        for (int offset = WARP_SIZE / 2; offset > 0; offset /= 2) {
            sums[tile] += __shfl_down_sync(0xffffffff, sums[tile], offset);
        }
    }

    // Lane 0 writes final result with bias and activation
    if (lane_id == 0) {
        for (int tile = 0; tile < TILE_SIZE; tile++) {
            int current_out = base_out + tile;
            if (current_out < out_features) {
                float result = sums[tile] + __ldg(bias + current_out);
                out[batch_idx * out_features + current_out] = (result > 0.0f) ? result : 0.0f;
            }
        }
    }
}

// Host launcher: Ensures balance by dynamically adjusting work distribution based
// on batch size and output dimensions

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

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

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

    // Adjust warps per block dynamically to ensure balanced computation
    const int warps_per_block = 8;
    const int threads_per_block = warps_per_block * WARP_SIZE;

    // Grid spacing along y-axis adjusted to cover all outputs
    int blocks_y = (out_features + (warps_per_block * TILE_SIZE) - 1) / (warps_per_block * TILE_SIZE);

    dim3 grid(batch_size, blocks_y);
    dim3 block(threads_per_block);

    cudaStream_t stream = c10::cuda::getCurrentCUDAStream();
    even_workload_distribution_kernel<<<grid, block, 0, stream>>>(
        x.data_ptr<float>(),
        weight.data_ptr<float>(),
        bias.data_ptr<float>(),
        out.data_ptr<float>(),
        in_features,
        out_features
    );

    return out;
}

PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
    m.def("forward", &even_workload_distribution_forward, "Even workload distribution (CUDA) with warp-level load balancing");
}
Performance Metrics
Metric Value Unit Variance Samples
Executed Ipc Active 1.554 inst/cycle 0.000 5
Executed Ipc Elapsed 1.372 inst/cycle 0.000 5
Issue Slots Busy 38.998 % 0.014 5
Issued Ipc Active 1.560 inst/cycle 0.000 5
SM Busy 38.998 % 0.014 5
Memory Throughput 81534652084.422 byte/second 177042702364717984.000 5
Mem Busy 70.364 % 0.162 5
Max Bandwidth 68.788 % 0.153 5
L1/TEX Hit Rate 57.928 % 0.139 5
L2 Hit Rate 94.534 % 49.132 5
Mem Pipes Busy 23.092 % 0.018 5
Warp Cycles Per Issued Instruction 24.044 cycle 0.002 5
Warp Cycles Per Executed Instruction 24.134 cycle 0.002 5
Avg. Active Threads Per Warp 30.460 0.000 5
Avg. Not Predicated Off Threads Per Warp 29.550 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 5.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 40.000 warp 0.000 5
Theoretical Occupancy 62.500 % 0.000 5
Achieved Occupancy 58.724 % 0.005 5
Achieved Active Warps Per SM 37.586 warp 0.002 5
Analysis Rules
Rule Description
INF HighPipeUtilization ALU is the highest-utilized pipeline (22.5%) 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 (62.5%) 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 472344.02 μs
Device Time 205.73 μs
Self CPU Time 60.25 μ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 472283.77 μs
Device Time 205.73 μs
Self CPU Time 132.35 μ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 471591.35 μs
Device Time 0.00 μs
Self CPU Time 156.63 μ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 468353.03 μs
Device Time 0.00 μs
Self CPU Time 468353.03 μ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 384633.12 μs
Device Time 548.35 μs
Self CPU Time 384633.12 μs
Self Device Time 548.35 μs
CPU Memory Usage 0 B
Device Memory Usage 0 B
Self CPU Memory Usage 0 B
Self Device Memory Usage 0 B
even_workload_distribution_kernel(float const*, float const*, float const*, float*, int, int)
CPU Time 0.00 μs
Device Time 176405.30 μs
Self CPU Time 0.00 μs
Self Device Time 176405.30 μ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 110878.11 μs
Device Time 496620.02 μs
Self CPU Time 18678.53 μ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 92201.54 μs
Device Time 496620.02 μs
Self CPU Time 25911.03 μs
Self Device Time 496620.02 μ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 496620.02 μs
Self CPU Time 0.00 μs
Self Device Time 496620.02 μ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
45316 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/20250203_optimize_b10_s4_e0_sweep/level_2/task_76/b9_s1_even_workload_dist_base/base/base.cu:11:51 bugprone-easily-swappable-parameters
11 | __global__ void even_workload_distribution_kernel(const float* __restrict__ x,
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~
12 | const float* __restrict__ weight,
| ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
13 | const float* __restrict__ bias,
| ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_76/b9_s1_even_workload_dist_base/base/base.cu:11:77: note: the first parameter in the range is 'x'
11 | __global__ void even_workload_distribution_kernel(const float* __restrict__ x,
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_76/b9_s1_even_workload_dist_base/base/base.cu:13:79: note: the last parameter in the range is 'bias'
13 | const float* __restrict__ bias,
| ^~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_76/b9_s1_even_workload_dist_base/base/base.cu:15:53: warning: 2 adjacent parameters of 'even_workload_distribution_kernel' of similar type ('int') are easily swapped by mistake [bugprone-easily-swappable-parameters]
15 | int in_features,
| ^~~~~~~~~~~~~~~~
16 | int out_features) {
| ~~~~~~~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_76/b9_s1_even_workload_dist_base/base/base.cu:15:57: note: the first parameter in the range is 'in_features'
15 | int in_features,
| ^~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_76/b9_s1_even_workload_dist_base/base/base.cu:16:57: note: the last parameter in the range is 'out_features'
16 | int out_features) {
| ^~~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_76/b9_s1_even_workload_dist_base/base/base.cu:18:21: warning: narrowing conversion from 'unsigned int' to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
18 | int batch_idx = blockIdx.x;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_76/b9_s1_even_workload_dist_base/base/base.cu:21:19: warning: narrowing conversion from 'unsigned int' to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
21 | int warp_id = threadIdx.x / WARP_SIZE;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_76/b9_s1_even_workload_dist_base/base/base.cu:22:19: warning: narrowing conversion from 'unsigned int' to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
22 | int lane_id = threadIdx.x % WARP_SIZE;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_76/b9_s1_even_workload_dist_base/base/base.cu:25:27: warning: narrowing conversion from 'unsigned int' to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
25 | int warps_per_block = blockDim.x / WARP_SIZE;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_76/b9_s1_even_workload_dist_base/base/base.cu:28:20: warning: narrowing conversion from 'unsigned int' to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
28 | int base_out = blockIdx.y * warps_per_block * TILE_SIZE + warp_id * TILE_SIZE;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_76/b9_s1_even_workload_dist_base/base/base.cu:32:26: 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]
32 | const float* x_row = x + batch_idx * in_features;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_76/b9_s1_even_workload_dist_base/base/base.cu:32:30: note: make conversion explicit to silence this warning
5 | const float* x_row = x + batch_idx * in_features;
| ^~~~~~~~~~~~~~~~~~~~~~~
| static_cast<ptrdiff_t>()
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_76/b9_s1_even_workload_dist_base/base/base.cu:32:30: note: perform multiplication in a wider type
32 | const float* x_row = x + batch_idx * in_features;
| ^~~~~~~~~
| static_cast<ptrdiff_t>()
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_76/b9_s1_even_workload_dist_base/base/base.cu:39:9: warning: Value stored to 'rem' during its initialization is never read [clang-analyzer-deadcode.DeadStores]
39 | int rem = in_features % 4; // Remainder handling
| ^~~ ~~~~~~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_76/b9_s1_even_workload_dist_base/base/base.cu:39:9: note: Value stored to 'rem' during its initialization is never read
39 | int rem = in_features % 4; // Remainder handling
| ^~~ ~~~~~~~~~~~~~~~
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_76/b9_s1_even_workload_dist_base/base/base.cu:46:34: 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]
46 | const float* w_row = weight + current_out * in_features;
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_76/b9_s1_even_workload_dist_base/base/base.cu:46:43: note: make conversion explicit to silence this warning
46 | const float* w_row = weight + current_out * in_features;
| ^~~~~~~~~~~~~~~~~~~~~~~~~
| static_cast<ptrdiff_t>( )
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_76/b9_s1_even_workload_dist_base/base/base.cu:46:43: note: perform multiplication in a wider type
46 | const float* w_row = weight + current_out * in_features;
| ^~~~~~~~~~~
| static_cast<ptrdiff_t>( )
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_76/b9_s1_even_workload_dist_base/base/base.cu:94:64: 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]
94 | torch::Tensor even_workload_distribution_forward(torch::Tensor x,
| ^
| const &
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_76/b9_s1_even_workload_dist_base/base/base.cu:95:67: 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]
95 | torch::Tensor weight,
| ^
| const &
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_76/b9_s1_even_workload_dist_base/base/base.cu:96:67: 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]
96 | torch::Tensor bias) {
| ^
| const &
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_76/b9_s1_even_workload_dist_base/base/base.cu:101:22: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
101 | int batch_size = x.size(0);
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_76/b9_s1_even_workload_dist_base/base/base.cu:102:23: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
102 | int in_features = x.size(1);
| ^
/home/robert_sakana_ai/llm_cuda/experiments/20250203_optimize_b10_s4_e0_sweep/level_2/task_76/b9_s1_even_workload_dist_base/base/base.cu:103:24: warning: narrowing conversion from 'int64_t' (aka 'long') to signed type 'int' is implementation-defined [bugprone-narrowing-conversions]
103 | int out_features = weight.size(0);
| ^