Merge pull request #162 from lukasm91/fix-overflow

Fix overflows up to at least tco1279

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
 

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, &
         & [C_SIZEOF(SRC(START_IN_BYTES:START_IN_BYTES+LENGTH_IN_BYTES-1))/C_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, &
         & [C_SIZEOF(SRC(START_IN_BYTES:START_IN_BYTES+LENGTH_IN_BYTES-1))/C_SIZEOF(DST(0))])

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, c_sizeof
-  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/c_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/c_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/c_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/c_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

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"
@@ -153,9 +156,9 @@ class cutlass_sgemm_grouped<CutlassType::cutlass_fp32, TransA, TransB> {
 template <cublasOperation_t TransA, cublasOperation_t TransB>
 void cutlass_sgemm_wrapper_grouped_op(int resol_id, 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;
@@ -180,9 +183,9 @@ void cutlass_sgemm_wrapper_grouped_op(int resol_id, int blas_id, int m, int *n,
 
 void cutlass_sgemm_wrapper_grouped(int resol_id, 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) {
 

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.
@@ -89,10 +90,10 @@ template <typename Gemm> void erase_from_caches(int resol_id) {
 // this version is using graphs and caches the graphs
 template <typename Gemm, typename Real>
 void run_group_graph(Gemm &&gemm, int resol_id, 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,
-                     int blas_id, void *growing_allocator) {
+                     Real alpha, 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_graph_cache<Gemm>);
 
@@ -163,9 +164,10 @@ void run_group_graph(Gemm &&gemm, int resol_id, int m, int *n, int *k,
 // stupid simple gemm calls
 template <typename Gemm, typename Real>
 void run_group(Gemm &&gemm, int resol_id, 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, int = -1) {
+               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 = -1) {
   for (int i = 0; i < batchCount; ++i) {
     if (m == 0 || n[i] == 0 || k[i] == 0)
       continue;
@@ -213,11 +215,14 @@ template <typename Real> struct hipblas_gemm_grouped {
 
 #ifndef USE_CUTLASS
 
-void hipblas_sgemm_wrapper_grouped(
-    int resol_id, 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,
-    int batchCount, hipStream_t stream, void *growing_allocator) {
+void hipblas_sgemm_wrapper_grouped(int resol_id, int blas_id, char transa,
+                                   char transb, int m, int *n, int *k,
+                                   float alpha, 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) {
 
   hipblasOperation_t op_t1 = HIPBLAS_OP_N, op_t2 = HIPBLAS_OP_N;
   if (transa == 'T' || transa == 't')
@@ -241,9 +246,9 @@ void hipblas_sgemm_wrapper_grouped(
 void hipblas_dgemm_wrapper_grouped(int resol_id, 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, int batchCount,
+                                   int64_t *offsetsA, const double *B, int ldb,
+                                   int64_t *offsetsB, double beta, 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;
@@ -311,10 +316,10 @@ void hipblas_sgemm_wrapper(char transa, char transb, int m, int n, int k,
 void hipblas_sgemm_wrapper_grouped(int resol_id, 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, int batchCount, size_t stream,
-                                   void *growing_allocator) {
+                                   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
   cutlass_sgemm_wrapper_grouped(resol_id, blas_id, transa, transb, m, n, k,
                                 alpha, A, lda, offsetsA, B, ldb, offsetsB, beta,
@@ -331,9 +336,9 @@ void hipblas_sgemm_wrapper_grouped(int resol_id, int blas_id, char transa,
 void hipblas_dgemm_wrapper_grouped(int resol_id, 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, int batchCount,
+                                   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(resol_id, blas_id, transa, transb, m, n, k,
                                 alpha, A, lda, offsetsA, B, ldb, offsetsB, beta,

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
 
@@ -81,10 +81,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  :: RESOL_ID, 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
@@ -101,10 +102,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 :: RESOL_ID, 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
@@ -227,14 +229,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
@@ -274,14 +276,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