@triton.jit
def geglu_backward_kernel(x_ptr, dx_ptr, dy_ptr, N, C, C2, BLK_C: tl.
    constexpr, BLK_N: tl.constexpr):
    pid_n = tl.program_id(0)
    pid_c = tl.program_id(1)
    offs_n = pid_n * BLK_N + tl.arange(0, BLK_N)
    offs_c = pid_c * BLK_C + tl.arange(0, BLK_C)
    mask_n = offs_n < N
    mask_c = offs_c < C2
    mask = mask_n[:, None] & mask_c[None, :]
    x_ptrs = x_ptr + offs_n[:, None] * C + offs_c[None, :]
    x1 = tl.load(x_ptrs, mask=mask)
    x2 = tl.load(x_ptrs + C2, mask=mask)
    dy_ptrs = dy_ptr + offs_n[:, None] * C2 + offs_c[None, :]
    dy = tl.load(dy_ptrs, mask=mask)
    dx1 = dy * gelu_forward(x2)
    dx2 = dy * x1
    dx2 *= gelu_backward(x2)
    dx_ptrs = dx_ptr + offs_n[:, None] * C + offs_c[None, :]
    tl.store(dx_ptrs, dx1, mask=mask)
    tl.store(dx_ptrs + C2, dx2, mask=mask)
