diff --git a/paddle/phi/kernels/funcs/broadcast_function.h b/paddle/phi/kernels/funcs/broadcast_function.h
index 5596e937d17cf6..ee4615ea49871c 100644
--- a/paddle/phi/kernels/funcs/broadcast_function.h
+++ b/paddle/phi/kernels/funcs/broadcast_function.h
@@ -14,6 +14,7 @@ limitations under the License. */
 
 #pragma once
 
+#include <limits>
 #include <sstream>
 #include "paddle/common/enforce.h"
 #include "paddle/phi/kernels/funcs/elementwise_base.h"
@@ -30,13 +31,33 @@ namespace funcs {
 
 enum BroadcastType { kMixed = 1, kBroadcast = 2, kElementwise = 3 };
 
+// On XPU the BroadcastConfig lives in datamover_primitives_xpu2.h and is not
+// templated; on GPU it is templated on the index type. This alias lets the
+// shared kernel signatures name the right type on each backend.
+#ifdef PADDLE_WITH_XPU_KP
+template <typename IndexT>
+using BroadcastConfigType = kps::details::BroadcastConfig;
+#else
+template <typename IndexT>
+using BroadcastConfigType = kps::details::BroadcastConfig<IndexT>;
+#endif
+
 template <typename OutT, typename Functor, int Arity, int NumOuts>
 struct BroadcastTypeClassifier {
   int64_t numel{0};
   int broadcast_num{0};              // Not used for XPU
   bool all_elementwise{true};        // Not used for XPU
   Array<bool, Arity> use_broadcast;  // Not used for XPU
+#ifdef PADDLE_WITH_XPU_KP
   Array<kps::details::BroadcastConfig, Arity> configs;
+#else
+  // The GPU BroadcastConfig is templated on the index type, which is only
+  // decided at launch time (uint32 vs uint64). Keep the simplified dims here
+  // and build the typed configs later in LaunchBroadcastKernel.
+  std::vector<int64_t> out_dims;
+  std::vector<std::vector<int64_t>> in_dims;
+  int rank{0};
+#endif
   Array<const _ptr_ char *__restrict__, Arity> ins_data;
   Array<_ptr_ OutT *, NumOuts> outs_data;
 
@@ -46,7 +67,12 @@ struct BroadcastTypeClassifier {
                           int axis) {
     numel = (*outs)[0]->numel();
 
-#ifndef PADDLE_WITH_XPU_KP
+#ifdef PADDLE_WITH_XPU_KP
+    // datamover_primitives_xpu2.h::BroadcastConfig (built inside
+    // InitBroadcastConfigs below) still computes strides/numel with 32-bit
+    // int, so the INT_MAX check must happen before that call, not after.
+    PADDLE_ENFORCE_LE_INT_MAX(numel, "BroadcastKernel numel (XPU)");
+#else
     for (size_t i = 0; i < ins.size(); ++i) {
       bool is_same_dim = ins[i]->numel() == numel;
       if (is_same_dim) {
@@ -96,14 +122,17 @@ struct BroadcastTypeClassifier {
         DimsSimplifiedLogger<int64_t>::Log(
             ins, outs, dims_simplifier, "BroadcastKernel");
       }
+      // Store the simplified dims; the typed BroadcastConfig is built once the
+      // index type is chosen in LaunchBroadcastKernel. An empty in_dims[i]
+      // means "no config needed" (mirrors the old `ins[i]->numel()` guard).
+      out_dims = dims_simplifier.out_dims;
+      rank = dims_simplifier.rank;
+      in_dims.resize(Arity);
       for (int i = 0; i < Arity; ++i) {
         // if data shape is[m, n], then you should set data_dim = {n, m}
         // eg: out's shape [3, 45, 1]. then out_dims = {1, 45, 3}
-        // if (ins[i]->numel() != (*outs)[0]->numel()) {
         if (ins[i]->numel()) {
-          configs[i] = kps::details::BroadcastConfig(dims_simplifier.out_dims,
-                                                     dims_simplifier.in_dims[i],
-                                                     dims_simplifier.rank);
+          in_dims[i] = dims_simplifier.in_dims[i];
         }
       }
     }
@@ -112,16 +141,20 @@ struct BroadcastTypeClassifier {
 };
 
 // Common broadcast/elementwise Loader.
-template <int Index, int VecSize, bool IsBoundary, int LoadType>
+template <typename IndexT,
+          int Index,
+          int VecSize,
+          bool IsBoundary,
+          int LoadType>
 struct BroadcastDataLoader {
   template <typename Array1, typename Array2, typename Array3, typename ArgsT>
   static __device__ __forceinline__ void Apply(const Array1 &ins,
                                                ArgsT *args,
                                                const Array2 &configs,
                                                const Array3 &use_broadcast,
-                                               const int block_offset,
+                                               IndexT block_offset,
                                                const int num,
-                                               const uint32_t numel,
+                                               IndexT numel,
                                                int read_lens) {
     using Type = std::tuple_element_t<Index, ArgsT>;
 #ifdef PADDLE_WITH_XPU_KP
@@ -145,7 +178,7 @@ struct BroadcastDataLoader {
 #else
     kps::Init<Type, ArgsT, Index, VecSize>(args, static_cast<Type>(1.0f));
     if (use_broadcast[Index]) {
-      kps::ReadDataBc<Type, VecSize, 1, ArgsT, Index, IsBoundary>(
+      kps::ReadDataBc<Type, IndexT, VecSize, 1, ArgsT, Index, IsBoundary>(
           args,
           reinterpret_cast<const _ptr_ Type *>(ins[Index]),
           block_offset,
@@ -170,23 +203,24 @@ struct BroadcastDataLoader {
 /* BroadcastDataLoaders Partial specialization */
 #ifndef PADDLE_WITH_XPU_KP
 // Scalar elementwise Loader with consideration of IsBoundary.
-template <int Index, int VecSize>
-struct BroadcastDataLoader<Index, VecSize, true, kElementwise> {
+template <typename IndexT, int Index, int VecSize>
+struct BroadcastDataLoader<IndexT, Index, VecSize, true, kElementwise> {
   template <typename Array1, typename Array2, typename Array3, typename ArgsT>
   static __device__ __forceinline__ void Apply(const Array1 &ins,
                                                ArgsT *args,
                                                const Array2 &configs,
                                                const Array3 &use_broadcast,
-                                               const int block_offset,
+                                               IndexT block_offset,
                                                const int num,
-                                               const uint32_t numel,
+                                               IndexT numel,
                                                int read_lens) {
     using Type = std::tuple_element_t<Index, ArgsT>;
-    int thread_offset = threadIdx.x * VecSize + block_offset;
+    IndexT thread_offset =
+        static_cast<IndexT>(threadIdx.x) * VecSize + block_offset;
 #pragma unroll
     for (int idx = 0; idx < VecSize; ++idx) {
       std::get<Index>(args[idx]) = static_cast<Type>(1);
-      int index = thread_offset + idx;
+      IndexT index = thread_offset + idx;
       if (index < numel) {
         std::get<Index>(args[idx]) =
             reinterpret_cast<const _ptr_ Type *>(ins[Index])[index];
@@ -196,24 +230,23 @@ struct BroadcastDataLoader<Index, VecSize, true, kElementwise> {
 };
 
 // Vectorized elementwise Loader without consideration of IsBoundary.
-template <int Index, int VecSize>
-struct BroadcastDataLoader<Index, VecSize, false, kElementwise> {
+template <typename IndexT, int Index, int VecSize>
+struct BroadcastDataLoader<IndexT, Index, VecSize, false, kElementwise> {
   template <typename Array1, typename Array2, typename Array3, typename ArgsT>
   static __device__ __forceinline__ void Apply(const Array1 &ins,
                                                ArgsT *args,
                                                const Array2 &configs,
                                                const Array3 &use_broadcast,
-                                               const int block_offset,
+                                               IndexT block_offset,
                                                const int num,
-                                               const uint32_t numel,
+                                               IndexT numel,
                                                int read_lens) {
     using Type = std::tuple_element_t<Index, ArgsT>;
     using VecType = phi::kps::details::VectorType<Type, VecSize>;
     VecType vec_temp;
 
-    int64_t thread_offset =
-        static_cast<int64_t>(threadIdx.x) +
-        static_cast<int64_t>(blockIdx.x) * static_cast<int64_t>(blockDim.x);
+    IndexT thread_offset =
+        block_offset / VecSize + static_cast<IndexT>(threadIdx.x);
     const VecType *__restrict__ vec_input =
         reinterpret_cast<const VecType *__restrict__>(ins[Index]);
     vec_temp = vec_input[thread_offset];
@@ -238,10 +271,10 @@ struct BroadcastDataInit {
 
 template <int Index, int VecSize>
 struct BroadcastDataSetter {
-  template <typename Array, typename ArgsT>
+  template <typename Array, typename ArgsT, typename IndexT>
   static __device__ __forceinline__ void Apply(const Array &ins,
                                                ArgsT *args,
-                                               uint32_t index_bc[][VecSize]) {
+                                               IndexT index_bc[][VecSize]) {
     using Type = std::tuple_element_t<Index, ArgsT>;
 #pragma unroll
     for (int k = 0; k < VecSize; ++k) {
@@ -254,8 +287,13 @@ struct BroadcastDataSetter {
 #endif
 
 // static broadcast unroller
-template <template <int Index, int VecSize, bool IsBoundary, int LoadType>
+template <template <typename IndexT,
+                    int Index,
+                    int VecSize,
+                    bool IsBoundary,
+                    int LoadType>
           typename Func,
+          typename IndexT,
           bool IsBoundary,
           int LoadType,
           int VecSize,
@@ -264,26 +302,32 @@ template <template <int Index, int VecSize, bool IsBoundary, int LoadType>
 struct BcUnroller {
   template <typename... Args>
   static HOSTDEVICE inline void step(Args &&...args) {
-    Func<Begin, VecSize, IsBoundary, LoadType>::Apply(
+    Func<IndexT, Begin, VecSize, IsBoundary, LoadType>::Apply(
         std::forward<Args>(args)...);
-    BcUnroller<Func, IsBoundary, LoadType, VecSize, End, Begin + 1>::step(
-        args...);
+    BcUnroller<Func, IndexT, IsBoundary, LoadType, VecSize, End, Begin + 1>::
+        step(args...);
   }
 };
 
-template <template <int Index, int VecSize, bool IsBoundary, int LoadType>
+template <template <typename IndexT,
+                    int Index,
+                    int VecSize,
+                    bool IsBoundary,
+                    int LoadType>
           typename Func,
+          typename IndexT,
           bool IsBoundary,
           int LoadType,
           int VecSize,
           int End>
-struct BcUnroller<Func, IsBoundary, LoadType, VecSize, End, End> {
+struct BcUnroller<Func, IndexT, IsBoundary, LoadType, VecSize, End, End> {
   template <typename... Args>
   static HOSTDEVICE inline void step(Args &&...args) {}
 };
 
 template <typename OutT,
           typename Functor,
+          typename IndexT,
           int Arity,
           int NumOuts,
           int VecSize,
@@ -293,10 +337,10 @@ __device__ void VectorizedBroadcastKernelImpl(
     const Array<const _ptr_ char *__restrict__, Arity> &ins,
     Array<_ptr_ OutT *, NumOuts> outs,
     const Array<bool, Arity> &use_broadcast,
-    const uint32_t numel,
-    const Array<kps::details::BroadcastConfig, Arity> &configs,
+    const IndexT numel,
+    const Array<BroadcastConfigType<IndexT>, Arity> &configs,
     uint32_t num,
-    uint32_t block_offset,
+    IndexT block_offset,
     int read_lens,
     Functor func) {
   using Traits = funcs::FunctionTraits<Functor>;
@@ -305,16 +349,28 @@ __device__ void VectorizedBroadcastKernelImpl(
   __simd__ ConditionalT<OutT, NumOuts> result[VecSize];
 
 #ifdef PADDLE_WITH_XPU_KP
-  BcUnroller<BroadcastDataLoader, IsBoundary, LoadType, VecSize, Arity>::step(
-      ins, args, configs, use_broadcast, block_offset, num, numel, read_lens);
+  BcUnroller<BroadcastDataLoader,
+             IndexT,
+             IsBoundary,
+             LoadType,
+             VecSize,
+             Arity>::step(ins,
+                          args,
+                          configs,
+                          use_broadcast,
+                          block_offset,
+                          num,
+                          numel,
+                          read_lens);
 #else
   if (LoadType == kBroadcast) {
-    uint32_t index_bc[Arity][VecSize] = {0};
+    IndexT index_bc[Arity][VecSize] = {0};
     Unroller<BroadcastDataInit, VecSize, Arity>::step(args);
-    uint32_t thread_offset = block_offset + threadIdx.x * VecSize;
+    IndexT thread_offset =
+        block_offset + static_cast<IndexT>(threadIdx.x) * VecSize;
 #pragma unroll
     for (int k = 0; k < VecSize; ++k) {
-      uint32_t idx = thread_offset + k;
+      IndexT idx = thread_offset + k;
       if (IsBoundary && idx == numel) break;
 #pragma unroll
       for (int i = 0; i < DDim::kMaxRank; ++i) {
@@ -329,8 +385,19 @@ __device__ void VectorizedBroadcastKernelImpl(
     }
     Unroller<BroadcastDataSetter, VecSize, Arity>::step(ins, args, index_bc);
   } else {
-    BcUnroller<BroadcastDataLoader, IsBoundary, LoadType, VecSize, Arity>::step(
-        ins, args, configs, use_broadcast, block_offset, num, numel, read_lens);
+    BcUnroller<BroadcastDataLoader,
+               IndexT,
+               IsBoundary,
+               LoadType,
+               VecSize,
+               Arity>::step(ins,
+                            args,
+                            configs,
+                            use_broadcast,
+                            block_offset,
+                            num,
+                            numel,
+                            read_lens);
   }
 #endif
   SameDimsElementwisePrimitiveCaller<ConditionalT<OutT, NumOuts>,
@@ -344,6 +411,7 @@ __device__ void VectorizedBroadcastKernelImpl(
 
 template <typename Functor,
           typename OutT,
+          typename IndexT,
           int Arity,
           int NumOuts,
           int VecSize,
@@ -352,19 +420,20 @@ __global__ void VectorizedBroadcastKernel(
     Array<const _ptr_ char *__restrict__, Arity> ins,
     Array<_ptr_ OutT *, NumOuts> outs,
     Array<bool, Arity> use_broadcast,
-    uint32_t numel,
-    Array<kps::details::BroadcastConfig, Arity> configs,
-    uint32_t main_offset,
-    uint32_t tail_tid,
+    IndexT numel,
+    Array<BroadcastConfigType<IndexT>, Arity> configs,
+    IndexT main_offset,
+    IndexT tail_tid,
     int read_lens,
     Functor func) {
 #ifdef PADDLE_WITH_XPU_KP
-  int64_t block_offset =
-      static_cast<int64_t>(BLOCK_ID_X) * BLOCK_NUM_X * read_lens;
-  int64_t stride = static_cast<int64_t>(BLOCK_NUM_X) * GRID_NUM_X * read_lens;
+  IndexT block_offset =
+      static_cast<IndexT>(BLOCK_ID_X) * BLOCK_NUM_X * read_lens;
+  IndexT stride = static_cast<IndexT>(BLOCK_NUM_X) * GRID_NUM_X * read_lens;
   for (; block_offset < main_offset; block_offset += stride) {
     VectorizedBroadcastKernelImpl<OutT,
                                   Functor,
+                                  IndexT,
                                   Arity,
                                   NumOuts,
                                   VecSize,
@@ -379,10 +448,15 @@ __global__ void VectorizedBroadcastKernel(
                                             read_lens,
                                             func);
   }
-  int64_t num = numel - block_offset;
-  if (num > 0) {
+  // `num` is the tail element count handled by a single block, always bounded
+  // by one block's span, so casting it to uint32 for the Impl is safe. Guard on
+  // `block_offset < numel` (not `num > 0`) so the subtraction never underflows
+  // when IndexT is unsigned.
+  if (block_offset < numel) {
+    IndexT num = numel - block_offset;
     VectorizedBroadcastKernelImpl<OutT,
                                   Functor,
+                                  IndexT,
                                   Arity,
                                   NumOuts,
                                   VecSize,
@@ -392,17 +466,18 @@ __global__ void VectorizedBroadcastKernel(
                                             use_broadcast,
                                             numel,
                                             configs,
-                                            num,
+                                            static_cast<uint32_t>(num),
                                             block_offset,
                                             read_lens,
                                             func);
   }
 #else
-  int64_t block_offset =
-      static_cast<int64_t>(BLOCK_ID_X) * BLOCK_NUM_X * VecSize;
-  if (block_offset < main_offset) {
+  IndexT block_offset = static_cast<IndexT>(BLOCK_ID_X) * BLOCK_NUM_X * VecSize;
+  IndexT stride = static_cast<IndexT>(BLOCK_NUM_X) * GRID_NUM_X * VecSize;
+  for (; block_offset < main_offset; block_offset += stride) {
     VectorizedBroadcastKernelImpl<OutT,
                                   Functor,
+                                  IndexT,
                                   Arity,
                                   NumOuts,
                                   VecSize,
@@ -416,9 +491,18 @@ __global__ void VectorizedBroadcastKernel(
                                             block_offset,
                                             read_lens,
                                             func);
-  } else {
+  }
+  // Tail block only: threads past `num` may have an out-of-range (and, when
+  // IndexT==uint32 and numel is near UINT32_MAX, wrapped) thread_offset, but
+  // their stores are masked out by `num` in WriteData, so the result stays
+  // correct. `num` is one block's tail count, so casting to uint32 is safe.
+  // Guard on `block_offset < numel` (not `num > 0`) so the subtraction never
+  // underflows when IndexT is unsigned.
+  if (block_offset < numel) {
+    IndexT num = numel - block_offset;
     VectorizedBroadcastKernelImpl<OutT,
                                   Functor,
+                                  IndexT,
                                   Arity,
                                   NumOuts,
                                   VecSize,
@@ -428,7 +512,7 @@ __global__ void VectorizedBroadcastKernel(
                                             use_broadcast,
                                             numel,
                                             configs,
-                                            tail_tid,
+                                            static_cast<uint32_t>(num),
                                             block_offset,
                                             read_lens,
                                             func);
@@ -441,52 +525,32 @@ void LaunchBroadcastKernel(
     const KPDevice &dev_ctx,
     const BroadcastTypeClassifier<OutT, Functor, Arity, NumOuts> &classifier,
     Functor func) {
+#ifndef PADDLE_WITH_XPU_KP
+  auto gpu_config = phi::backends::gpu::GetGpuLaunchConfig1D(
+      dev_ctx, classifier.numel, VecSize);
+#endif
+  auto launch_with_index_t = [&](auto index_tag) {
+    using IndexT = decltype(index_tag);
+    // Although this runs on the host and int64 would be numerically fine, we
+    // keep IndexT here so the offsets/main_offset/tail_tid passed to the kernel
+    // match the IndexT the device kernel is instantiated with.
 #ifdef PADDLE_WITH_XPU_KP
-  const int64_t numel_64 = classifier.numel;
-  PADDLE_ENFORCE_LE_UINT32_MAX(numel_64, "XPU broadcast kernel numel");
-  const uint32_t numel = static_cast<uint32_t>(numel_64);
-  const int threads = 64;
-  const int blocks = 8;
-  int read_lens = configs[0].buf_len;
-  auto stream = dev_ctx.x_context()->xpu_stream;
-  const int64_t block_len = static_cast<int64_t>(read_lens) * threads;
-  const int64_t main_offset_64 = (numel_64 / block_len) * block_len;
-  const int64_t tail_tid_64 = numel_64 % block_len;
-  const uint32_t main_offset = static_cast<uint32_t>(main_offset_64);
-  const uint32_t tail_tid = static_cast<uint32_t>(tail_tid_64);
-
-  VectorizedBroadcastKernel<Functor, OutT, Arity, NumOuts, VecSize, false>
-      <<<blocks, threads, 0, stream>>>(classifier.ins_data,
-                                       classifier.outs_data,
-                                       classifier.use_broadcast,
-                                       numel,
-                                       classifier.configs,
-                                       main_offset,
-                                       tail_tid,
-                                       read_lens,
-                                       func);
-#else
-  const int64_t numel_64 = classifier.numel;
-  auto gpu_config =
-      phi::backends::gpu::GetGpuLaunchConfig1D(dev_ctx, numel_64, VecSize);
-  auto stream = dev_ctx.stream();
-  uint32_t threads = static_cast<uint32_t>(gpu_config.GetBlockSize());
-  auto blocks = gpu_config.block_per_grid;
-  PADDLE_ENFORCE_LE_UINT32_MAX(numel_64, "numel");
-  const uint32_t numel = static_cast<uint32_t>(numel_64);
-  const int64_t block_len = static_cast<int64_t>(VecSize) * threads;
-  const int64_t main_offset_64 = (numel_64 / block_len) * block_len;
-  const int64_t tail_tid_64 = numel_64 % block_len;
-  uint32_t main_offset = static_cast<uint32_t>(main_offset_64);
-  uint32_t tail_tid = static_cast<uint32_t>(tail_tid_64);
-
-  if (classifier.all_elementwise) {
+    const IndexT numel = static_cast<IndexT>(classifier.numel);
+    const int threads = 64;
+    const int blocks = 8;
+    int read_lens = classifier.configs[0].buf_len;
+    auto stream = dev_ctx.x_context()->xpu_stream;
+    const IndexT block_len = static_cast<IndexT>(read_lens) * threads;
+    const IndexT main_offset = (numel / block_len) * block_len;
+    const IndexT tail_tid = numel % block_len;
+
     VectorizedBroadcastKernel<Functor,
                               OutT,
+                              IndexT,
                               Arity,
                               NumOuts,
                               VecSize,
-                              kElementwise>
+                              false>
         <<<blocks, threads, 0, stream>>>(classifier.ins_data,
                                          classifier.outs_data,
                                          classifier.use_broadcast,
@@ -494,31 +558,107 @@ void LaunchBroadcastKernel(
                                          classifier.configs,
                                          main_offset,
                                          tail_tid,
-                                         VecSize,
-                                         func);
-  } else if (classifier.broadcast_num > (Arity >> 1)) {
-    constexpr BroadcastType type_ = (Arity > 1) ? kBroadcast : kMixed;
-    VectorizedBroadcastKernel<Functor, OutT, Arity, NumOuts, VecSize, type_>
-        <<<blocks, threads, 0, stream>>>(classifier.ins_data,
-                                         classifier.outs_data,
-                                         classifier.use_broadcast,
-                                         numel,
-                                         classifier.configs,
-                                         main_offset,
-                                         tail_tid,
-                                         VecSize,
+                                         read_lens,
                                          func);
+#else
+    const IndexT numel = static_cast<IndexT>(classifier.numel);
+    auto stream = dev_ctx.stream();
+    uint32_t threads = static_cast<uint32_t>(gpu_config.GetBlockSize());
+    auto blocks = gpu_config.block_per_grid;
+    const IndexT block_len = static_cast<IndexT>(VecSize) * threads;
+    const IndexT main_offset = (numel / block_len) * block_len;
+    const IndexT tail_tid = numel % block_len;
+
+    // Build the index-typed configs now that IndexT is known. An empty
+    // in_dims[i] means the input needs no broadcast config (see
+    // InitBroadcastConfigs).
+    Array<BroadcastConfigType<IndexT>, Arity> configs;
+    if (!classifier.all_elementwise) {
+      for (int i = 0; i < Arity; ++i) {
+        if (!classifier.in_dims[i].empty()) {
+          configs[i] = kps::details::BroadcastConfig<IndexT>(
+              classifier.out_dims, classifier.in_dims[i], classifier.rank);
+        }
+      }
+    }
+
+    if (classifier.all_elementwise) {
+      VectorizedBroadcastKernel<Functor,
+                                OutT,
+                                IndexT,
+                                Arity,
+                                NumOuts,
+                                VecSize,
+                                kElementwise>
+          <<<blocks, threads, 0, stream>>>(classifier.ins_data,
+                                           classifier.outs_data,
+                                           classifier.use_broadcast,
+                                           numel,
+                                           configs,
+                                           main_offset,
+                                           tail_tid,
+                                           VecSize,
+                                           func);
+    } else if (classifier.broadcast_num > (Arity >> 1)) {
+      constexpr BroadcastType type_ = (Arity > 1) ? kBroadcast : kMixed;
+      VectorizedBroadcastKernel<Functor,
+                                OutT,
+                                IndexT,
+                                Arity,
+                                NumOuts,
+                                VecSize,
+                                type_>
+          <<<blocks, threads, 0, stream>>>(classifier.ins_data,
+                                           classifier.outs_data,
+                                           classifier.use_broadcast,
+                                           numel,
+                                           configs,
+                                           main_offset,
+                                           tail_tid,
+                                           VecSize,
+                                           func);
+    } else {
+      VectorizedBroadcastKernel<Functor,
+                                OutT,
+                                IndexT,
+                                Arity,
+                                NumOuts,
+                                VecSize,
+                                kMixed>
+          <<<blocks, threads, 0, stream>>>(classifier.ins_data,
+                                           classifier.outs_data,
+                                           classifier.use_broadcast,
+                                           numel,
+                                           configs,
+                                           main_offset,
+                                           tail_tid,
+                                           VecSize,
+                                           func);
+    }
+#endif
+  };
+
+#ifdef PADDLE_WITH_XPU_KP
+  // XPU's BroadcastConfig is 32-bit and numel is already capped at INT_MAX in
+  // BroadcastTypeClassifier, so only the uint32 path is ever reachable. Calling
+  // it explicitly avoids instantiating an unused uint64 XPU kernel.
+  launch_with_index_t(uint32_t{0});
+#else
+  // Use the cheaper uint32 indexing (and 32-bit FastDivMod inside
+  // BroadcastConfig) on GPU only when the element count and the grid stride
+  // both fit. The bound is INT_MAX rather than UINT32_MAX because the 32-bit
+  // FastDivMod constructor computes `1 << shift_val` in signed 32-bit and would
+  // overflow for a divisor (an out dim) larger than INT_MAX. The grid_stride
+  // bound mirrors `stride` inside the kernel.
+  const uint64_t index_limit = std::numeric_limits<int32_t>::max();
+  const uint64_t grid_stride =
+      static_cast<uint64_t>(gpu_config.block_per_grid.x) *
+      static_cast<uint64_t>(gpu_config.GetBlockSize()) * VecSize;
+  if (static_cast<uint64_t>(classifier.numel) <= index_limit &&
+      grid_stride <= index_limit) {
+    launch_with_index_t(uint32_t{0});
   } else {
-    VectorizedBroadcastKernel<Functor, OutT, Arity, NumOuts, VecSize, kMixed>
-        <<<blocks, threads, 0, stream>>>(classifier.ins_data,
-                                         classifier.outs_data,
-                                         classifier.use_broadcast,
-                                         numel,
-                                         classifier.configs,
-                                         main_offset,
-                                         tail_tid,
-                                         VecSize,
-                                         func);
+    launch_with_index_t(uint64_t{0});
   }
 #endif
 }
diff --git a/paddle/phi/kernels/primitive/datamover_primitives.h b/paddle/phi/kernels/primitive/datamover_primitives.h
index 92f76f1c806f82..60ef9913faec9a 100644
--- a/paddle/phi/kernels/primitive/datamover_primitives.h
+++ b/paddle/phi/kernels/primitive/datamover_primitives.h
@@ -20,6 +20,8 @@
 #ifdef PADDLE_WITH_HIP
 #include <hip/hip_fp16.h>
 #endif
+#include <type_traits>
+
 #include "paddle/common/ddim.h"
 #include "paddle/common/enforce.h"
 #include "paddle/phi/kernels/funcs/fast_divmod.h"
@@ -40,9 +42,15 @@ struct alignas(sizeof(T) * VecSize) VectorType {
  * index of the output data. if input or output shape is [dim0, dim1] then dims
  * must be [dim1, dim0].
  */
+template <typename IndexT = int64_t>
 struct BroadcastConfig {
-  funcs::FastDivMod<int> divmoders[DDim::kMaxRank];
-  uint64_t strides[DDim::kMaxRank];
+  // FastDivMod has a 32-bit primary template and a 64-bit int64_t
+  // specialization (there is no unsigned specialization), so pick the matching
+  // signed divmod type according to the index width.
+  using DivModT = std::conditional_t<(sizeof(IndexT) > 4), int64_t, int>;
+
+  funcs::FastDivMod<DivModT> divmoders[DDim::kMaxRank];
+  IndexT strides[DDim::kMaxRank];
   int rank{0};
 
   // BroadcastConfig should be defined on host used on device.
@@ -52,8 +60,7 @@ 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<DivModT>(out_dims[i]);
     }
 
     for (int i = 0; i < dim_size; ++i) {
@@ -61,8 +68,8 @@ struct BroadcastConfig {
       strides[i] = (i != 0 && strides[i] != 0)
                        ? std::accumulate(in_dims.begin(),
                                          in_dims.begin() + i,
-                                         1,
-                                         std::multiplies<int64_t>())
+                                         IndexT{1},
+                                         std::multiplies<IndexT>())
                        : strides[i];
     }
     rank = dim_size;
@@ -410,23 +417,25 @@ __device__ __forceinline__ void ReadData(ArgsT* dst,
  * stride_nx: Each read one element stride stride_nx elements in the last dim.
  * stride_ny: Each read one element stride stride_ny elements in the first dim.
  */
-template <typename T, int NX, int NY, bool IsBoundary = false>
+template <typename T, typename IndexT, int NX, int NY, bool IsBoundary = false>
 __device__ __forceinline__ void ReadDataBc(
     T* dst,
     const T* __restrict__ src,
-    uint32_t block_offset,
-    const details::BroadcastConfig& config,
-    int total_num_output,
+    IndexT block_offset,
+    const details::BroadcastConfig<IndexT>& config,
+    IndexT total_num_output,
     int stride_nx,
     int stride_ny) {
-  uint32_t thread_offset = block_offset + threadIdx.x;
-  uint32_t index_src = 0;
+  IndexT thread_offset = block_offset + static_cast<IndexT>(threadIdx.x);
+  IndexT 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;
+      IndexT index_output = thread_offset +
+                            static_cast<IndexT>(ny) * stride_ny +
+                            static_cast<IndexT>(nx) * stride_nx;
       index_src = 0;
       if (IsBoundary) {
         if (index_output >= total_num_output) {
@@ -746,20 +755,20 @@ __device__ __forceinline__ void Init(T* dst, T* init_data, int num) {
  * coordinate mapping relationship between output data and input data.
  * total_num_output: Total number of original output.
  */
-template <typename T, int NX, int NY, bool IsBoundary = false>
+template <typename T, typename IndexT, int NX, int NY, bool IsBoundary = false>
 __device__ __forceinline__ void ReadDataBc(
     T* dst,
     const T* __restrict__ src,
-    uint32_t block_offset,
-    const details::BroadcastConfig& config,
-    int total_num_output,
+    IndexT block_offset,
+    const details::BroadcastConfig<IndexT>& config,
+    IndexT total_num_output,
     int read_lens = NX) {
-  uint32_t thread_offset = block_offset + threadIdx.x * NX;
-  uint32_t index_src = 0;
+  IndexT thread_offset = block_offset + static_cast<IndexT>(threadIdx.x) * NX;
+  IndexT index_src = 0;
 
 #pragma unroll
   for (uint32_t nx = 0; nx < NX; ++nx) {
-    uint32_t index_output = thread_offset + nx;
+    IndexT index_output = thread_offset + nx;
     index_src = 0;
     if (IsBoundary) {
       if (index_output >= total_num_output) {
@@ -803,6 +812,7 @@ __device__ __forceinline__ void ReadDataBc(
  */
 
 template <typename T,
+          typename IndexT,
           int NX,
           int NY,
           typename ArgsT,
@@ -811,16 +821,16 @@ template <typename T,
 __device__ __forceinline__ void ReadDataBc(
     ArgsT* dst,
     const T* __restrict__ src,
-    uint32_t block_offset,
-    const details::BroadcastConfig& config,
-    int total_num_output,
+    IndexT block_offset,
+    const details::BroadcastConfig<IndexT>& config,
+    IndexT total_num_output,
     int read_lens = NX) {
-  uint32_t thread_offset = block_offset + threadIdx.x * NX;
-  uint32_t index_src = 0;
+  IndexT thread_offset = block_offset + static_cast<IndexT>(threadIdx.x) * NX;
+  IndexT index_src = 0;
 
 #pragma unroll
   for (uint32_t nx = 0; nx < NX; ++nx) {
-    uint32_t index_output = thread_offset + nx;
+    IndexT index_output = thread_offset + nx;
     index_src = 0;
     if (IsBoundary) {
       if (index_output >= total_num_output) {