Study Plan

Masked Triton GELU Kernel

Activation FunctionsGPU Programming

Easy

Implement the tanh approximation of GELU in a Triton kernel. For each input value $x$ , the output is:

\operatorname{GELU}(x)=\frac{x}{2}\left(1+\tanh\left(\sqrt{\frac{2}{\pi}}\left(x+0.044715x^3\right)\right)\right).

Here, $\pi$ is pi and tanh is the hyperbolic tangent. Complete only the empty kernel body; the supplied solve launcher is already implemented. The pointers identify the input and output buffers, n_elements is their flattened length, and BLOCK_SIZE is the compile-time number of lanes per program block. Return no value; write GELU values into the supplied output tensor with the same shape and storage dtype as the input.

Examples

Input: x = [-2, -1, 0, 1, 2], dtype = float32

Output: [-0.045402, -0.158808, 0, 0.841192, 1.954598]

Explanation: GELU scales each value independently, attenuating negative values.

Input: x = [-3, -0.5, 0, 0.25, 1, 3, 8], dtype = float16

Output: [-0.003637, -0.154297, 0, 0.149658, 0.841309, 2.996094, 8]

Input: x = [1.5], dtype = bfloat16

Output: [1.398438]

Hint 1

tl.program_id identifies the current program block.

Hint 2

tl.load and tl.store accept masks for partial blocks.

Hint 3

libdevice.tanh evaluates the hyperbolic tangent.

Requirements

Use Triton and evaluate the stated approximation in float32.
Write every valid output element without accessing memory outside the buffers.
Preserve the input tensor.

Constraints

x and out are separate contiguous CUDA tensors with equal shape, device, and storage dtype.
Supported storage dtypes are float32, float16, and bfloat16; input values are finite.
The flattened length may be zero or not divisible by BLOCK_SIZE.
The supplied launcher handles empty inputs and converts bfloat16 through temporary float32 buffers.

Try Similar Problems

Triton Row-Wise SoftmaxMedium Fused Tiled Matmul and ReLUHard Implement Swish (SiLU) as a CUDA KernelEasy ReLU as a CUDA KernelEasy Implement Leaky ReLU as a CUDA KernelEasy Implement Sigmoid as a CUDA KernelEasy Implement Exact GELU as a CUDA KernelEasy GELU in TritonEasy ReLU in TritonEasy SiLU in TritonEasy Tanh as a CUDA KernelEasy Tiled Triton Row SumMedium

Masked Triton GELU Kernel

Activation FunctionsGPU Programming

Easy

Implement the tanh approximation of GELU in a Triton kernel. For each input value $x$ , the output is:

\operatorname{GELU}(x)=\frac{x}{2}\left(1+\tanh\left(\sqrt{\frac{2}{\pi}}\left(x+0.044715x^3\right)\right)\right).

Examples

Input: x = [-2, -1, 0, 1, 2], dtype = float32

Output: [-0.045402, -0.158808, 0, 0.841192, 1.954598]

Explanation: GELU scales each value independently, attenuating negative values.

Input: x = [-3, -0.5, 0, 0.25, 1, 3, 8], dtype = float16

Output: [-0.003637, -0.154297, 0, 0.149658, 0.841309, 2.996094, 8]

Input: x = [1.5], dtype = bfloat16

Output: [1.398438]

Hint 1

tl.program_id identifies the current program block.

Hint 2

tl.load and tl.store accept masks for partial blocks.

Hint 3

libdevice.tanh evaluates the hyperbolic tangent.

Requirements

Use Triton and evaluate the stated approximation in float32.
Write every valid output element without accessing memory outside the buffers.
Preserve the input tensor.

Constraints

x and out are separate contiguous CUDA tensors with equal shape, device, and storage dtype.
Supported storage dtypes are float32, float16, and bfloat16; input values are finite.
The flattened length may be zero or not divisible by BLOCK_SIZE.
The supplied launcher handles empty inputs and converts bfloat16 through temporary float32 buffers.

Try Similar Problems

x =

Accepts: array

dtype =

Accepts: string

Examples

Requirements

Constraints

Kernel performance

Masked Triton GELU Kernel

Examples

Requirements

Constraints

Kernel performance