@triton.jit
def rope_bw(dx_ptr, f_ptr, dt_ptr, dx_s_stride, f_s_stride, dt_s_stride, d,
    d2, BLOCK_SIZE: tl.constexpr):
    s_idx = tl.program_id(0)
    bh_idx = tl.program_id(1)
    dx_start_ptr = dx_ptr + s_idx * dx_s_stride
    f_start_ptr = f_ptr + s_idx * f_s_stride
    dt_start_ptr = dt_ptr + s_idx * dt_s_stride
    d2_half = d2 // 2
    col_offsets = tl.arange(0, BLOCK_SIZE)
    mask = col_offsets < d2_half
    f0_ptrs = f_start_ptr + col_offsets
    f1_ptrs = f_start_ptr + col_offsets + d2_half
    f0 = tl.load(f0_ptrs, mask=mask, other=0.0)
    cos0 = tl.cos(f0)
    sin0 = tl.sin(f0)
    f1 = tl.load(f1_ptrs, mask=mask, other=0.0)
    cos1 = tl.cos(f1)
    sin1 = tl.sin(f1)
    dx0_ptrs = dx_start_ptr + bh_idx * d + col_offsets
    dx1_ptrs = dx_start_ptr + bh_idx * d + col_offsets + d2_half
    dx0 = tl.load(dx0_ptrs, mask=mask, other=0.0)
    dx1 = tl.load(dx1_ptrs, mask=mask, other=0.0)
    dt0 = dx0 * cos0 + dx1 * sin1
    dt1 = dx1 * cos1 - dx0 * sin0
    dt0_ptrs = dt_start_ptr + bh_idx * d + col_offsets
    dt1_ptrs = dt_start_ptr + bh_idx * d + col_offsets + d2_half
    tl.store(dt0_ptrs, dt0, mask=mask)
    tl.store(dt1_ptrs, dt1, mask=mask)
    if d2 < d:
        remainder = d - d2
        q, r = remainder // BLOCK_SIZE, remainder % BLOCK_SIZE
        for i in range(q):
            dx2_ptrs = (dx_start_ptr + bh_idx * d + col_offsets + d2 + 
                BLOCK_SIZE * i)
            dt2_ptrs = (dt_start_ptr + bh_idx * d + col_offsets + d2 + 
                BLOCK_SIZE * i)
            dx2 = tl.load(dx2_ptrs)
            tl.store(dt2_ptrs, dx2)
        if r > 0:
            dx2_ptrs = (dx_start_ptr + bh_idx * d + col_offsets + d2 + 
                BLOCK_SIZE * q)
            dt2_ptrs = (dt_start_ptr + bh_idx * d + col_offsets + d2 + 
                BLOCK_SIZE * q)
            mask = col_offsets < r
            dx2 = tl.load(dx2_ptrs, mask=mask, other=0.0)
            tl.store(dt2_ptrs, dx2, mask=mask)
