   }
 }
 
+template <typename Context, typename T>
+typename std::enable_if<!std::is_same<T, phi::dtype::float16>::value>::type
+ThrustCumsumKernel(const Context& dev_ctx,
+                   const T* in_data,
+                   T* out_data,
+                   int64_t size,
+                   bool reverse,
+                   bool exclusive) {
+#ifdef __HIPCC__
+  const auto& policy = thrust::hip::par.on(dev_ctx.stream());
+#else
+  const auto& policy = thrust::cuda::par.on(dev_ctx.stream());
+#endif
+  if (reverse) {
+    thrust::reverse_iterator<thrust::device_ptr<const T>> reversed_in(
+        thrust::device_pointer_cast(in_data) + size);
+    thrust::reverse_iterator<thrust::device_ptr<T>> reversed_out(
+        thrust::device_pointer_cast(out_data) + size);
+    if (exclusive) {
+      thrust::exclusive_scan(
+          policy, reversed_in, reversed_in + size, reversed_out);
+    } else {
+      thrust::inclusive_scan(
+          policy, reversed_in, reversed_in + size, reversed_out);
+    }
+  } else {
+    if (exclusive) {
+      thrust::exclusive_scan(policy, in_data, in_data + size, out_data);
+    } else {
+      thrust::inclusive_scan(policy, in_data, in_data + size, out_data);
+    }
+  }
+
+  return;
+}
+
+template <typename Context, typename T>
+typename std::enable_if<std::is_same<T, phi::dtype::float16>::value>::type
+ThrustCumsumKernel(const Context& dev_ctx,
+                   const phi::dtype::float16* in_data,
+                   phi::dtype::float16* out_data,
+                   int64_t size,
+                   bool reverse,
+                   bool exclusive) {}
+
 template <typename T, typename Context, typename Op>
 void ScanKernel(const Context& dev_ctx,
                 const DenseTensor& x,

#include "paddle/phi/core/kernel_registry.h"

namespace phi {

template <typename T>
class CumTypeTrait {
 public:
  using Type = T;
};

template <>
class CumTypeTrait<phi::dtype::float16> {
 public:
  using Type = __half;
};

template <typename T, int BLOCK_SIZE>
__device__ void BlockReverse(

  // length of the ‘axis’ dimension.
  if (!std::is_same<T, phi::dtype::float16>::value &&
      std::is_same<Op, cub::Sum>::value && size == out_dims[axis]) {
    ThrustCumsumKernel<Context, T>(
        dev_ctx, in_data, out_data, size, reverse, exclusive);
    return;
  }

