1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
373
374
375
376
377
378
379
380
381
382
383
384
385
386
387
388
389
390
391
392
393
394
395
396
397
398
399
400
401
402
403
404
405
406
407
408
409
410
411
412
413
414
415
416
417
418
419
420
421
422
423
424
425
426
427
428
429
430
431
432
433
434
435
436
437
438
439
440
441
442
443
444
445
446
447
448
449
450
451
452
453
454
455
456
457
458
459
460
461
462
463
464
465
466
467
468
469
470
471
472
473
474
475
476
477
478
479
480
481
482
483
484
485
486
487
488
489
490
491
492
493
494
495
496
497
498
499
500
501
502
503
504
505
506
507
508
509
510
511
512
513
514
515
516
517
518
519
520
521
522
523
524
525
526
|
//------------------------------------------------------------------------------
// GB_emult_02: C = A.*B where A is sparse/hyper and B is bitmap/full
//------------------------------------------------------------------------------
// SuiteSparse:GraphBLAS, Timothy A. Davis, (c) 2017-2022, All Rights Reserved.
// SPDX-License-Identifier: Apache-2.0
//------------------------------------------------------------------------------
// C = A.*B where A is sparse/hyper and B is bitmap/full constructs C with
// the same sparsity structure as A. This method can also be called with
// the two input matrices swapped, with flipxy true, to handle the case
// where A is bitmap/full and B is sparse/hyper.
// When no mask is present, or the mask is applied later, this method handles
// the following cases:
// ------------------------------------------
// C = A .* B
// ------------------------------------------
// sparse . sparse bitmap
// sparse . sparse full
// sparse . bitmap sparse
// sparse . full sparse
// If M is sparse/hyper and complemented, it is not passed here:
// ------------------------------------------
// C <!M>= A .* B
// ------------------------------------------
// sparse sparse sparse bitmap (mask later)
// sparse sparse sparse full (mask later)
// sparse sparse bitmap sparse (mask later)
// sparse sparse full sparse (mask later)
// If M is present, it is bitmap/full:
// ------------------------------------------
// C <M> = A .* B
// ------------------------------------------
// sparse bitmap sparse bitmap
// sparse bitmap sparse full
// sparse bitmap bitmap sparse
// sparse bitmap full sparse
// ------------------------------------------
// C <M> = A .* B
// ------------------------------------------
// sparse full sparse bitmap
// sparse full sparse full
// sparse full bitmap sparse
// sparse full full sparse
// ------------------------------------------
// C <!M> = A .* B
// ------------------------------------------
// sparse bitmap sparse bitmap
// sparse bitmap sparse full
// sparse bitmap bitmap sparse
// sparse bitmap full sparse
// ------------------------------------------
// C <!M> = A .* B
// ------------------------------------------
// sparse full sparse bitmap
// sparse full sparse full
// sparse full bitmap sparse
// sparse full full sparse
#include "GB_ewise.h"
#include "GB_emult.h"
#include "GB_binop.h"
#include "GB_unused.h"
#ifndef GBCUDA_DEV
#include "GB_binop__include.h"
#endif
#define GB_FREE_WORKSPACE \
{ \
GB_WERK_POP (Work, int64_t) ; \
GB_WERK_POP (A_ek_slicing, int64_t) ; \
}
#define GB_FREE_ALL \
{ \
GB_FREE_WORKSPACE ; \
GB_phbix_free (C) ; \
}
GrB_Info GB_emult_02 // C=A.*B when A is sparse/hyper, B bitmap/full
(
GrB_Matrix C, // output matrix, static header
const GrB_Type ctype, // type of output matrix C
const bool C_is_csc, // format of output matrix C
const GrB_Matrix M, // optional mask, unused if NULL
const bool Mask_struct, // if true, use the only structure of M
const bool Mask_comp, // if true, use !M
const GrB_Matrix A, // input A matrix (sparse/hyper)
const GrB_Matrix B, // input B matrix (bitmap/full)
GrB_BinaryOp op, // op to perform C = op (A,B)
bool flipxy, // if true use fmult(y,x) else fmult(x,y)
GB_Context Context
)
{
//--------------------------------------------------------------------------
// check inputs
//--------------------------------------------------------------------------
GrB_Info info ;
ASSERT (C != NULL && (C->static_header || GBNSTATIC)) ;
ASSERT_MATRIX_OK_OR_NULL (M, "M for emult_02", GB0) ;
ASSERT_MATRIX_OK (A, "A for emult_02", GB0) ;
ASSERT_MATRIX_OK (B, "B for emult_02", GB0) ;
ASSERT_BINARYOP_OK (op, "op for emult_02", GB0) ;
ASSERT_TYPE_OK (ctype, "ctype for emult_02", GB0) ;
ASSERT (GB_IS_SPARSE (A) || GB_IS_HYPERSPARSE (A)) ;
ASSERT (!GB_PENDING (A)) ;
ASSERT (GB_JUMBLED_OK (A)) ;
ASSERT (!GB_ZOMBIES (A)) ;
ASSERT (GB_IS_BITMAP (B) || GB_IS_FULL (B)) ;
ASSERT (M == NULL || GB_IS_BITMAP (B) || GB_IS_FULL (B)) ;
int C_sparsity = GB_sparsity (A) ;
if (M == NULL)
{
GBURBLE ("emult_02:(%s=%s.*%s)",
GB_sparsity_char (C_sparsity),
GB_sparsity_char_matrix (A),
GB_sparsity_char_matrix (B)) ;
}
else
{
GBURBLE ("emult_02:(%s<%s%s%s>=%s.*%s) ",
GB_sparsity_char (C_sparsity),
Mask_comp ? "!" : "",
GB_sparsity_char_matrix (M),
Mask_struct ? ",struct" : "",
GB_sparsity_char_matrix (A),
GB_sparsity_char_matrix (B)) ;
}
//--------------------------------------------------------------------------
// revise the operator to handle flipxy
//--------------------------------------------------------------------------
// Replace the ANY operator with SECOND. ANY and SECOND give the same
// result if flipxy is false. However, SECOND is changed to FIRST if
// flipxy is true. This ensures that the results do not depend on the
// sparsity structures of A and B.
if (op->opcode == GB_ANY_binop_code)
{
switch (op->xtype->code)
{
case GB_BOOL_code : op = GrB_SECOND_BOOL ; break ;
case GB_INT8_code : op = GrB_SECOND_INT8 ; break ;
case GB_INT16_code : op = GrB_SECOND_INT16 ; break ;
case GB_INT32_code : op = GrB_SECOND_INT32 ; break ;
case GB_INT64_code : op = GrB_SECOND_INT64 ; break ;
case GB_UINT8_code : op = GrB_SECOND_UINT8 ; break ;
case GB_UINT16_code : op = GrB_SECOND_UINT16 ; break ;
case GB_UINT32_code : op = GrB_SECOND_UINT32 ; break ;
case GB_UINT64_code : op = GrB_SECOND_UINT64 ; break ;
case GB_FP32_code : op = GrB_SECOND_FP32 ; break ;
case GB_FP64_code : op = GrB_SECOND_FP64 ; break ;
case GB_FC32_code : op = GxB_SECOND_FC32 ; break ;
case GB_FC64_code : op = GxB_SECOND_FC64 ; break ;
default: ;
}
}
if (flipxy)
{
bool handled ;
op = GB_flip_op (op, &handled) ;
if (handled) flipxy = false ;
}
ASSERT_BINARYOP_OK (op, "final op for emult_02", GB0) ;
//--------------------------------------------------------------------------
// declare workspace
//--------------------------------------------------------------------------
GB_WERK_DECLARE (Work, int64_t) ;
int64_t *restrict Wfirst = NULL ;
int64_t *restrict Wlast = NULL ;
int64_t *restrict Cp_kfirst = NULL ;
GB_WERK_DECLARE (A_ek_slicing, int64_t) ;
//--------------------------------------------------------------------------
// get M, A, and B
//--------------------------------------------------------------------------
const int8_t *restrict Mb = (M == NULL) ? NULL : M->b ;
const GB_void *restrict Mx = (M == NULL || Mask_struct) ? NULL :
(const GB_void *) M->x ;
const size_t msize = (M == NULL) ? 0 : M->type->size ;
const int64_t *restrict Ap = A->p ;
const int64_t *restrict Ah = A->h ;
const int64_t *restrict Ai = A->i ;
const int64_t vlen = A->vlen ;
const int64_t vdim = A->vdim ;
const int64_t nvec = A->nvec ;
const int64_t anz = GB_nnz (A) ;
const int8_t *restrict Bb = B->b ;
const bool B_is_bitmap = GB_IS_BITMAP (B) ;
//--------------------------------------------------------------------------
// check if C is iso and compute its iso value if it is
//--------------------------------------------------------------------------
const size_t csize = ctype->size ;
GB_void cscalar [GB_VLA(csize)] ;
bool C_iso = GB_iso_emult (cscalar, ctype, A, B, op) ;
//--------------------------------------------------------------------------
// allocate C->p and C->h
//--------------------------------------------------------------------------
GB_OK (GB_new (&C, // sparse or hyper (same as A), existing header
ctype, vlen, vdim, GB_Ap_calloc, C_is_csc,
C_sparsity, A->hyper_switch, nvec, Context)) ;
int64_t *restrict Cp = C->p ;
//--------------------------------------------------------------------------
// slice the input matrix A
//--------------------------------------------------------------------------
int A_nthreads, A_ntasks ;
GB_GET_NTHREADS_MAX (nthreads_max, chunk, Context) ;
GB_SLICE_MATRIX (A, 8, chunk) ;
//--------------------------------------------------------------------------
// count entries in C
//--------------------------------------------------------------------------
C->nvec_nonempty = A->nvec_nonempty ;
C->nvec = nvec ;
const bool C_has_pattern_of_A = !B_is_bitmap && (M == NULL) ;
if (!C_has_pattern_of_A)
{
//----------------------------------------------------------------------
// allocate workspace
//----------------------------------------------------------------------
GB_WERK_PUSH (Work, 3*A_ntasks, int64_t) ;
if (Work == NULL)
{
// out of memory
GB_FREE_ALL ;
return (GrB_OUT_OF_MEMORY) ;
}
Wfirst = Work ;
Wlast = Work + A_ntasks ;
Cp_kfirst = Work + A_ntasks * 2 ;
//----------------------------------------------------------------------
// count entries in C
//----------------------------------------------------------------------
// This phase is very similar to GB_select_phase1 (GB_ENTRY_SELECTOR).
if (M == NULL)
{
//------------------------------------------------------------------
// Method2(a): C = A.*B where A is sparse/hyper and B is bitmap
//------------------------------------------------------------------
ASSERT (B_is_bitmap) ;
int tid ;
#pragma omp parallel for num_threads(A_nthreads) schedule(dynamic,1)
for (tid = 0 ; tid < A_ntasks ; tid++)
{
int64_t kfirst = kfirst_Aslice [tid] ;
int64_t klast = klast_Aslice [tid] ;
Wfirst [tid] = 0 ;
Wlast [tid] = 0 ;
for (int64_t k = kfirst ; k <= klast ; k++)
{
// count the entries in C(:,j)
int64_t j = GBH (Ah, k) ;
int64_t pB_start = j * vlen ;
int64_t pA, pA_end ;
GB_get_pA (&pA, &pA_end, tid, k,
kfirst, klast, pstart_Aslice, Ap, vlen) ;
int64_t cjnz = 0 ;
for ( ; pA < pA_end ; pA++)
{
cjnz += Bb [pB_start + Ai [pA]] ;
}
if (k == kfirst)
{
Wfirst [tid] = cjnz ;
}
else if (k == klast)
{
Wlast [tid] = cjnz ;
}
else
{
Cp [k] = cjnz ;
}
}
}
}
else
{
//------------------------------------------------------------------
// Method2(c): C<#M> = A.*B; M, B bitmap/full, A is sparse/hyper
//------------------------------------------------------------------
ASSERT (M != NULL) ;
int tid ;
#pragma omp parallel for num_threads(A_nthreads) schedule(dynamic,1)
for (tid = 0 ; tid < A_ntasks ; tid++)
{
int64_t kfirst = kfirst_Aslice [tid] ;
int64_t klast = klast_Aslice [tid] ;
Wfirst [tid] = 0 ;
Wlast [tid] = 0 ;
for (int64_t k = kfirst ; k <= klast ; k++)
{
// count the entries in C(:,j)
int64_t j = GBH (Ah, k) ;
int64_t pB_start = j * vlen ;
int64_t pA, pA_end ;
GB_get_pA (&pA, &pA_end, tid, k,
kfirst, klast, pstart_Aslice, Ap, vlen) ;
int64_t cjnz = 0 ;
for ( ; pA < pA_end ; pA++)
{
int64_t i = Ai [pA] ;
int64_t pB = pB_start + i ;
bool mij = GBB (Mb, pB) && GB_mcast (Mx, pB, msize) ;
mij = mij ^ Mask_comp ;
cjnz += (mij && GBB (Bb, pB)) ;
}
if (k == kfirst)
{
Wfirst [tid] = cjnz ;
}
else if (k == klast)
{
Wlast [tid] = cjnz ;
}
else
{
Cp [k] = cjnz ;
}
}
}
}
//----------------------------------------------------------------------
// finalize Cp, cumulative sum of Cp and compute Cp_kfirst
//----------------------------------------------------------------------
GB_ek_slice_merge1 (Cp, Wfirst, Wlast, A_ek_slicing, A_ntasks) ;
GB_ek_slice_merge2 (&(C->nvec_nonempty), Cp_kfirst, Cp, nvec,
Wfirst, Wlast, A_ek_slicing, A_ntasks, A_nthreads, Context) ;
}
//--------------------------------------------------------------------------
// allocate C->i and C->x
//--------------------------------------------------------------------------
int64_t cnz = (C_has_pattern_of_A) ? anz : Cp [nvec] ;
// set C->iso = C_iso OK
GB_OK (GB_bix_alloc (C, cnz, GxB_SPARSE, false, true, C_iso, Context)) ;
//--------------------------------------------------------------------------
// copy pattern into C
//--------------------------------------------------------------------------
// TODO: could make these components of C shallow instead of memcpy
if (GB_IS_HYPERSPARSE (A))
{
// copy A->h into C->h
GB_memcpy (C->h, Ah, nvec * sizeof (int64_t), A_nthreads) ;
}
if (C_has_pattern_of_A)
{
// Method2(b): B is full and no mask present, so the pattern of C is
// the same as the pattern of A
GB_memcpy (Cp, Ap, (nvec+1) * sizeof (int64_t), A_nthreads) ;
GB_memcpy (C->i, Ai, cnz * sizeof (int64_t), A_nthreads) ;
}
C->jumbled = A->jumbled ;
C->magic = GB_MAGIC ;
//--------------------------------------------------------------------------
// get the opcode
//--------------------------------------------------------------------------
// if flipxy was true on input and the op is positional, FIRST, SECOND, or
// PAIR, the op has already been flipped, so these tests do not have to
// consider that case.
GB_Opcode opcode = op->opcode ;
bool op_is_positional = GB_OPCODE_IS_POSITIONAL (opcode) ;
bool op_is_first = (opcode == GB_FIRST_binop_code) ;
bool op_is_second = (opcode == GB_SECOND_binop_code) ;
bool op_is_pair = (opcode == GB_PAIR_binop_code) ;
GB_Type_code ccode = ctype->code ;
//--------------------------------------------------------------------------
// check if the values of A and/or B are ignored
//--------------------------------------------------------------------------
// With C = ewisemult (A,B), only the intersection of A and B is used.
// If op is SECOND or PAIR, the values of A are never accessed.
// If op is FIRST or PAIR, the values of B are never accessed.
// If op is PAIR, the values of A and B are never accessed.
// Contrast with ewiseadd.
// A is passed as x, and B as y, in z = op(x,y)
bool A_is_pattern = op_is_second || op_is_pair || op_is_positional ;
bool B_is_pattern = op_is_first || op_is_pair || op_is_positional ;
//--------------------------------------------------------------------------
// using a built-in binary operator (except for positional operators)
//--------------------------------------------------------------------------
#define GB_PHASE_2_OF_2
bool done = false ;
if (C_iso)
{
//----------------------------------------------------------------------
// C is iso
//----------------------------------------------------------------------
// Cx [0] = cscalar = op (A,B)
GB_BURBLE_MATRIX (C, "(iso emult) ") ;
memcpy (C->x, cscalar, csize) ;
// pattern of C = set intersection of pattern of A and B
// flipxy is ignored since the operator is not applied
#define GB_ISO_EMULT
#include "GB_emult_02_template.c"
done = true ;
}
else
{
#ifndef GBCUDA_DEV
//------------------------------------------------------------------
// define the worker for the switch factory
//------------------------------------------------------------------
#define GB_AemultB_02(mult,xname) GB (_AemultB_02_ ## mult ## xname)
#define GB_BINOP_WORKER(mult,xname) \
{ \
info = GB_AemultB_02(mult,xname) (C, \
M, Mask_struct, Mask_comp, A, B, flipxy, \
Cp_kfirst, A_ek_slicing, A_ntasks, A_nthreads) ; \
done = (info != GrB_NO_VALUE) ; \
} \
break ;
//------------------------------------------------------------------
// launch the switch factory
//------------------------------------------------------------------
GB_Type_code xcode, ycode, zcode ;
if (!op_is_positional &&
GB_binop_builtin (A->type, A_is_pattern, B->type, B_is_pattern,
op, false, &opcode, &xcode, &ycode, &zcode) && ccode == zcode)
{
#define GB_NO_PAIR
#include "GB_binop_factory.c"
}
#endif
}
//--------------------------------------------------------------------------
// generic worker
//--------------------------------------------------------------------------
if (!done)
{
GB_BURBLE_MATRIX (C, "(generic emult_02: %s) ", op->name) ;
int ewise_method = flipxy ? GB_EMULT_METHOD3 : GB_EMULT_METHOD2 ;
GB_ewise_generic (C, op, NULL, 0, 0,
NULL, NULL, NULL, C_sparsity, ewise_method, Cp_kfirst,
NULL, 0, 0, A_ek_slicing, A_ntasks, A_nthreads, NULL, 0, 0,
M, Mask_struct, Mask_comp, A, B, Context) ;
}
//--------------------------------------------------------------------------
// remove empty vectors from C, if hypersparse
//--------------------------------------------------------------------------
GB_OK (GB_hypermatrix_prune (C, Context)) ;
//--------------------------------------------------------------------------
// free workspace and return result
//--------------------------------------------------------------------------
GB_FREE_WORKSPACE ;
ASSERT_MATRIX_OK (C, "C output for emult_02", GB0) ;
return (GrB_SUCCESS) ;
}
|
