fix overflows

ecmwf-ifs · Oct 16, 2024 · dd1e3bb · dd1e3bb
1 parent ed88c5a
commit dd1e3bb
Show file tree

Hide file tree

Showing 23 changed files with 325 additions and 241 deletions.
diff --git a/src/trans/common/internal/tpm_distr.F90 b/src/trans/common/internal/tpm_distr.F90
@@ -12,7 +12,7 @@ MODULE TPM_DISTR
 
 ! Module for distributed memory environment.
 
-USE EC_PARKIND  ,ONLY : JPIM     ,JPRD
+USE EC_PARKIND  ,ONLY : JPIM     ,JPRD, JPIB
 
 IMPLICIT NONE
 
@@ -97,7 +97,7 @@ MODULE TPM_DISTR
 INTEGER(KIND=JPIM) :: NDGL_FS ! Number of rows of latitudes for which this process is
                      ! performing Fourier Space calculations
 
-INTEGER(KIND=JPIM) ,ALLOCATABLE  :: NSTAGTF(:) ! Offset for specific latitude in 
+INTEGER(KIND=JPIB) ,ALLOCATABLE  :: NSTAGTF(:) ! Offset for specific latitude in 
                                   ! Fourier/gridpoint buffer
 INTEGER(KIND=JPIM) :: NLENGTF ! Second dimension of Fourier/gridpoint buffer 
                      ! (sum of (NLOEN+3) over local latitudes)
@@ -171,7 +171,7 @@ MODULE TPM_DISTR
 REAL(KIND=JPRD) ,ALLOCATABLE :: RWEIGHT(:) ! Weight per grid-point (if weighted distribution)
 INTEGER(KIND=JPIM) ,ALLOCATABLE :: NPROCA_GP(:) ! Number of grid-points per a-set
 
-INTEGER(KIND=JPIM), ALLOCATABLE :: OFFSETS_GEMM1(:), OFFSETS_GEMM2(:)
+INTEGER(KIND=JPIB), ALLOCATABLE :: OFFSETS_GEMM1(:), OFFSETS_GEMM2(:)
 
 END TYPE DISTR_TYPE
 
@@ -188,7 +188,7 @@ MODULE TPM_DISTR
 INTEGER(KIND=JPIM) ,ALLOCATABLE :: D_NPROCL(:) ! Process responsible for each lat. (F.S)
 INTEGER(KIND=JPIM) ,ALLOCATABLE :: D_NPNTGTB1(:,:)
 INTEGER(KIND=JPIM) ,ALLOCATABLE :: D_NASM0(:)  ! Address in a spectral array of (m, n=m)
-INTEGER(KIND=JPIM) ,ALLOCATABLE  :: D_NSTAGTF(:) ! Offset for specific latitude in 
+INTEGER(KIND=JPIB) ,ALLOCATABLE  :: D_NSTAGTF(:) ! Offset for specific latitude in 
 INTEGER(KIND=JPIM) :: D_NDGL_FS     ! Number of rows of latitudes for which this process is
                                   ! performing Fourier Space calculations
 INTEGER(KIND=JPIM) ,ALLOCATABLE :: D_MSTABF(:)
@@ -200,7 +200,7 @@ MODULE TPM_DISTR
 ! The offsets in the input and output arrays to the gemms.
 ! (1) are the offsets in the "inputs" of dirtrans ("outputs" invtrans)
 ! (2) are the offsets in the "outputs" of invtrans ("inputs" dirtrans)
-INTEGER(KIND=JPIM), POINTER :: D_OFFSETS_GEMM1(:), D_OFFSETS_GEMM2(:)
+INTEGER(KIND=JPIB), POINTER :: D_OFFSETS_GEMM1(:), D_OFFSETS_GEMM2(:)
 
 END MODULE TPM_DISTR
 
diff --git a/src/trans/gpu/algor/buffered_allocator_mod.F90 b/src/trans/gpu/algor/buffered_allocator_mod.F90
@@ -69,10 +69,11 @@ FUNCTION MAKE_BUFFERED_ALLOCATOR()
     MAKE_BUFFERED_ALLOCATOR%NEXT_BUF = 0
   END FUNCTION MAKE_BUFFERED_ALLOCATOR
 
-  FUNCTION RESERVE(ALLOCATOR, SZ)
+  FUNCTION RESERVE(ALLOCATOR, SZ, WHO)
     IMPLICIT NONE
     TYPE(BUFFERED_ALLOCATOR), INTENT(INOUT) :: ALLOCATOR
     INTEGER(KIND=C_SIZE_T), INTENT(IN) :: SZ
+    CHARACTER(*), INTENT(IN), OPTIONAL :: WHO
 
     TYPE(ALLOCATION_RESERVATION_HANDLE) :: RESERVE
 
@@ -88,7 +89,7 @@ SUBROUTINE INSTANTIATE_ALLOCATOR(ALLOCATOR, GROWING_ALLOCATION)
     IMPLICIT NONE
     TYPE(BUFFERED_ALLOCATOR), INTENT(INOUT) :: ALLOCATOR
     !!TYPE(GROWING_ALLOCATION_TYPE), INTENT(IN), POINTER :: GROWING_ALLOCATION
-    TYPE(GROWING_ALLOCATION_TYPE), target, INTENT(INout) :: GROWING_ALLOCATION
+    TYPE(GROWING_ALLOCATION_TYPE), TARGET, INTENT(INOUT) :: GROWING_ALLOCATION
     INTEGER :: I
 
     DO I = 0, NBUF-1
@@ -126,10 +127,13 @@ SUBROUTINE ASSIGN_PTR_FLOAT(DST, SRC, START_IN_BYTES, LENGTH_IN_BYTES, SET_VALUE
     INTEGER(KIND=4), INTENT(IN), OPTIONAL :: SET_STREAM
     LOGICAL :: SET_VALUE_EFF
     INTEGER(KIND=4) :: SET_STREAM_EFF
-    INTEGER(KIND=C_SIZE_T) :: START_IN_BYTES, LENGTH_IN_BYTES
+    INTEGER(KIND=C_SIZE_T) :: START_IN_BYTES, LENGTH_IN_BYTES, END_IN_BYTES, J
     IF (START_IN_BYTES + LENGTH_IN_BYTES - 1 > SIZE(SRC, KIND=C_SIZE_T)) THEN
       CALL ABORT_TRANS("Logical Error in ASSIGN_PTR - OOB assignment")
     ENDIF
+    IF (START_IN_BYTES < 1) THEN
+      CALL ABORT_TRANS("Logical Error in ASSIGN_PTR - OOB assignment")
+    ENDIF
     IF (PRESENT(SET_VALUE)) THEN
       SET_VALUE_EFF = SET_VALUE
     ELSE
@@ -143,9 +147,11 @@ SUBROUTINE ASSIGN_PTR_FLOAT(DST, SRC, START_IN_BYTES, LENGTH_IN_BYTES, SET_VALUE
     IF (SET_VALUE_EFF .AND. LENGTH_IN_BYTES > 0) THEN
       ! This option is turned off by default, but for experimentation we can turn it on. This is
       ! setting all bits to 1 (meaning NaN in floating point)
-      !$ACC KERNELS PRESENT(SRC) ASYNC(SET_STREAM_EFF)
-      SRC(START_IN_BYTES:START_IN_BYTES+LENGTH_IN_BYTES-1) = -1
-      !$ACC END KERNELS!! LOOP
+      !$ACC PARALLEL PRESENT(SRC) ASYNC(SET_STREAM_EFF)
+      DO J=1_C_SIZE_T,LENGTH_IN_BYTES
+        SRC(J) = -1
+      ENDDO
+      !$ACC END PARALLEL
     ENDIF
     CALL C_F_POINTER(C_LOC(SRC(START_IN_BYTES:START_IN_BYTES+LENGTH_IN_BYTES-1)), DST, &
         & [SIZEOF(SRC(START_IN_BYTES:START_IN_BYTES+LENGTH_IN_BYTES-1))/SIZEOF(DST(0))])
@@ -159,10 +165,13 @@ SUBROUTINE ASSIGN_PTR_DOUBLE(DST, SRC, START_IN_BYTES, LENGTH_IN_BYTES, SET_VALU
     INTEGER(KIND=4), INTENT(IN), OPTIONAL :: SET_STREAM
     LOGICAL :: SET_VALUE_EFF
     INTEGER(KIND=4) :: SET_STREAM_EFF
-    INTEGER(KIND=C_SIZE_T) :: START_IN_BYTES, LENGTH_IN_BYTES
+    INTEGER(KIND=C_SIZE_T) :: START_IN_BYTES, LENGTH_IN_BYTES, END_IN_BYTES, J
     IF (START_IN_BYTES + LENGTH_IN_BYTES - 1 > SIZE(SRC, KIND=C_SIZE_T)) THEN
       CALL ABORT_TRANS("Logical Error in ASSIGN_PTR - OOB assignment")
     ENDIF
+    IF (START_IN_BYTES < 1) THEN
+      CALL ABORT_TRANS("Logical Error in ASSIGN_PTR - OOB assignment")
+    ENDIF
     IF (PRESENT(SET_VALUE)) THEN
       SET_VALUE_EFF = SET_VALUE
     ELSE
@@ -176,9 +185,12 @@ SUBROUTINE ASSIGN_PTR_DOUBLE(DST, SRC, START_IN_BYTES, LENGTH_IN_BYTES, SET_VALU
     IF (SET_VALUE_EFF .AND. LENGTH_IN_BYTES > 0) THEN
       ! This option is turned off by default, but for experimentation we can turn it on. This is
       ! setting all bits to 1 (meaning NaN in floating point)
-      !$ACC KERNELS PRESENT(SRC) ASYNC(SET_STREAM_EFF)
-      SRC(START_IN_BYTES:START_IN_BYTES+LENGTH_IN_BYTES-1) = -1
-      !$ACC END KERNELS!! LOOP
+      END_IN_BYTES=START_IN_BYTES+LENGTH_IN_BYTES-1
+      !$ACC PARALLEL PRESENT(SRC) ASYNC(SET_STREAM_EFF)
+      DO J=1_C_SIZE_T,LENGTH_IN_BYTES
+        SRC(J) = -1
+      ENDDO
+      !$ACC END PARALLEL
     ENDIF
     CALL C_F_POINTER(C_LOC(SRC(START_IN_BYTES:START_IN_BYTES+LENGTH_IN_BYTES-1)), DST, &
         & [SIZEOF(SRC(START_IN_BYTES:START_IN_BYTES+LENGTH_IN_BYTES-1))/SIZEOF(DST(0))])

diff --git a/src/trans/gpu/algor/ext_acc.F90 b/src/trans/gpu/algor/ext_acc.F90
@@ -18,7 +18,7 @@ module openacc_ext_type
 end module
 module openacc_ext
   use iso_c_binding, only: c_ptr, c_size_t, c_loc
-  use openacc, only: acc_create, acc_copyin, acc_handle_kind
+  use openacc, only: acc_handle_kind
   use openacc_ext_type, only: ext_acc_arr_desc
   implicit none
 
@@ -247,7 +247,7 @@ function get_common_pointers(in_ptrs, out_ptrs) result(num_ranges)
     enddo
   end function
   subroutine ext_acc_create(ptrs, stream)
-    use openacc, only: acc_create, acc_async_sync
+    use openacc, only: acc_async_sync
     use iso_fortran_env, only: int32
     implicit none
     type(ext_acc_arr_desc), intent(in) :: ptrs(:)
@@ -269,8 +269,7 @@ subroutine ext_acc_create(ptrs, stream)
 
     do i = 1, num_ranges
       call c_f_pointer(common_ptrs(i)%ptr, pp, shape=[common_ptrs(i)%sz/sizeof(pp(1))])
-      !!call acc_create_async(pp, common_ptrs(i)%sz, async=stream_act)
-      call acc_create(pp, int(common_ptrs(i)%sz))
+      !$acc enter data create(pp) async(stream_act)
     enddo
   end subroutine
   subroutine ext_acc_copyin(ptrs, stream)
@@ -296,12 +295,11 @@ subroutine ext_acc_copyin(ptrs, stream)
 
     do i = 1, num_ranges
       call c_f_pointer(common_ptrs(i)%ptr, pp, shape=[common_ptrs(i)%sz/sizeof(pp(1))])
-      !!call acc_copyin_async(pp, common_ptrs(i)%sz, async=stream_act)
-      call acc_copyin(pp, int(common_ptrs(i)%sz))
+      !$acc enter data copyin(pp) async(stream_act)
     enddo
   end subroutine
   subroutine ext_acc_copyout(ptrs, stream)
-    use openacc, only: acc_async_sync, acc_copyout
+    use openacc, only: acc_async_sync
     implicit none
     type(ext_acc_arr_desc), intent(in) :: ptrs(:)
     integer(acc_handle_kind), optional :: stream
@@ -323,12 +321,11 @@ subroutine ext_acc_copyout(ptrs, stream)
 
     do i = 1, num_ranges
       call c_f_pointer(common_ptrs(i)%ptr, pp, shape=[common_ptrs(i)%sz/sizeof(pp(1))])
-      !!call acc_copyout_async(pp, common_ptrs(i)%sz, async=stream_act)
-      call acc_copyout(pp, int(common_ptrs(i)%sz))
+      !$acc exit data copyout(pp) async(stream_act)
     enddo
   end subroutine
   subroutine ext_acc_delete(ptrs, stream)
-    use openacc, only: acc_async_sync, acc_delete
+    use openacc, only: acc_async_sync
     implicit none
     type(ext_acc_arr_desc), intent(in) :: ptrs(:)
     integer(acc_handle_kind), optional :: stream
@@ -350,8 +347,7 @@ subroutine ext_acc_delete(ptrs, stream)
 
     do i = 1, num_ranges
       call c_f_pointer(common_ptrs(i)%ptr, pp, shape=[common_ptrs(i)%sz/sizeof(pp(1))])
-      !!call acc_delete_async(pp, common_ptrs(i)%sz, async=stream_act)
-      call acc_delete(pp, int(common_ptrs(i)%sz))
+      !$acc exit data delete(pp) async(stream_act)
     enddo
   end subroutine
 end module
diff --git a/src/trans/gpu/algor/hicblas_cutlass.cuda.h b/src/trans/gpu/algor/hicblas_cutlass.cuda.h
@@ -1,3 +1,6 @@
+// (C) Copyright 2000- ECMWF.
+// (C) Copyright 2024- NVIDIA.
+
 #ifdef USE_CUTLASS
 //#include "hicblas.h"
 #include "cutlass/gemm/device/gemm.h"
@@ -151,9 +154,9 @@ class cutlass_sgemm_grouped<CutlassType::cutlass_fp32, TransA, TransB> {
 template <cublasOperation_t TransA, cublasOperation_t TransB>
 void cutlass_sgemm_wrapper_grouped_op(int blas_id, int m, int *n, int *k,
                                       float alpha, const float *A, int lda,
-                                      int *offsetsA, const float *B, int ldb,
-                                      int *offsetsB, float beta, float *C,
-                                      int ldc, int *offsetsC, int batchCount,
+                                      int64_t *offsetsA, const float *B, int ldb,
+                                      int64_t *offsetsB, float beta, float *C,
+                                      int ldc, int64_t *offsetsC, int batchCount,
                                       cudaStream_t stream,
                                       void *growing_allocator) {
   using namespace detail;
@@ -178,9 +181,9 @@ void cutlass_sgemm_wrapper_grouped_op(int blas_id, int m, int *n, int *k,
 
 void cutlass_sgemm_wrapper_grouped(int blas_id, char transa, char transb,
                                    int m, int *n, int *k, float alpha,
-                                   const float *A, int lda, int *offsetsA,
-                                   const float *B, int ldb, int *offsetsB, float beta,
-                                   float *C, int ldc, int *offsetsC,
+                                   const float *A, int lda, int64_t *offsetsA,
+                                   const float *B, int ldb, int64_t *offsetsB, float beta,
+                                   float *C, int ldc, int64_t *offsetsC,
                                    int batchCount, cudaStream_t stream,
                                    void *growing_allocator) {
 

diff --git a/src/trans/gpu/algor/hicblas_gemm.hip.cpp b/src/trans/gpu/algor/hicblas_gemm.hip.cpp
@@ -1,4 +1,5 @@
 // (C) Copyright 2000- ECMWF.
+// (C) Copyright 2024- NVIDIA.
 //
 // This software is licensed under the terms of the Apache Licence Version 2.0
 // which can be obtained at http://www.apache.org/licenses/LICENSE-2.0.
@@ -64,9 +65,9 @@ template <typename Gemm, typename Real> void free_gemm_cache(float *, size_t) {
 // this version is using graphs and caches the graphs
 template <typename Gemm, typename Real>
 void run_group_graph(Gemm &&gemm, int m, int *n, int *k, Real alpha,
-                     const Real *A, int lda, int *offsetsA, const Real *B,
-                     int ldb, int *offsetsB, Real beta, Real *C, int ldc,
-                     int *offsetsC, int batchCount, hipStream_t stream,
+                     const Real *A, int lda, int64_t *offsetsA, const Real *B,
+                     int ldb, int64_t *offsetsB, Real beta, Real *C, int ldc,
+                     int64_t *offsetsC, int batchCount, hipStream_t stream,
                      int blas_id, void *growing_allocator) {
   growing_allocator_register_free_c(growing_allocator,
                                     free_gemm_cache<Gemm, Real>);
@@ -133,8 +134,8 @@ void run_group_graph(Gemm &&gemm, int m, int *n, int *k, Real alpha,
 // stupid simple gemm calls
 template <typename Gemm, typename Real>
 void run_group(Gemm &&gemm, int m, int *n, int *k, Real alpha, const Real *A,
-               int lda, int *offsetsA, const Real *B, int ldb, int *offsetsB,
-               Real beta, Real *C, int ldc, int *offsetsC, int batchCount,
+               int lda, int64_t *offsetsA, const Real *B, int ldb, int64_t *offsetsB,
+               Real beta, Real *C, int ldc, int64_t *offsetsC, int batchCount,
                hipStream_t stream, int = -1) {
   for (int i = 0; i < batchCount; ++i) {
     if (m == 0 || n[i] == 0 || k[i] == 0)
@@ -186,9 +187,9 @@ template <typename Real> struct hipblas_gemm_grouped {
 
 void hipblas_sgemm_wrapper_grouped(int blas_id, char transa, char transb,
                                    int m, int *n, int *k, float alpha,
-                                   const float *A, int lda, int *offsetsA,
-                                   const float *B, int ldb, int *offsetsB, float beta,
-                                   float *C, int ldc, int *offsetsC,
+                                   const float *A, int lda, int64_t *offsetsA,
+                                   const float *B, int ldb, int64_t *offsetsB, float beta,
+                                   float *C, int ldc, int64_t *offsetsC,
                                    int batchCount, hipStream_t stream,
                                   void *growing_allocator) {
 
@@ -215,10 +216,10 @@ void hipblas_sgemm_wrapper_grouped(int blas_id, char transa, char transb,
 void hipblas_dgemm_wrapper_grouped(int blas_id, char transa, char transb,
                                    int m, int *n, int *k,
                                    double alpha,
-                                   const double *A, int lda, int *offsetsA,
-                                   const double *B, int ldb, int *offsetsB,
+                                   const double *A, int lda, int64_t *offsetsA,
+                                   const double *B, int ldb, int64_t *offsetsB,
                                    double beta,
-                                   double *C, int ldc, int *offsetsC,
+                                   double *C, int ldc, int64_t *offsetsC,
                                    int batchCount, hipStream_t stream, void *) {
 
   hipblasOperation_t op_t1=HIPBLAS_OP_N, op_t2=HIPBLAS_OP_N;
@@ -292,9 +293,9 @@ void hipblas_sgemm_wrapper (char transa, char transb,
 
 void hipblas_sgemm_wrapper_grouped(int blas_id, char transa, char transb,
                                 int m, int *n, int *k, float alpha,
-                                const float *A, int lda, int *offsetsA,
-                                const float *B, int ldb, int *offsetsB, float beta,
-                                float *C, int ldc, int *offsetsC,
+                                const float *A, int lda, int64_t *offsetsA,
+                                const float *B, int ldb, int64_t *offsetsB, float beta,
+                                float *C, int ldc, int64_t *offsetsC,
                                 int batchCount, size_t stream,
                                 void *growing_allocator) {
 #ifdef USE_CUTLASS
@@ -313,9 +314,9 @@ void hipblas_sgemm_wrapper_grouped(int blas_id, char transa, char transb,
 
 void hipblas_dgemm_wrapper_grouped(int blas_id, char transa, char transb,
                                 int m, int *n, int *k, double alpha,
-                                const double *A, int lda, int *offsetsA,
-                                const double *B, int ldb, int *offsetsB, double beta,
-                                double *C, int ldc, int *offsetsC,
+                                const double *A, int lda, int64_t *offsetsA,
+                                const double *B, int ldb, int64_t *offsetsB, double beta,
+                                double *C, int ldc, int64_t *offsetsC,
                                 int batchCount, size_t stream,
                                 void *growing_allocator) {
     hipblas_dgemm_wrapper_grouped(blas_id, transa, transb, m, n, k, alpha, A, lda, offsetsA, B,

diff --git a/src/trans/gpu/algor/hicblas_mod.F90 b/src/trans/gpu/algor/hicblas_mod.F90
@@ -14,7 +14,7 @@
 
 MODULE HICBLAS_MOD
 
-USE EC_PARKIND, ONLY: JPIM, JPRM, JPRD
+USE EC_PARKIND, ONLY: JPIM, JPRM, JPRD, JPIB
 USE GROWING_ALLOCATOR_MOD, ONLY: GROWING_ALLOCATION_TYPE
 USE OPENACC_LIB, ONLY: ACC_GET_HIP_STREAM
 
@@ -75,10 +75,11 @@ SUBROUTINE HIP_DGEMM_GROUPED( &
     & C, LDC, OFFSETC,          &
     & BATCHCOUNT, STREAM, ALLOC &
   &) BIND(C, NAME='hipblas_dgemm_wrapper_grouped')
-    USE ISO_C_BINDING, ONLY: C_CHAR, C_INT, C_DOUBLE, C_SIZE_T, C_PTR
+    USE ISO_C_BINDING, ONLY: C_CHAR, C_INT, C_DOUBLE, C_SIZE_T, C_PTR, C_INT64_T
     CHARACTER(1,C_CHAR), VALUE :: CTA, CTB
     INTEGER(C_INT), VALUE  :: BLAS_ID, M, LDA, LDB, LDC, BATCHCOUNT
-    INTEGER(C_INT)         :: N(*), K(*), OFFSETA(*), OFFSETB(*), OFFSETC(*)
+    INTEGER(C_INT)         :: N(*), K(*)
+    INTEGER(C_INT64_T)     :: OFFSETA(*), OFFSETB(*), OFFSETC(*)
     REAL(C_DOUBLE), VALUE  :: ALPHA,BETA
     REAL(C_DOUBLE)         :: A(*), B(*), C(*)
     INTEGER(KIND=C_SIZE_T) :: STREAM
@@ -95,10 +96,11 @@ SUBROUTINE HIP_SGEMM_GROUPED( &
     & C, LDC, OFFSETC,          &
     & BATCHCOUNT, STREAM, ALLOC &
   &) BIND(C, NAME='hipblas_sgemm_wrapper_grouped')
-    USE ISO_C_BINDING, ONLY: C_CHAR, C_INT, C_FLOAT, C_SIZE_T, C_PTR
+    USE ISO_C_BINDING, ONLY: C_CHAR, C_INT, C_FLOAT, C_SIZE_T, C_PTR, C_INT64_T
     CHARACTER(1,C_CHAR), VALUE :: CTA, CTB
     INTEGER(C_INT), VALUE :: BLAS_ID, M, LDA, LDB, LDC, BATCHCOUNT
-    INTEGER(C_INT)        :: N(*), K(*), OFFSETA(*), OFFSETB(*), OFFSETC(*)
+    INTEGER(C_INT)        :: N(*), K(*)
+    INTEGER(C_INT64_T)    :: OFFSETA(*), OFFSETB(*), OFFSETC(*)
     REAL(C_FLOAT), VALUE  :: ALPHA,BETA
     REAL(C_FLOAT)         :: A(*), B(*), C(*)
     INTEGER(KIND=C_SIZE_T) :: STREAM
@@ -220,14 +222,14 @@ SUBROUTINE HIP_DGEMM_GROUPED_OVERLOAD( &
   REAL(KIND=JPRD) :: ALPHA
   REAL(KIND=JPRD), DIMENSION(:) :: AARRAY
   INTEGER(KIND=JPIM) :: LDA
-  INTEGER(KIND=JPIM) :: OFFSETA(:)
+  INTEGER(KIND=JPIB) :: OFFSETA(:)
   REAL(KIND=JPRD), DIMENSION(*) :: BARRAY
   INTEGER(KIND=JPIM) :: LDB
-  INTEGER(KIND=JPIM) :: OFFSETB(:)
+  INTEGER(KIND=JPIB) :: OFFSETB(:)
   REAL(KIND=JPRD) :: BETA
   REAL(KIND=JPRD), DIMENSION(:) :: CARRAY
   INTEGER(KIND=JPIM) :: LDC
-  INTEGER(KIND=JPIM) :: OFFSETC(:)
+  INTEGER(KIND=JPIB) :: OFFSETC(:)
   INTEGER(KIND=JPIM) :: BATCHCOUNT
   INTEGER(KIND=C_INT) :: STREAM
   TYPE(GROWING_ALLOCATION_TYPE), INTENT(IN) :: ALLOC
@@ -266,14 +268,14 @@ SUBROUTINE HIP_SGEMM_GROUPED_OVERLOAD(&
   REAL(KIND=JPRM) :: ALPHA
   REAL(KIND=JPRM), DIMENSION(:) :: AARRAY
   INTEGER(KIND=JPIM) :: LDA
-  INTEGER(KIND=JPIM) :: OFFSETA(:)
+  INTEGER(KIND=JPIB) :: OFFSETA(:)
   REAL(KIND=JPRM), DIMENSION(:,:,:) :: BARRAY
   INTEGER(KIND=JPIM) :: LDB
-  INTEGER(KIND=JPIM) :: OFFSETB(:)
+  INTEGER(KIND=JPIB) :: OFFSETB(:)
   REAL(KIND=JPRM) :: BETA
   REAL(KIND=JPRM), DIMENSION(:) :: CARRAY
   INTEGER(KIND=JPIM) :: LDC
-  INTEGER(KIND=JPIM) :: OFFSETC(:)
+  INTEGER(KIND=JPIB) :: OFFSETC(:)
   INTEGER(KIND=JPIM) :: BATCHCOUNT
   INTEGER(KIND=C_INT) :: STREAM
   TYPE(GROWING_ALLOCATION_TYPE), INTENT(IN) :: ALLOC