 
 #include "paddle/fluid/operators/temporal_shift_op.h"
 #include "paddle/fluid/platform/cuda_primitives.h"
+#include "paddle/fluid/platform/gpu_launch_config.h"
 
 namespace paddle {
 namespace operators {

    T* output_data = output->mutable_data<T>({nt, c, h, w}, ctx.GetPlace());

    int pixelNum = nt * chw;
    int grid_dim = (pixelNum + 512 - 1) / 512;
    grid_dim = grid_dim > 8 ? 8 : grid_dim;

    KeTemporalShiftFw<
        T><<<grid_dim, 512, 0, ctx.cuda_device_context().stream()>>>(
        input_data, output_data, ntchw, tchw, chw, hw, w, t, c, shift_ratio);
  }
};

        static_cast<T>(0));

    int pixelNum = nt * chw;
    platform::GpuLaunchConfig config =
        platform::GetGpuLaunchConfig1D(ctx.cuda_device_context(), pixelNum);

    KeTemporalShiftBw<T><<<config.block_per_grid, config.thread_per_block, 0,
                           ctx.cuda_device_context().stream()>>>(
        output_grad_data, input_grad_data, ntchw, tchw, chw, hw, w, t, c,
        shift_ratio);
  }
