[Operator Mechanism] fix BroadcastConfig do not support big_tensor

paddle/phi/kernels/primitive/datamover_primitives.h

@@ -41,7 +41,7 @@ struct alignas(sizeof(T) * VecSize) VectorType {
  * must be [dim1, dim0].
  */
 struct BroadcastConfig {
-  funcs::FastDivMod<int> divmoders[DDim::kMaxRank];
+  funcs::FastDivMod<int64_t> divmoders[DDim::kMaxRank];
   uint64_t strides[DDim::kMaxRank];
   int rank{0};
 
@@ -52,16 +52,15 @@ struct BroadcastConfig {
                   const std::vector<int64_t>& in_dims,
                   int dim_size) {
     for (int i = 0; i < dim_size; ++i) {
-      PADDLE_ENFORCE_LE_INT_MAX(out_dims[i], "out_dim");
-      divmoders[i] = funcs::FastDivMod<int>(static_cast<int>(out_dims[i]));
+      divmoders[i] = funcs::FastDivMod<int64_t>(out_dims[i]);
     }
 
     for (int i = 0; i < dim_size; ++i) {
       strides[i] = in_dims[i] == 1 ? 0 : 1;
       strides[i] = (i != 0 && strides[i] != 0)
                        ? std::accumulate(in_dims.begin(),
                                          in_dims.begin() + i,
-                                         1,
+                                         int64_t{1},
                                          std::multiplies<int64_t>())
                        : strides[i];
     }
@@ -416,17 +415,19 @@ __device__ __forceinline__ void ReadDataBc(
     const T* __restrict__ src,
     uint32_t block_offset,
     const details::BroadcastConfig& config,
-    int total_num_output,
+    uint64_t total_num_output,
     int stride_nx,
     int stride_ny) {
   uint32_t thread_offset = block_offset + threadIdx.x;
-  uint32_t index_src = 0;
+  uint64_t index_src = 0;
 
 #pragma unroll
   for (int ny = 0; ny < NY; ++ny) {
 #pragma unroll
     for (uint32_t nx = 0; nx < NX; ++nx) {
-      uint32_t index_output = thread_offset + ny * stride_ny + nx * stride_nx;
+      uint64_t index_output = thread_offset +
+                              static_cast<uint64_t>(ny) * stride_ny +
+                              static_cast<uint64_t>(nx) * stride_nx;
       index_src = 0;
       if (IsBoundary) {
         if (index_output >= total_num_output) {
@@ -752,14 +753,14 @@ __device__ __forceinline__ void ReadDataBc(
     const T* __restrict__ src,
     uint32_t block_offset,
     const details::BroadcastConfig& config,
-    int total_num_output,
+    uint64_t total_num_output,
     int read_lens = NX) {
   uint32_t thread_offset = block_offset + threadIdx.x * NX;
-  uint32_t index_src = 0;
+  uint64_t index_src = 0;
 
 #pragma unroll
   for (uint32_t nx = 0; nx < NX; ++nx) {
-    uint32_t index_output = thread_offset + nx;
+    uint64_t index_output = thread_offset + nx;
     index_src = 0;
     if (IsBoundary) {
       if (index_output >= total_num_output) {
@@ -813,14 +814,14 @@ __device__ __forceinline__ void ReadDataBc(
     const T* __restrict__ src,
     uint32_t block_offset,
     const details::BroadcastConfig& config,
-    int total_num_output,
+    uint64_t total_num_output,
     int read_lens = NX) {
   uint32_t thread_offset = block_offset + threadIdx.x * NX;
-  uint32_t index_src = 0;
+  uint64_t index_src = 0;
 
 #pragma unroll
   for (uint32_t nx = 0; nx < NX; ++nx) {
-    uint32_t index_output = thread_offset + nx;
+    uint64_t index_output = thread_offset + nx;
     index_src = 0;
     if (IsBoundary) {
       if (index_output >= total_num_output) {