Study Plan

Tiled Triton Row Sum

GPU Programming

Medium

Implement a tiled row reduction in Triton. For an input matrix $x$ with $R$ rows and $C$ columns, its output $y$ contains one float32 sum per row:

y_r=\sum_{c=0}^{C-1}\operatorname{float32}(x_{r,c}).

Here, $r$ and $c$ index rows and columns, and float32 denotes conversion before summation. Complete only the empty kernel body; the supplied solve launcher is implemented. The pointers identify input and output buffers, n_cols is the column count, row_stride and col_stride are input strides in elements, and BLOCK_SIZE is the compile-time tile width. Return no value; write a contiguous float32 vector of length R into the supplied output buffer.

Examples

Input: x = [[1, 1, 1], [1, 1, 1]], dtype = float32, tile_size = 32

Output: [3, 3]

Explanation: Each row contains three ones, so each row sum is 3.

Input: x has shape (1, 33) with every value equal to 1, dtype = float16, tile_size = 32

Output: [33]

Input: x = [[1, 2, -1], [0.5, 1.5, 2]], dtype = bfloat16, tile_size = 64

Output: [2, 4]

Hint 1

tl.load accepts a mask and a fallback value for incomplete tiles.

Hint 2

tl.sum reduces values along a block axis.

Hint 3

tl.program_id identifies the current row program.

Requirements

Use Triton with float32 accumulation and fixed-width tiles.
Support positive input strides without modifying the input.
Zero-width rows produce zero sums; zero rows require no writes.

Constraints

x is a two-dimensional CUDA tensor with finite float32, float16, or bfloat16 values.
out is a separate contiguous float32 CUDA vector with one entry per input row, on the same device.
tile_size is a power of two from 32 through 1024.
The supplied launcher handles invalid tile sizes and empty dimensions, and converts bfloat16 inputs to float32.

Try Similar Problems

Triton Row-Wise SoftmaxMedium Masked Triton GELU KernelEasy Fused Tiled Matmul and ReLUHard Blockwise Online SoftmaxHard Argmax as a CUDA KernelMedium Row-Wise LogSumExp in TritonMedium RMS Normalization as a CUDA KernelMedium Mean and Variance as a CUDA KernelMedium Scaled Dot-Product Attention in TritonMedium Block-Pointer Matmul in TritonMedium Layer Normalization as a CUDA KernelMedium Single-Pass Mean and Variance in TritonMedium

Tiled Triton Row Sum

GPU Programming

Medium

Implement a tiled row reduction in Triton. For an input matrix $x$ with $R$ rows and $C$ columns, its output $y$ contains one float32 sum per row:

y_r=\sum_{c=0}^{C-1}\operatorname{float32}(x_{r,c}).

Examples

Input: x = [[1, 1, 1], [1, 1, 1]], dtype = float32, tile_size = 32

Output: [3, 3]

Explanation: Each row contains three ones, so each row sum is 3.

Input: x has shape (1, 33) with every value equal to 1, dtype = float16, tile_size = 32

Output: [33]

Input: x = [[1, 2, -1], [0.5, 1.5, 2]], dtype = bfloat16, tile_size = 64

Output: [2, 4]

Hint 1

tl.load accepts a mask and a fallback value for incomplete tiles.

Hint 2

tl.sum reduces values along a block axis.

Hint 3

tl.program_id identifies the current row program.

Requirements

Use Triton with float32 accumulation and fixed-width tiles.
Support positive input strides without modifying the input.
Zero-width rows produce zero sums; zero rows require no writes.

Constraints

x is a two-dimensional CUDA tensor with finite float32, float16, or bfloat16 values.
out is a separate contiguous float32 CUDA vector with one entry per input row, on the same device.
tile_size is a power of two from 32 through 1024.
The supplied launcher handles invalid tile sizes and empty dimensions, and converts bfloat16 inputs to float32.

Try Similar Problems

x =

Accepts: array

dtype =

Accepts: string

tile_size =

Accepts: number

Examples

Requirements

Constraints

Kernel performance

Tiled Triton Row Sum

Examples

Requirements

Constraints

Kernel performance