Skip to content

Commit 587eeaa

Browse files
DrTimothyAldenDavisDrTimothyAldenDavis
committed
revert stable branch to v8.3.0
1 parent 337b580 commit 587eeaa

File tree

259 files changed

+23203
-38366
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

259 files changed

+23203
-38366
lines changed

‎CMakeLists.txt‎

Lines changed: 203 additions & 103 deletions
Large diffs are not rendered by default.

‎CUDA/CMakeLists.txt‎

Lines changed: 44 additions & 34 deletions
Original file line numberDiff line numberDiff line change
@@ -55,6 +55,7 @@ set ( GRAPHBLAS_CUDA_INCLUDES
5555
../Source/Shared
5656
../Source/Template
5757
../Source/Factories
58+
Template
5859
../Include
5960
../CUDA )
6061

@@ -64,18 +65,26 @@ message ( STATUS "GraphBLAS CUDA includes: ${GRAPHBLAS_CUDA_INCLUDES}" )
6465
# GraphBLAS_CUDA properties
6566
#-------------------------------------------------------------------------------
6667

67-
target_include_directories(GraphBLAS_CUDA PRIVATE
68-
${CUDAToolkit_INCLUDE_DIRS}
68+
target_include_directories ( GraphBLAS_CUDA PRIVATE
6969
${GRAPHBLAS_CUDA_INCLUDES})
70-
set_target_properties(GraphBLAS_CUDA PROPERTIES POSITION_INDEPENDENT_CODE ON)
71-
set_target_properties(GraphBLAS_CUDA PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
70+
set_target_properties ( GraphBLAS_CUDA PROPERTIES POSITION_INDEPENDENT_CODE ON )
71+
set_target_properties ( GraphBLAS_CUDA PROPERTIES CUDA_SEPARABLE_COMPILATION ON )
7272
# FIXME: use SUITESPARSE_CUDA_ARCHITECTURES
73-
set_target_properties(GraphBLAS_CUDA PROPERTIES CUDA_ARCHITECTURES "52;75;80" )
73+
set_target_properties ( GraphBLAS_CUDA PROPERTIES CUDA_ARCHITECTURES "52;75;80" )
7474

75-
target_link_libraries(GraphBLAS_CUDA CUDA::nvrtc CUDA::cudart_static CUDA::cuda_driver CUDA::nvToolsExt )
75+
target_link_libraries ( GraphBLAS_CUDA PRIVATE CUDA::nvrtc CUDA::cudart_static CUDA::cuda_driver )
76+
77+
if ( TARGET CUDA::nvtx3 )
78+
target_link_libraries ( GraphBLAS_CUDA PRIVATE CUDA::nvtx3 )
79+
target_compile_definitions ( GraphBLAS_CUDA PRIVATE GBNVTX )
80+
endif ( )
7681

7782
target_compile_definitions ( GraphBLAS_CUDA PUBLIC "SUITESPARSE_CUDA" )
7883

84+
if ( OpenMP_CXX_FOUND )
85+
target_include_directories ( GraphBLAS_CUDA PRIVATE OpenMP::OpenMP_CXX )
86+
endif ( )
87+
7988
target_include_directories ( GraphBLAS_CUDA
8089
INTERFACE $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}>
8190
$<INSTALL_INTERFACE:${SUITESPARSE_INCLUDEDIR}> )
@@ -101,12 +110,12 @@ export ( EXPORT GraphBLAS_CUDATargets
101110
# install export target and config for find_package
102111
install ( EXPORT GraphBLAS_CUDATargets
103112
NAMESPACE SuiteSparse::
104-
DESTINATION ${SUITESPARSE_LIBDIR}/cmake/GraphBLAS )
113+
DESTINATION ${SUITESPARSE_PKGFILEDIR}/cmake/GraphBLAS )
105114

106115
configure_package_config_file (
107116
Config/GraphBLAS_CUDAConfig.cmake.in
108117
${CMAKE_CURRENT_BINARY_DIR}/GraphBLAS_CUDAConfig.cmake
109-
INSTALL_DESTINATION ${SUITESPARSE_LIBDIR}/cmake/GraphBLAS )
118+
INSTALL_DESTINATION ${SUITESPARSE_PKGFILEDIR}/cmake/GraphBLAS )
110119

111120
write_basic_package_version_file (
112121
${CMAKE_CURRENT_BINARY_DIR}/GraphBLAS_CUDAConfigVersion.cmake
@@ -115,7 +124,7 @@ write_basic_package_version_file (
115124
install ( FILES
116125
${CMAKE_CURRENT_BINARY_DIR}/GraphBLAS_CUDAConfig.cmake
117126
${CMAKE_CURRENT_BINARY_DIR}/GraphBLAS_CUDAConfigVersion.cmake
118-
DESTINATION ${SUITESPARSE_LIBDIR}/cmake/GraphBLAS )
127+
DESTINATION ${SUITESPARSE_PKGFILEDIR}/cmake/GraphBLAS )
119128

120129
#-------------------------------------------------------------------------------
121130
# create pkg-config file
@@ -143,13 +152,15 @@ if ( NOT MSVC )
143152
NEWLINE_STYLE LF )
144153
install ( FILES
145154
${CMAKE_CURRENT_BINARY_DIR}/GraphBLAS_CUDA.pc
146-
DESTINATION ${SUITESPARSE_LIBDIR}/pkgconfig )
155+
DESTINATION ${SUITESPARSE_PKGFILEDIR}/pkgconfig )
147156
endif ( )
148157

149158
#-------------------------------------------------------------------------------
150159
# test suite for the CUDA kernels
151160
#-------------------------------------------------------------------------------
152161

162+
if ( 0 )
163+
153164
# 1. Execute enumify/stringify/jitify logic to compile ptx kernels and
154165
# compile/link w/ relevant *.cu files.
155166

@@ -244,28 +255,27 @@ set_target_properties(graphblascuda_test PROPERTIES CUDA_ARCHITECTURES "52;75;80
244255

245256
include(GoogleTest)
246257

247-
add_dependencies(graphblascuda_test GraphBLAS)
248-
add_dependencies(graphblascuda_test GraphBLAS_CUDA)
249-
add_dependencies(graphblascuda_test gtest_main)
250-
add_dependencies(graphblascuda_test rmm_wrap)
251-
252-
target_link_libraries(graphblascuda_test
253-
PUBLIC
254-
GraphBLAS
255-
GraphBLAS_CUDA
256-
rmm_wrap
257-
CUDA::cudart_static
258-
CUDA::nvrtc
259-
${ADDITIONAL_DEPS}
260-
PRIVATE
261-
gtest_main)
262-
263-
target_include_directories(graphblascuda_test
264-
PUBLIC
265-
${CMAKE_CURRENT_SOURCE_DIR}/../Include
266-
${CMAKE_CURRENT_SOURCE_DIR}/../rmm_wrap
267-
rmm_wrap
268-
${ADDITIONAL_INCLUDES}
269-
${CUDAToolkit_INCLUDE_DIRS}
270-
${GRAPHBLAS_CUDA_INCLUDES})
258+
if ( ENABLE_SHARED_LIBS )
259+
target_link_libraries ( graphblascuda_test PUBLIC GraphBLAS )
260+
else ( )
261+
target_link_libraries ( graphblascuda_test PUBLIC GraphBLAS_static )
262+
endif ( )
271263

264+
target_link_libraries ( graphblascuda_test
265+
PUBLIC
266+
GraphBLAS_CUDA
267+
RMM_wrap
268+
CUDA::cudart_static
269+
CUDA::nvrtc
270+
${ADDITIONAL_DEPS}
271+
PRIVATE
272+
gtest_main )
273+
274+
target_include_directories ( graphblascuda_test
275+
PUBLIC
276+
rmm_wrap
277+
${ADDITIONAL_INCLUDES}
278+
${CUDAToolkit_INCLUDE_DIRS}
279+
${GRAPHBLAS_CUDA_INCLUDES} )
280+
281+
endif ( )

‎CUDA/GB_cuda_init.c‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,8 +1,8 @@
11
//------------------------------------------------------------------------------
2-
// GB_cuda_init: initialize the GPUs for use by GraphBLAS
2+
// GraphBLAS/CUDA/GB_cuda_init: initialize the GPUs for use by GraphBLAS
33
//------------------------------------------------------------------------------
44

5-
// SuiteSparse:GraphBLAS, Timothy A. Davis, (c) 2017-2022, All Rights Reserved.
5+
// SuiteSparse:GraphBLAS, Timothy A. Davis, (c) 2017-2023, All Rights Reserved.
66
// SPDX-License-Identifier: Apache-2.0
77

88
//------------------------------------------------------------------------------

‎CUDA/GB_cuda_kernel.h‎

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -75,9 +75,6 @@ typedef struct GB_Semiring_opaque *GrB_Semiring ;
7575
typedef struct GB_Scalar_opaque *GrB_Scalar ;
7676
typedef struct GB_Vector_opaque *GrB_Vector ;
7777
typedef struct GB_Matrix_opaque *GrB_Matrix ;
78-
typedef struct GB_Context_opaque *GxB_Context ;
79-
typedef struct GB_Global_opaque *GrB_Global ;
80-
typedef struct GB_Iterator_opaque *GxB_Iterator ;
8178

8279
#define GxB_HYPERSPARSE 1 // store matrix in hypersparse form
8380
#define GxB_SPARSE 2 // store matrix as sparse form (compressed vector)

‎CUDA/GB_cuda_type_bits.c‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,6 @@
11
// SPDX-License-Identifier: Apache-2.0
22

3-
#undef GBCUDA_CPLUSPLUS
3+
#if 0
44
#include "GB.h"
55

66
size_t GB_cuda_type_bits (GB_Type_code);
@@ -26,3 +26,4 @@ size_t GB_cuda_type_bits (GB_Type_code type_code)
2626
}
2727
}
2828

29+
#endif

‎CUDA/templates/GB_jit_AxB_dot3_phase1.cuh‎

Lines changed: 14 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -80,36 +80,32 @@ __global__ void GB_jit_AxB_dot3_phase1
8080
const bool M_is_hyper = M->h != NULL ;
8181
ASSERT (GB_M_IS_SPARSE || GB_M_IS_HYPER) ;
8282

83+
const int64_t *__restrict__ Ah = A->h ;
8384
const int64_t *__restrict__ Ap = A->p ;
8485
const int64_t *__restrict__ Ai = A->i ;
8586
const int64_t avlen = A->vlen ;
8687
// const int64_t anz = GB_nnz(A) ;
8788
const GB_A_NVALS (anz) ;
8889

90+
const int64_t *__restrict__ Bh = B->h ;
8991
const int64_t *__restrict__ Bp = B->p ;
9092
const int64_t *__restrict__ Bi = B->i ;
9193
const int64_t bvlen = B->vlen ;
9294
// const int64_t bnz = GB_nnz(B);
9395
const GB_B_NVALS (bnz) ;
9496

9597
#if GB_A_IS_HYPER
96-
const int64_t anvec = A->nvec ;
97-
const int64_t *__restrict__ Ah = A->h ;
98-
const int64_t *__restrict__ A_Yp = (A->Y == NULL) ? NULL : A->Y->p ;
99-
const int64_t *__restrict__ A_Yi = (A->Y == NULL) ? NULL : A->Y->i ;
100-
const int64_t *__restrict__ A_Yx = (int64_t *)
101-
((A->Y == NULL) ? NULL : A->Y->x) ;
102-
const int64_t A_hash_bits = (A->Y == NULL) ? 0 : (A->Y->vdim - 1) ;
98+
const int64_t *__restrict__ A_Yp = A->Y->p ;
99+
const int64_t *__restrict__ A_Yi = A->Y->i ;
100+
const int64_t *__restrict__ A_Yx = (int64_t *) A->Y->x ;
101+
const int64_t A_hash_bits = A->Y->vdim - 1 ;
103102
#endif
104103

105104
#if GB_B_IS_HYPER
106-
const int64_t bnvec = B->nvec ;
107-
const int64_t *__restrict__ Bh = B->h ;
108-
const int64_t *__restrict__ B_Yp = (B->Y == NULL) ? NULL : B->Y->p ;
109-
const int64_t *__restrict__ B_Yi = (B->Y == NULL) ? NULL : B->Y->i ;
110-
const int64_t *__restrict__ B_Yx = (int64_t *)
111-
((B->Y == NULL) ? NULL : B->Y->x) ;
112-
const int64_t B_hash_bits = (B->Y == NULL) ? 0 : (B->Y->vdim - 1) ;
105+
const int64_t *__restrict__ B_Yp = B->Y->p ;
106+
const int64_t *__restrict__ B_Yi = B->Y->i ;
107+
const int64_t *__restrict__ B_Yx = (int64_t *) B->Y->x ;
108+
const int64_t B_hash_bits = B->Y->vdim - 1 ;
113109
#endif
114110

115111
// int64_t *restrict Cp = C->p ; // copy of Mp
@@ -224,8 +220,8 @@ __global__ void GB_jit_AxB_dot3_phase1
224220

225221
int64_t pB, pB_end ;
226222
#if GB_B_IS_HYPER
227-
GB_hyper_hash_lookup (Bh, bnvec, Bp, B_Yp, B_Yi, B_Yx,
228-
B_hash_bits, j, &pB, &pB_end) ;
223+
GB_hyper_hash_lookup (Bp, B_Yp, B_Yi, B_Yx, B_hash_bits,
224+
j, &pB, &pB_end) ;
229225
#elif GB_B_IS_SPARSE
230226
pB = Bp[j] ;
231227
pB_end = Bp[j+1] ;
@@ -245,8 +241,8 @@ __global__ void GB_jit_AxB_dot3_phase1
245241

246242
int64_t pA, pA_end ;
247243
#if GB_A_IS_HYPER
248-
GB_hyper_hash_lookup (Ah, anvec, Ap, A_Yp, A_Yi, A_Yx,
249-
A_hash_bits, i, &pA, &pA_end) ;
244+
GB_hyper_hash_lookup (Ap, A_Yp, A_Yi, A_Yx, A_hash_bits,
245+
i, &pA, &pA_end) ;
250246
#elif GB_A_IS_SPARSE
251247
pA = Ap[i] ;
252248
pA_end = Ap[i+1] ;

‎CUDA/templates/GB_jit_AxB_dot3_phase3_mp.cuh‎

Lines changed: 10 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -114,23 +114,17 @@ __global__ void AxB_dot3_phase3_mp
114114
ASSERT (GB_B_IS_HYPER || GB_B_IS_SPARSE) ;
115115

116116
#if GB_A_IS_HYPER
117-
const int64_t anvec = A->nvec ;
118-
const int64_t *__restrict__ Ah = A->h ;
119-
const int64_t *__restrict__ A_Yp = (A->Y == NULL) ? NULL : A->Y->p ;
120-
const int64_t *__restrict__ A_Yi = (A->Y == NULL) ? NULL : A->Y->i ;
121-
const int64_t *__restrict__ A_Yx = (int64_t *)
122-
((A->Y == NULL) ? NULL : A->Y->x) ;
123-
const int64_t A_hash_bits = (A->Y == NULL) ? 0 : (A->Y->vdim - 1) ;
117+
const int64_t *__restrict__ A_Yp = A->Y->p ;
118+
const int64_t *__restrict__ A_Yi = A->Y->i ;
119+
const int64_t *__restrict__ A_Yx = (int64_t *) A->Y->x ;
120+
const int64_t A_hash_bits = A->Y->vdim - 1 ;
124121
#endif
125122

126123
#if GB_B_IS_HYPER
127-
const int64_t bnvec = B->nvec ;
128-
const int64_t *__restrict__ Bh = B->h ;
129-
const int64_t *__restrict__ B_Yp = (B->Y == NULL) ? NULL : B->Y->p ;
130-
const int64_t *__restrict__ B_Yi = (B->Y == NULL) ? NULL : B->Y->i ;
131-
const int64_t *__restrict__ B_Yx = (int64_t *)
132-
((B->Y == NULL) ? NULL : B->Y->x) ;
133-
const int64_t B_hash_bits = (B->Y == NULL) ? 0 : (B->Y->vdim - 1) ;
124+
const int64_t *__restrict__ B_Yp = B->Y->p ;
125+
const int64_t *__restrict__ B_Yi = B->Y->i ;
126+
const int64_t *__restrict__ B_Yx = (int64_t *) B->Y->x ;
127+
const int64_t B_hash_bits = B->Y->vdim - 1 ;
134128
#endif
135129

136130
// zombie count
@@ -168,7 +162,7 @@ __global__ void AxB_dot3_phase3_mp
168162
// find A(:,i)
169163
int64_t pA_start, pA_end ;
170164
#if GB_A_IS_HYPER
171-
GB_hyper_hash_lookup (Ah, anvec, Ap, A_Yp, A_Yi, A_Yx, A_hash_bits,
165+
GB_hyper_hash_lookup (Ap, A_Yp, A_Yi, A_Yx, A_hash_bits,
172166
i, &pA_start, &pA_end) ;
173167
#else
174168
pA_start = Ap[i] ;
@@ -198,7 +192,7 @@ __global__ void AxB_dot3_phase3_mp
198192
// find B(:,j)
199193
int64_t pB_start, pB_end ;
200194
#if GB_B_IS_HYPER
201-
GB_hyper_hash_lookup (Bh, bnvec, Bp, B_Yp, B_Yi, B_Yx, B_hash_bits,
195+
GB_hyper_hash_lookup (Bp, B_Yp, B_Yi, B_Yx, B_hash_bits,
202196
j, &pB_start, &pB_end) ;
203197
#else
204198
pB_start = Bp[j] ;

‎CUDA/templates/GB_jit_AxB_dot3_phase3_spdn.cuh‎

Lines changed: 10 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -113,23 +113,17 @@ __global__ void AxB_dot3_phase3_spdn
113113
#endif
114114

115115
#if GB_A_IS_HYPER
116-
const int64_t anvec = A->nvec ;
117-
const int64_t *__restrict__ Ah = A->h ;
118-
const int64_t *__restrict__ A_Yp = (A->Y == NULL) ? NULL : A->Y->p ;
119-
const int64_t *__restrict__ A_Yi = (A->Y == NULL) ? NULL : A->Y->i ;
120-
const int64_t *__restrict__ A_Yx = (int64_t *)
121-
((A->Y == NULL) ? NULL : A->Y->x) ;
122-
const int64_t A_hash_bits = (A->Y == NULL) ? 0 : (A->Y->vdim - 1) ;
116+
const int64_t *__restrict__ A_Yp = A->Y->p ;
117+
const int64_t *__restrict__ A_Yi = A->Y->i ;
118+
const int64_t *__restrict__ A_Yx = (int64_t *) A->Y->x ;
119+
const int64_t A_hash_bits = A->Y->vdim - 1 ;
123120
#endif
124121

125122
#if GB_B_IS_HYPER
126-
const int64_t bnvec = B->nvec ;
127-
const int64_t *__restrict__ Bh = B->h ;
128-
const int64_t *__restrict__ B_Yp = (B->Y == NULL) ? NULL : B->Y->p ;
129-
const int64_t *__restrict__ B_Yi = (B->Y == NULL) ? NULL : B->Y->i ;
130-
const int64_t *__restrict__ B_Yx = (int64_t *)
131-
((B->Y == NULL) ? NULL : B->Y->x) ;
132-
const int64_t B_hash_bits = (B->Y == NULL) ? 0 : (B->Y->vdim - 1) ;
123+
const int64_t *__restrict__ B_Yp = B->Y->p ;
124+
const int64_t *__restrict__ B_Yi = B->Y->i ;
125+
const int64_t *__restrict__ B_Yx = (int64_t *) B->Y->x ;
126+
const int64_t B_hash_bits = B->Y->vdim - 1 ;
133127
#endif
134128

135129
// zombie count
@@ -157,7 +151,7 @@ __global__ void AxB_dot3_phase3_spdn
157151
// find A(:,i)
158152
int64_t pA, pA_end ;
159153
#if GB_A_IS_HYPER
160-
GB_hyper_hash_lookup (Ah, anvec, Ap, A_Yp, A_Yi, A_Yx, A_hash_bits,
154+
GB_hyper_hash_lookup (Ap, A_Yp, A_Yi, A_Yx, A_hash_bits,
161155
i, &pA, &pA_end) ;
162156
#elif GB_A_IS_SPARSE
163157
pA = Ap[i] ;
@@ -177,7 +171,7 @@ __global__ void AxB_dot3_phase3_spdn
177171
// find B(:,j)
178172
int64_t pB, pB_end ;
179173
#if GB_B_IS_HYPER
180-
GB_hyper_hash_lookup (Bh, bnvec, Bp, B_Yp, B_Yi, B_Yx, B_hash_bits,
174+
GB_hyper_hash_lookup (Bp, B_Yp, B_Yi, B_Yx, B_hash_bits,
181175
j, &pB, &pB_end) ;
182176
#elif GB_B_IS_SPARSE
183177
pB = Bp[j] ;

‎CUDA/templates/GB_jit_AxB_dot3_phase3_vsdn.cuh‎

Lines changed: 10 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -98,23 +98,17 @@ __global__ void AxB_dot3_phase3_vsdn
9898
#endif
9999

100100
#if GB_A_IS_HYPER
101-
const int64_t anvec = A->nvec ;
102-
const int64_t *__restrict__ Ah = A->h ;
103-
const int64_t *__restrict__ A_Yp = (A->Y == NULL) ? NULL : A->Y->p ;
104-
const int64_t *__restrict__ A_Yi = (A->Y == NULL) ? NULL : A->Y->i ;
105-
const int64_t *__restrict__ A_Yx = (int64_t *)
106-
((A->Y == NULL) ? NULL : A->Y->x) ;
107-
const int64_t A_hash_bits = (A->Y == NULL) ? 0 : (A->Y->vdim - 1) ;
101+
const int64_t *__restrict__ A_Yp = A->Y->p ;
102+
const int64_t *__restrict__ A_Yi = A->Y->i ;
103+
const int64_t *__restrict__ A_Yx = (int64_t *) A->Y->x ;
104+
const int64_t A_hash_bits = A->Y->vdim - 1 ;
108105
#endif
109106

110107
#if GB_B_IS_HYPER
111-
const int64_t bnvec = B->nvec ;
112-
const int64_t *__restrict__ Bh = B->h ;
113-
const int64_t *__restrict__ B_Yp = (B->Y == NULL) ? NULL : B->Y->p ;
114-
const int64_t *__restrict__ B_Yi = (B->Y == NULL) ? NULL : B->Y->i ;
115-
const int64_t *__restrict__ B_Yx = (int64_t *)
116-
((B->Y == NULL) ? NULL : B->Y->x) ;
117-
const int64_t B_hash_bits = (B->Y == NULL) ? 0 : (B->Y->vdim - 1) ;
108+
const int64_t *__restrict__ B_Yp = B->Y->p ;
109+
const int64_t *__restrict__ B_Yi = B->Y->i ;
110+
const int64_t *__restrict__ B_Yx = (int64_t *) B->Y->x ;
111+
const int64_t B_hash_bits = B->Y->vdim - 1 ;
118112
#endif
119113

120114
// typedef cub::BlockReduce<int, 32> BlockReduce;
@@ -153,7 +147,7 @@ __global__ void AxB_dot3_phase3_vsdn
153147
// find A(:,i)
154148
int64_t pA, pA_end ;
155149
#if GB_A_IS_HYPER
156-
GB_hyper_hash_lookup (Ah, anvec, Ap, A_Yp, A_Yi, A_Yx, A_hash_bits,
150+
GB_hyper_hash_lookup (Ap, A_Yp, A_Yi, A_Yx, A_hash_bits,
157151
i, &pA, &pA_end) ;
158152
#elif GB_A_IS_SPARSE
159153
pA = Ap[i] ;
@@ -167,7 +161,7 @@ __global__ void AxB_dot3_phase3_vsdn
167161
// find B(:,j)
168162
int64_t pB, pB_end ;
169163
#if GB_B_IS_HYPER
170-
GB_hyper_hash_lookup (Bh, bnvec, Bp, B_Yp, B_Yi, B_Yx, B_hash_bits,
164+
GB_hyper_hash_lookup (Bp, B_Yp, B_Yi, B_Yx, B_hash_bits,
171165
j, &pB, &pB_end) ;
172166
#elif GB_B_IS_SPARSE
173167
pB = Bp[j]; // col of C

0 commit comments

Comments
 (0)