@triton.jit
def _AvgPoolCeilStats_cl3d_impl(X_ptr, Out_ptr, Mean_ptr, Sqmean_ptr,
    h_input, w_input, d_input, d_output, batch_stride_input, H_stride_input,
    W_stride_input, batch_stride_output, H_stride_output, W_stride_output,
    numel_no_channels_output, num_channels: tl.constexpr, almost_half_d: tl
    .constexpr):
    batch = tl.program_id(0)
    H = tl.program_id(1)
    W = tl.program_id(2)
    Out_ptr += (batch * batch_stride_output + H * H_stride_output + W *
        W_stride_output)
    output = tl.zeros([almost_half_d, num_channels], dtype=tl.float32)
    pair_offset = tl.arange(0, 2)
    channels_offset = tl.arange(0, num_channels)
    d_offset = tl.arange(0, almost_half_d)
    offset = d_offset[:, None, None] * (2 * num_channels) + channels_offset[
        None, :, None] + pair_offset[None, None, :] * num_channels
    output_offset = d_offset[:, None] * num_channels + channels_offset[None, :]
    mask_input = offset < d_input * num_channels
    output_mask = output_offset < d_output * num_channels
    norm_step = tl.sum(mask_input.to(tl.float32), axis=2).to(tl.float32)
    norm_step = tl.where(norm_step != 0, norm_step, 1.0)
    num_norm = 1
    Temp_ptr = (X_ptr + batch * batch_stride_input + 2 * H * H_stride_input +
        2 * W * W_stride_input)
    x = tl.load(Temp_ptr + offset, mask=mask_input, other=0.0).to(tl.float32)
    x = tl.sum(x, axis=2)
    output += x
    W_skip = False
    if 2 * (W + 1) > w_input:
        W_skip = True
    else:
        Temp_ptr = (X_ptr + batch * batch_stride_input + 2 * H *
            H_stride_input + (2 * W + 1) * W_stride_input)
        x = tl.load(Temp_ptr + offset, mask=mask_input, other=0.0).to(tl.
            float32)
        x = tl.sum(x, axis=2)
        output += x
        num_norm += 1
    H_skip = False
    if 2 * (H + 1) > h_input:
        H_skip = True
    else:
        Temp_ptr = X_ptr + batch * batch_stride_input + (2 * H + 1
            ) * H_stride_input + 2 * W * W_stride_input
        x = tl.load(Temp_ptr + offset, mask=mask_input, other=0.0).to(tl.
            float32)
        x = tl.sum(x, axis=2)
        output += x
        num_norm += 1
    if not H_skip and not W_skip:
        Temp_ptr = X_ptr + batch * batch_stride_input + (2 * H + 1
            ) * H_stride_input + (2 * W + 1) * W_stride_input
        x = tl.load(Temp_ptr + offset, mask=mask_input, other=0.0).to(tl.
            float32)
        x = tl.sum(x, axis=2)
        output += x
        num_norm += 1
    output = output / (norm_step * num_norm)
    tl.store(Out_ptr + output_offset, output, mask=output_mask)
    output = tl.trans(output)
    mean = tl.sum(output, axis=1) / numel_no_channels_output
    sqmean = tl.sum(output * output, axis=1) / numel_no_channels_output
    tl.atomic_add(Mean_ptr + channels_offset, mean)
    tl.atomic_add(Sqmean_ptr + channels_offset, sqmean)
