Skip to content

Commit 4a03d25

Browse files
VidithMVidithM
committed
Make row/colscale, select_bitmap more memory-friendly for CUDA
1 parent 19b8df3 commit 4a03d25

7 files changed

Lines changed: 110 additions & 33 deletions

‎CUDA/template/GB_jit_kernel_cuda_colscale.cu‎

Lines changed: 22 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -19,8 +19,10 @@ __global__ void GB_cuda_colscale_kernel
1919
GB_C_TYPE *__restrict__ Cx = (GB_C_TYPE *) C->x ;
2020

2121
#if ( GB_A_IS_SPARSE || GB_A_IS_HYPER )
22+
GB_Ap_TYPE *__restrict__ Cp = (GB_Ap_TYPE *) C->p ;
2223
const GB_Ap_TYPE *__restrict__ Ap = (GB_Ap_TYPE *) A->p ;
2324
#if ( GB_A_IS_HYPER )
25+
GB_Aj_TYPE *__restrict__ Ch = (int64_t *) C->h ;
2426
const GB_Aj_TYPE *__restrict__ Ah = (GB_Aj_TYPE *) A->h ;
2527
#endif
2628
#endif
@@ -31,16 +33,17 @@ __global__ void GB_cuda_colscale_kernel
3133

3234
GB_A_NHELD (anz) ;
3335

36+
int nthreads = blockDim.x * gridDim.x ;
37+
int tid = blockIdx.x * blockDim.x + threadIdx.x ;
38+
3439
#if (GB_A_IS_BITMAP || GB_A_IS_FULL)
3540
const int64_t avlen = A->vlen ;
3641
// bitmap/full case
37-
int nthreads_in_entire_grid = blockDim.x * gridDim.x ;
38-
int tid = blockIdx.x * blockDim.x + threadIdx.x ;
39-
for (int64_t p = tid ; p < anz ; p += nthreads_in_entire_grid)
42+
for (int64_t p = tid ; p < anz ; p += nthreads)
4043
{
4144
if (!GBb_A (Ab, p)) continue ;
4245
// the pth entry in A is A(i,j) where i = p%avlen and j = p/avlen
43-
int64_t col_idx = p / avlen ;
46+
GB_Aj_TYPE col_idx = p / avlen ;
4447
// int64_t row_idx = p % avlen ;
4548
GB_DECLAREB (djj) ;
4649
GB_GETB (djj, Dx, col_idx, ) ;
@@ -51,7 +54,19 @@ __global__ void GB_cuda_colscale_kernel
5154
}
5255

5356
#else
54-
const int64_t anvec = A->nvec ;
57+
const GB_Aj_TYPE anvec = A->nvec ;
58+
// Copy A->p, A->h to C->p, C->h here instead of using
59+
// GB_dup_worker on the CPU, so A->p, A->h stay on the GPU
60+
// if they were already there
61+
for (GB_Aj_TYPE kA = tid ; kA < anvec ; kA += nthreads)
62+
{
63+
Cp [kA] = Ap [kA] ;
64+
#if ( GB_A_IS_HYPER )
65+
Ch [kA] = Ah [kA] ;
66+
#endif
67+
}
68+
Cp [anvec] = Ap [anvec] ;
69+
5570
// sparse/hypersparse case (cuda_ek_slice only works for sparse/hypersparse)
5671
for (int64_t pfirst = blockIdx.x << log2_chunk_size ;
5772
pfirst < anz ;
@@ -65,8 +80,8 @@ __global__ void GB_cuda_colscale_kernel
6580
for (int64_t pdelta = threadIdx.x ; pdelta < my_chunk_size ; pdelta += blockDim.x)
6681
{
6782
int64_t p_final ;
68-
int64_t k = GB_cuda_ek_slice_entry (&p_final, pdelta, pfirst, Ap, anvec_sub1, kfirst, slope) ;
69-
int64_t j = GBh_A (Ah, k) ;
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) ;
7085

7186
GB_DECLAREB (djj) ;
7287
GB_GETB (djj, Dx, j, ) ;

‎CUDA/template/GB_jit_kernel_cuda_rowscale.cu‎

Lines changed: 7 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -16,6 +16,7 @@ __global__ void GB_cuda_rowscale_kernel
1616
#define B_iso GB_B_ISO
1717

1818
#if ( GB_B_IS_SPARSE || GB_B_IS_HYPER )
19+
GB_Bi_TYPE *__restrict__ Ci = (GB_Bi_TYPE *) C->i ;
1920
const GB_Bi_TYPE *__restrict__ Bi = (GB_Bi_TYPE *) B->i ;
2021
#endif
2122

@@ -29,14 +30,18 @@ __global__ void GB_cuda_rowscale_kernel
2930
const int64_t bvlen = B->vlen ;
3031
#endif
3132

32-
int ntasks = gridDim.x * blockDim.x;
33+
int nthreads = gridDim.x * blockDim.x;
3334
int tid = blockIdx.x * blockDim.x + threadIdx.x;
3435

35-
for (int64_t p = tid ; p < bnz ; p += ntasks)
36+
for (int64_t p = tid ; p < bnz ; p += nthreads)
3637
{
3738
if (!GBb_B (Bb, p)) { continue ; }
3839

3940
int64_t i = GBi_B (Bi, p, bvlen) ; // get row index of B(i,j)
41+
#if ( GB_B_IS_SPARSE || GB_B_IS_HYPER )
42+
// Copy B->i to C->i here instead of using GB_dup_worker
43+
Ci [p] = i ;
44+
#endif
4045
GB_DECLAREA (dii) ;
4146
GB_GETA (dii, Dx, i, D_iso) ; // dii = D(i,i)
4247
GB_DECLAREB (bij) ;

‎CUDA/template/GB_jit_kernel_cuda_select_bitmap.cu‎

Lines changed: 9 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -19,10 +19,14 @@ __global__ void GB_cuda_select_bitmap_kernel
1919
{
2020
int8_t *Cb_out = C->b ;
2121

22-
#if ( GB_DEPENDS_ON_X )
22+
#if ( GB_DEPENDS_ON_X || !GB_ISO_SELECT )
2323
const GB_A_TYPE *__restrict__ Ax = (GB_A_TYPE *) A->x ;
2424
#endif
2525

26+
#if ( !GB_ISO_SELECT )
27+
GB_C_TYPE *__restrict__ Cx = (GB_C_TYPE *) C->x ;
28+
#endif
29+
2630
#if ( GB_A_IS_BITMAP )
2731
const int8_t *__restrict__ Ab = A->b ;
2832
#endif
@@ -39,6 +43,10 @@ __global__ void GB_cuda_select_bitmap_kernel
3943
int nthreads = blockDim.x * gridDim.x ;
4044
for (int64_t p = tid ; p < anz ; p += nthreads)
4145
{
46+
#if ( !GB_ISO_SELECT )
47+
Cx [p] = Ax [p] ;
48+
#endif
49+
4250
Cb_out [p] = 0 ;
4351
if (!GBb_A (Ab, p)) { continue; }
4452

‎CUDA/template/GB_jit_kernel_cuda_select_sparse.cu‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -317,7 +317,7 @@ GB_JIT_CUDA_KERNEL_SELECT_SPARSE_PROTO (GB_jit_kernel)
317317

318318
GB_cuda_select_sparse_phase2 <<<grid, block, 0, stream>>>
319319
(Map, A, Ak_keep, (GB_Ci_TYPE *) C->i, (GB_C_TYPE *) C->x) ;
320-
CUDA_OK (cudaGetLastError ( )) ;
320+
// CUDA_OK (cudaGetLastError ( )) ;
321321
CUDA_OK (cudaStreamSynchronize (stream)) ;
322322

323323
//--------------------------------------------------------------------------

‎Source/mxm/GB_colscale.c‎

Lines changed: 39 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -92,21 +92,17 @@ GrB_Info GB_colscale // C = A*D, column scale with diagonal D
9292
GB_void cscalar [GB_VLA(zsize)] ;
9393
bool C_iso = GB_AxB_iso (cscalar, A, D, A->vdim, semiring, flipxy, true) ;
9494

95-
//--------------------------------------------------------------------------
96-
// copy the pattern of A into C
97-
//--------------------------------------------------------------------------
98-
99-
// allocate C->x but do not initialize it
100-
GB_OK (GB_dup_worker (&C, C_iso, A, false, ztype)) ;
10195
info = GrB_NO_VALUE ;
102-
ASSERT (C->type == ztype) ;
10396

10497
//--------------------------------------------------------------------------
10598
// C = A*D, column scale, compute numerical values
10699
//--------------------------------------------------------------------------
107100

108101
if (GB_IS_BUILTIN_BINOP_CODE_POSITIONAL (opcode))
109-
{
102+
{
103+
// Copy the pattern of A into C. Allocates, but does not initialize C->x.
104+
GB_OK (GB_dup_worker (&C, C_iso, A, false, ztype)) ;
105+
ASSERT (C->type == ztype) ;
110106

111107
//----------------------------------------------------------------------
112108
// apply a positional operator: convert C=A*D to C=op(A)
@@ -157,7 +153,10 @@ GrB_Info GB_colscale // C = A*D, column scale with diagonal D
157153

158154
}
159155
else if (C_iso)
160-
{
156+
{
157+
// Copy the pattern of A into C. Allocates, but does not initialize C->x.
158+
GB_OK (GB_dup_worker (&C, C_iso, A, false, ztype)) ;
159+
ASSERT (C->type == ztype) ;
161160

162161
//----------------------------------------------------------------------
163162
// via the iso kernel
@@ -179,6 +178,17 @@ GrB_Info GB_colscale // C = A*D, column scale with diagonal D
179178
// determine if the values are accessed
180179
//----------------------------------------------------------------------
181180

181+
// Do not dup A->p, A->h into C yet; if we use CUDA, we'll do it on the
182+
// GPU
183+
int64_t *tmp_Ap = A->p ;
184+
int64_t *tmp_Ah = A->h ;
185+
A->p = NULL ;
186+
A->h = NULL ;
187+
GB_OK (GB_dup_worker (&C, C_iso, A, false, ztype)) ;
188+
A->p = tmp_Ap ;
189+
A->h = tmp_Ah ;
190+
ASSERT (C->type == ztype) ;
191+
182192
ASSERT (fmult != NULL) ;
183193
bool op_is_first = (opcode == GB_FIRST_binop_code) ;
184194
bool op_is_second = (opcode == GB_SECOND_binop_code) ;
@@ -217,6 +227,26 @@ GrB_Info GB_colscale // C = A*D, column scale with diagonal D
217227
}
218228
#endif
219229

230+
// We are using the CPU. Finish the dup from A -> C.
231+
if (info == GrB_NO_VALUE)
232+
{
233+
// copy A->p, A->h into C->p, C->h
234+
size_t psize = A->p_is_32 ?
235+
sizeof (uint32_t) : sizeof (uint64_t) ;
236+
size_t isize = A->i_is_32 ?
237+
sizeof (uint32_t) : sizeof (uint64_t) ;
238+
int64_t anvec = A->nvec ;
239+
int nthreads_max = GB_Context_nthreads_max ( ) ;
240+
241+
if (A->p != NULL)
242+
{
243+
GB_memcpy (C->p, A->p, (anvec+1) * psize, nthreads_max) ;
244+
}
245+
if (A->h != NULL)
246+
{
247+
GB_memcpy (C->h, A->h, anvec * isize, nthreads_max) ;
248+
}
249+
}
220250
//----------------------------------------------------------------------
221251
// determine the number of threads to use
222252
//----------------------------------------------------------------------

‎Source/mxm/GB_rowscale.c‎

Lines changed: 27 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -81,21 +81,17 @@ GrB_Info GB_rowscale // C = D*B, row scale with diagonal D
8181
GB_void cscalar [GB_VLA(zsize)] ;
8282
bool C_iso = GB_AxB_iso (cscalar, D, B, D->vdim, semiring, flipxy, true) ;
8383

84-
//--------------------------------------------------------------------------
85-
// copy the pattern of B into C
86-
//--------------------------------------------------------------------------
87-
88-
// allocate C->x but do not initialize it
89-
GB_OK (GB_dup_worker (&C, C_iso, B, false, ztype)) ;
9084
info = GrB_NO_VALUE ;
91-
ASSERT (C->type == ztype) ;
9285

9386
//--------------------------------------------------------------------------
9487
// C = D*B, row scale, compute numerical values
9588
//--------------------------------------------------------------------------
9689

9790
if (GB_IS_BUILTIN_BINOP_CODE_POSITIONAL (opcode))
9891
{
92+
// Copy the pattern of B into C. Allocates, but does not initialize C->x.
93+
GB_OK (GB_dup_worker (&C, C_iso, B, false, ztype)) ;
94+
ASSERT (C->type == ztype) ;
9995

10096
//----------------------------------------------------------------------
10197
// apply a positional operator: convert C=D*B to C=op(B)
@@ -147,6 +143,9 @@ GrB_Info GB_rowscale // C = D*B, row scale with diagonal D
147143
}
148144
else if (C_iso)
149145
{
146+
// Copy the pattern of B into C. Allocates, but does not initialize C->x.
147+
GB_OK (GB_dup_worker (&C, C_iso, B, false, ztype)) ;
148+
ASSERT (C->type == ztype) ;
150149

151150
//----------------------------------------------------------------------
152151
// via the iso kernel
@@ -169,6 +168,14 @@ GrB_Info GB_rowscale // C = D*B, row scale with diagonal D
169168
// determine if the values are accessed
170169
//----------------------------------------------------------------------
171170

171+
// Do not dup B->i into C yet; if we use CUDA, we'll do it on the GPU
172+
int64_t *tmp_Bi = B->i ;
173+
B->i = NULL ;
174+
GB_OK (GB_dup_worker (&C, C_iso, B, false, ztype)) ;
175+
B->i = tmp_Bi ;
176+
ASSERT (C->type == ztype) ;
177+
178+
172179
ASSERT (fmult != NULL) ;
173180
bool op_is_first = (opcode == GB_FIRST_binop_code) ;
174181
bool op_is_second = (opcode == GB_SECOND_binop_code) ;
@@ -207,6 +214,19 @@ GrB_Info GB_rowscale // C = D*B, row scale with diagonal D
207214
}
208215
#endif
209216

217+
// We are using the CPU. Finish the dup from B -> C
218+
if (info == GrB_NO_VALUE)
219+
{
220+
// Copy in B->i
221+
int64_t bnz = GB_nnz_held (B) ;
222+
size_t isize = B->i_is_32 ? sizeof (uint32_t) : sizeof (uint64_t) ;
223+
int nthreads_max = GB_Context_nthreads_max ( ) ;
224+
225+
if (B->i != NULL)
226+
{
227+
GB_memcpy (C->i, B->i, bnz * isize, nthreads_max) ;
228+
}
229+
}
210230
//----------------------------------------------------------------------
211231
// determine the number of threads to use
212232
//----------------------------------------------------------------------

‎Source/select/GB_select_bitmap.c‎

Lines changed: 5 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -83,12 +83,6 @@ GrB_Info GB_select_bitmap
8383
// Cx [0] = Ax [0] or (A->type) thunk
8484
GB_select_iso (C->x, opcode, athunk, A->x, asize) ;
8585
}
86-
else
87-
{
88-
// Cx [0:anz-1] = Ax [0:anz-1]
89-
// Fixme for CUDA: do this on the GPU if appropriate
90-
GB_memcpy (C->x, A->x, anz * asize, nthreads) ;
91-
}
9286

9387
//--------------------------------------------------------------------------
9488
// bitmap selector kernel
@@ -105,6 +99,11 @@ GrB_Info GB_select_bitmap
10599

106100
if (info == GrB_NO_VALUE)
107101
{
102+
if (!C_iso) {
103+
// Cx [0:anz-1] = Ax [0:anz-1]
104+
GB_memcpy (C->x, A->x, anz * asize, nthreads) ;
105+
}
106+
108107
if (GB_IS_INDEXUNARYOP_CODE_POSITIONAL (opcode))
109108
{
110109

0 commit comments

Comments
 (0)