@@ -43,7 +43,7 @@ __global__ void GB_cuda_colscale_kernel
4343 {
4444 if (!GBb_A (Ab, p)) continue ;
4545 // the pth entry in A is A(i,j) where i = p%avlen and j = p/avlen
46- GB_Aj_TYPE col_idx = p / avlen ;
46+ int64_t col_idx = p / avlen ;
4747 // int64_t row_idx = p % avlen ;
4848 GB_DECLAREB (djj) ;
4949 GB_GETB (djj, Dx, col_idx, ) ;
@@ -54,11 +54,11 @@ __global__ void GB_cuda_colscale_kernel
5454 }
5555
5656 #else
57- const GB_Aj_TYPE anvec = A->nvec ;
57+ const int64_t anvec = A->nvec ;
5858 // Copy A->p, A->h to C->p, C->h here instead of using
5959 // GB_dup_worker on the CPU, so A->p, A->h stay on the GPU
6060 // if they were already there
61- for (GB_Aj_TYPE kA = tid ; kA < anvec ; kA += nthreads)
61+ for (int64_t kA = tid ; kA < anvec ; kA += nthreads)
6262 {
6363 Cp [kA ] = Ap [kA ] ;
6464 #if ( GB_A_IS_HYPER )
@@ -80,8 +80,8 @@ __global__ void GB_cuda_colscale_kernel
8080 for (int64_t pdelta = threadIdx .x ; pdelta < my_chunk_size ; pdelta += blockDim .x )
8181 {
8282 int64_t p_final ;
83- GB_Aj_TYPE k = GB_cuda_ek_slice_entry (&p_final, pdelta, pfirst, Ap, anvec_sub1, kfirst, slope) ;
84- GB_Aj_TYPE j = GBh_A (Ah, k) ;
83+ int64_t k = GB_cuda_ek_slice_entry (&p_final, pdelta, pfirst, Ap, anvec_sub1, kfirst, slope) ;
84+ int64_t j = GBh_A (Ah, k) ;
8585
8686 GB_DECLAREB (djj) ;
8787 GB_GETB (djj, Dx, j, ) ;
0 commit comments