From b9bd4bc0b46c8de8d31eb84a9120032875c07ce6 Mon Sep 17 00:00:00 2001 From: Richard Lettich Date: Mon, 21 Jun 2021 19:09:59 -0700 Subject: [PATCH 01/12] Integration WIP, it builds --- .gitmodules | 3 + CMakeLists.txt | 8 +- Makefile | 2 +- build/outtt | 385 ++++++++++++++++++++++++ ext/GALATIC | 1 + ext/cub | 1 - graphblas/backend/cuda/operations.hpp | 7 +- graphblas/backend/cuda/reduce.hpp | 2 +- graphblas/backend/cuda/spgemm.hpp | 174 +++++++++++ graphblas/backend/cuda/spmspv_inner.hpp | 2 +- graphblas/backend/cuda/spmv.hpp | 2 +- graphblas/stddef.hpp | 12 +- 12 files changed, 584 insertions(+), 15 deletions(-) create mode 100644 build/outtt create mode 160000 ext/GALATIC delete mode 160000 ext/cub diff --git a/.gitmodules b/.gitmodules index 792535b..e1d59bf 100644 --- a/.gitmodules +++ b/.gitmodules @@ -4,3 +4,6 @@ [submodule "ext/cub"] path = ext/cub url = https://ctcyang@github.com/NVlabs/cub.git +[submodule "ext/GALATIC"] + path = ext/GALATIC + url = git@github.com:richardlett/GALATIC.git diff --git a/CMakeLists.txt b/CMakeLists.txt index 9e7d002..dad5c8d 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1,7 +1,7 @@ # CMakeLists.txt for unittest # required cmakeversion -cmake_minimum_required(VERSION 2.8) +cmake_minimum_required(VERSION 3.20) # packages find_package(CUDA) find_package(Boost COMPONENTS program_options REQUIRED) @@ -13,7 +13,7 @@ set( PROJ_PATH ${CMAKE_SOURCE_DIR}) set( PROJ_OUT_PATH ${CMAKE_BINARY_DIR}) set( PROJ_HEADERS "" ) set( PROJ_LIBRARIES "" ) -set( PROJ_INCLUDES "./" "ext/moderngpu/include" "ext/cub/cub") +set( PROJ_INCLUDES "./" "ext/moderngpu/include") set( mgpu_SRC_FILES "ext/moderngpu/src/mgpucontext.cu" "ext/moderngpu/src/mgpuutil.cpp") set( CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/../bin ) #set( CUDA_CURAND_LIBRARY "$ENV{CUDA_HOME}/lib64/libcurand.so" ) @@ -23,12 +23,12 @@ set( CUDA_CUSPARSE_LIBRARY "$ENV{CUDA_HOME}/lib64/libcusparse.so" ) #FILE( GLOB_RECURSE PROJ_LIBRARIES ext/cublas1.1/*.cu ) FILE( GLOB_RECURSE PROJ_HEADERS graphblas/*.hpp) # nvcc flags -set(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS} -arch=sm_35 -lineinfo -O3 -use_fast_math -Xptxas=-v") +set(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS} -arch=sm_50 -lineinfo -O3 -use_fast_math -Xptxas=-v --expt-relaxed-constexpr ") #set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-fpermissive;-arch=sm_35;-lineinfo;-Xptxas=-v;-dlcm=ca;-maxrregcount=64) #set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-gencode arch=compute_20,code=sm_21) # needed for cudamalloc set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS}") -set(CMAKE_CXX_FLAGS "-fpermissive -g -m64 -std=c++11" ) +set(CMAKE_CXX_FLAGS "-fpermissive -g -m64 -std=c++14" ) #set(CMAKE_CXX_FLAGS "-fpermissive -pg -m64 -std=c++11" ) #set(CMAKE_CXX_FLAGS "-fpermissive -g -m64 -std=c++11 -H" ) set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${OpenMP_CXX_FLAGS}") diff --git a/Makefile b/Makefile index b683a75..bf0e8e7 100644 --- a/Makefile +++ b/Makefile @@ -5,7 +5,7 @@ include common.mk #------------------------------------------------------------------------------- # Includes -INC += -I$(MGPU_DIR) -I$(CUB_DIR) -I$(BOOST_DIR) -I$(GRB_DIR) +INC += -I$(MGPU_DIR) -I$(BOOST_DIR) -I$(GRB_DIR) #------------------------------------------------------------------------------- # Dependency Lists diff --git a/build/outtt b/build/outtt new file mode 100644 index 0000000..17b909b --- /dev/null +++ b/build/outtt @@ -0,0 +1,385 @@ +loading A +Undirected due to mtx: 1 +Undirected due to cmd: 0 +Undirected: 1 +Remove self-loop: 1 +Reading ../data/small/.chesapeake.mtx.ud.nosl.bin +Allocate 40 +Allocate 408 +Allocate 408 +csrColInd: +[0]:6 [1]:7 [2]:10 [3]:11 [4]:12 [5]:21 [6]:22 [7]:33 [8]:34 [9]:36 [10]:38 [11]:6 [12]:7 [13]:8 [14]:10 [15]:11 [16]:12 [17]:21 [18]:22 [19]:34 [20]:35 [21]:38 [22]:13 [23]:14 [24]:15 [25]:16 [26]:17 [27]:35 [28]:38 [29]:16 [30]:35 [31]:36 [32]:38 [33]:5 [34]:33 [35]:34 [36]:38 [37]:4 [38]:6 [39]:34 +csrRowPtr: +[0]:0 [1]:11 [2]:22 [3]:29 [4]:33 [5]:37 [6]:41 [7]:51 [8]:64 [9]:71 [10]:76 [11]:83 [12]:92 [13]:99 [14]:107 [15]:116 [16]:121 [17]:125 [18]:134 [19]:143 [20]:147 [21]:152 [22]:165 [23]:175 [24]:179 [25]:185 [26]:192 [27]:201 [28]:207 [29]:213 [30]:219 [31]:224 [32]:231 [33]:238 [34]:241 [35]:256 [36]:285 [37]:289 [38]:307 [39]:340 +csrVal: +[0]:1 [1]:1 [2]:1 [3]:1 [4]:1 [5]:1 [6]:1 [7]:1 [8]:1 [9]:1 [10]:1 [11]:1 [12]:1 [13]:1 [14]:1 [15]:1 [16]:1 [17]:1 [18]:1 [19]:1 [20]:1 [21]:1 [22]:1 [23]:1 [24]:1 [25]:1 [26]:1 [27]:1 [28]:1 [29]:1 [30]:1 [31]:1 [32]:1 [33]:1 [34]:1 [35]:1 [36]:1 [37]:1 [38]:1 [39]:1 +pretty print: +0 0 0 0 0 0 x x 0 0 x x x 0 0 0 0 0 0 0 +0 0 0 0 0 0 x x x 0 x x x 0 0 0 0 0 0 0 +0 0 0 0 0 0 0 0 0 0 0 0 0 x x x x x 0 0 +0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 0 0 +0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 0 0 x 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 +x x 0 0 0 x 0 x x 0 x x x 0 0 0 0 0 0 0 +x x 0 0 0 0 x 0 x x 0 0 0 0 0 0 0 0 0 x +0 x 0 0 0 0 x x 0 x 0 0 0 0 0 0 0 0 0 0 +0 0 0 0 0 0 0 x x 0 0 0 0 0 0 0 0 0 0 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 x x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 0 0 0 0 0 0 0 0 x x 0 0 x x 0 x 0 0 +0 0 0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 +cscRowInd: +[0]:6 [1]:7 [2]:10 [3]:11 [4]:12 [5]:21 [6]:22 [7]:33 [8]:34 [9]:36 [10]:38 [11]:6 [12]:7 [13]:8 [14]:10 [15]:11 [16]:12 [17]:21 [18]:22 [19]:34 [20]:35 [21]:38 [22]:13 [23]:14 [24]:15 [25]:16 [26]:17 [27]:35 [28]:38 [29]:16 [30]:35 [31]:36 [32]:38 [33]:5 [34]:33 [35]:34 [36]:38 [37]:4 [38]:6 [39]:34 +cscColPtr: +[0]:0 [1]:11 [2]:22 [3]:29 [4]:33 [5]:37 [6]:41 [7]:51 [8]:64 [9]:71 [10]:76 [11]:83 [12]:92 [13]:99 [14]:107 [15]:116 [16]:121 [17]:125 [18]:134 [19]:143 [20]:147 [21]:152 [22]:165 [23]:175 [24]:179 [25]:185 [26]:192 [27]:201 [28]:207 [29]:213 [30]:219 [31]:224 [32]:231 [33]:238 [34]:241 [35]:256 [36]:285 [37]:289 [38]:307 [39]:340 +cscVal: +[0]:1 [1]:1 [2]:1 [3]:1 [4]:1 [5]:1 [6]:1 [7]:1 [8]:1 [9]:1 [10]:1 [11]:1 [12]:1 [13]:1 [14]:1 [15]:1 [16]:1 [17]:1 [18]:1 [19]:1 [20]:1 [21]:1 [22]:1 [23]:1 [24]:1 [25]:1 [26]:1 [27]:1 [28]:1 [29]:1 [30]:1 [31]:1 [32]:1 [33]:1 [34]:1 [35]:1 [36]:1 [37]:1 [38]:1 [39]:1 +pretty print: +0 0 0 0 0 0 x x 0 0 x x x 0 0 0 0 0 0 0 +0 0 0 0 0 0 x x x 0 x x x 0 0 0 0 0 0 0 +0 0 0 0 0 0 0 0 0 0 0 0 0 x x x x x 0 0 +0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 0 0 +0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 0 0 x 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 +x x 0 0 0 x 0 x x 0 x x x 0 0 0 0 0 0 0 +x x 0 0 0 0 x 0 x x 0 0 0 0 0 0 0 0 0 x +0 x 0 0 0 0 x x 0 x 0 0 0 0 0 0 0 0 0 0 +0 0 0 0 0 0 0 x x 0 0 0 0 0 0 0 0 0 0 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 x x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 0 0 0 0 0 0 0 0 x x 0 0 x x 0 x 0 0 +0 0 0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 +loading B +Undirected due to mtx: 1 +Undirected due to cmd: 0 +Undirected: 1 +Remove self-loop: 1 +Reading ../data/small/.chesapeake.mtx.ud.nosl.bin +Allocate 40 +Allocate 408 +Allocate 408 +csrColInd: +[0]:6 [1]:7 [2]:10 [3]:11 [4]:12 [5]:21 [6]:22 [7]:33 [8]:34 [9]:36 [10]:38 [11]:6 [12]:7 [13]:8 [14]:10 [15]:11 [16]:12 [17]:21 [18]:22 [19]:34 [20]:35 [21]:38 [22]:13 [23]:14 [24]:15 [25]:16 [26]:17 [27]:35 [28]:38 [29]:16 [30]:35 [31]:36 [32]:38 [33]:5 [34]:33 [35]:34 [36]:38 [37]:4 [38]:6 [39]:34 +csrRowPtr: +[0]:0 [1]:11 [2]:22 [3]:29 [4]:33 [5]:37 [6]:41 [7]:51 [8]:64 [9]:71 [10]:76 [11]:83 [12]:92 [13]:99 [14]:107 [15]:116 [16]:121 [17]:125 [18]:134 [19]:143 [20]:147 [21]:152 [22]:165 [23]:175 [24]:179 [25]:185 [26]:192 [27]:201 [28]:207 [29]:213 [30]:219 [31]:224 [32]:231 [33]:238 [34]:241 [35]:256 [36]:285 [37]:289 [38]:307 [39]:340 +csrVal: +[0]:1 [1]:1 [2]:1 [3]:1 [4]:1 [5]:1 [6]:1 [7]:1 [8]:1 [9]:1 [10]:1 [11]:1 [12]:1 [13]:1 [14]:1 [15]:1 [16]:1 [17]:1 [18]:1 [19]:1 [20]:1 [21]:1 [22]:1 [23]:1 [24]:1 [25]:1 [26]:1 [27]:1 [28]:1 [29]:1 [30]:1 [31]:1 [32]:1 [33]:1 [34]:1 [35]:1 [36]:1 [37]:1 [38]:1 [39]:1 +pretty print: +0 0 0 0 0 0 x x 0 0 x x x 0 0 0 0 0 0 0 +0 0 0 0 0 0 x x x 0 x x x 0 0 0 0 0 0 0 +0 0 0 0 0 0 0 0 0 0 0 0 0 x x x x x 0 0 +0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 0 0 +0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 0 0 x 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 +x x 0 0 0 x 0 x x 0 x x x 0 0 0 0 0 0 0 +x x 0 0 0 0 x 0 x x 0 0 0 0 0 0 0 0 0 x +0 x 0 0 0 0 x x 0 x 0 0 0 0 0 0 0 0 0 0 +0 0 0 0 0 0 0 x x 0 0 0 0 0 0 0 0 0 0 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 x x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 0 0 0 0 0 0 0 0 x x 0 0 x x 0 x 0 0 +0 0 0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 +cscRowInd: +[0]:6 [1]:7 [2]:10 [3]:11 [4]:12 [5]:21 [6]:22 [7]:33 [8]:34 [9]:36 [10]:38 [11]:6 [12]:7 [13]:8 [14]:10 [15]:11 [16]:12 [17]:21 [18]:22 [19]:34 [20]:35 [21]:38 [22]:13 [23]:14 [24]:15 [25]:16 [26]:17 [27]:35 [28]:38 [29]:16 [30]:35 [31]:36 [32]:38 [33]:5 [34]:33 [35]:34 [36]:38 [37]:4 [38]:6 [39]:34 +cscColPtr: +[0]:0 [1]:11 [2]:22 [3]:29 [4]:33 [5]:37 [6]:41 [7]:51 [8]:64 [9]:71 [10]:76 [11]:83 [12]:92 [13]:99 [14]:107 [15]:116 [16]:121 [17]:125 [18]:134 [19]:143 [20]:147 [21]:152 [22]:165 [23]:175 [24]:179 [25]:185 [26]:192 [27]:201 [28]:207 [29]:213 [30]:219 [31]:224 [32]:231 [33]:238 [34]:241 [35]:256 [36]:285 [37]:289 [38]:307 [39]:340 +cscVal: +[0]:1 [1]:1 [2]:1 [3]:1 [4]:1 [5]:1 [6]:1 [7]:1 [8]:1 [9]:1 [10]:1 [11]:1 [12]:1 [13]:1 [14]:1 [15]:1 [16]:1 [17]:1 [18]:1 [19]:1 [20]:1 [21]:1 [22]:1 [23]:1 [24]:1 [25]:1 [26]:1 [27]:1 [28]:1 [29]:1 [30]:1 [31]:1 [32]:1 [33]:1 [34]:1 [35]:1 [36]:1 [37]:1 [38]:1 [39]:1 +pretty print: +0 0 0 0 0 0 x x 0 0 x x x 0 0 0 0 0 0 0 +0 0 0 0 0 0 x x x 0 x x x 0 0 0 0 0 0 0 +0 0 0 0 0 0 0 0 0 0 0 0 0 x x x x x 0 0 +0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 0 0 +0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 0 0 x 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 +x x 0 0 0 x 0 x x 0 x x x 0 0 0 0 0 0 0 +x x 0 0 0 0 x 0 x x 0 0 0 0 0 0 0 0 0 x +0 x 0 0 0 0 x x 0 x 0 0 0 0 0 0 0 0 0 0 +0 0 0 0 0 0 0 x x 0 0 0 0 0 0 0 0 0 0 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 x x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 0 0 0 0 0 0 0 0 x x 0 0 x x 0 x 0 0 +0 0 0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 +===Begin mxm=== +csrColInd: +[0]:6 [1]:7 [2]:10 [3]:11 [4]:12 [5]:21 [6]:22 [7]:33 [8]:34 [9]:36 [10]:38 [11]:6 [12]:7 [13]:8 [14]:10 [15]:11 [16]:12 [17]:21 [18]:22 [19]:34 [20]:35 [21]:38 [22]:13 [23]:14 [24]:15 [25]:16 [26]:17 [27]:35 [28]:38 [29]:16 [30]:35 [31]:36 [32]:38 [33]:5 [34]:33 [35]:34 [36]:38 [37]:4 [38]:6 [39]:34 +csrRowPtr: +[0]:0 [1]:11 [2]:22 [3]:29 [4]:33 [5]:37 [6]:41 [7]:51 [8]:64 [9]:71 [10]:76 [11]:83 [12]:92 [13]:99 [14]:107 [15]:116 [16]:121 [17]:125 [18]:134 [19]:143 [20]:147 [21]:152 [22]:165 [23]:175 [24]:179 [25]:185 [26]:192 [27]:201 [28]:207 [29]:213 [30]:219 [31]:224 [32]:231 [33]:238 [34]:241 [35]:256 [36]:285 [37]:289 [38]:307 [39]:340 +csrVal: +[0]:1 [1]:1 [2]:1 [3]:1 [4]:1 [5]:1 [6]:1 [7]:1 [8]:1 [9]:1 [10]:1 [11]:1 [12]:1 [13]:1 [14]:1 [15]:1 [16]:1 [17]:1 [18]:1 [19]:1 [20]:1 [21]:1 [22]:1 [23]:1 [24]:1 [25]:1 [26]:1 [27]:1 [28]:1 [29]:1 [30]:1 [31]:1 [32]:1 [33]:1 [34]:1 [35]:1 [36]:1 [37]:1 [38]:1 [39]:1 +pretty print: +0 0 0 0 0 0 x x 0 0 x x x 0 0 0 0 0 0 0 +0 0 0 0 0 0 x x x 0 x x x 0 0 0 0 0 0 0 +0 0 0 0 0 0 0 0 0 0 0 0 0 x x x x x 0 0 +0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 0 0 +0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 0 0 x 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 +x x 0 0 0 x 0 x x 0 x x x 0 0 0 0 0 0 0 +x x 0 0 0 0 x 0 x x 0 0 0 0 0 0 0 0 0 x +0 x 0 0 0 0 x x 0 x 0 0 0 0 0 0 0 0 0 0 +0 0 0 0 0 0 0 x x 0 0 0 0 0 0 0 0 0 0 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 x x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 0 0 0 0 0 0 0 0 x x 0 0 x x 0 x 0 0 +0 0 0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 +cscRowInd: +[0]:6 [1]:7 [2]:10 [3]:11 [4]:12 [5]:21 [6]:22 [7]:33 [8]:34 [9]:36 [10]:38 [11]:6 [12]:7 [13]:8 [14]:10 [15]:11 [16]:12 [17]:21 [18]:22 [19]:34 [20]:35 [21]:38 [22]:13 [23]:14 [24]:15 [25]:16 [26]:17 [27]:35 [28]:38 [29]:16 [30]:35 [31]:36 [32]:38 [33]:5 [34]:33 [35]:34 [36]:38 [37]:4 [38]:6 [39]:34 +cscColPtr: +[0]:0 [1]:11 [2]:22 [3]:29 [4]:33 [5]:37 [6]:41 [7]:51 [8]:64 [9]:71 [10]:76 [11]:83 [12]:92 [13]:99 [14]:107 [15]:116 [16]:121 [17]:125 [18]:134 [19]:143 [20]:147 [21]:152 [22]:165 [23]:175 [24]:179 [25]:185 [26]:192 [27]:201 [28]:207 [29]:213 [30]:219 [31]:224 [32]:231 [33]:238 [34]:241 [35]:256 [36]:285 [37]:289 [38]:307 [39]:340 +cscVal: +[0]:1 [1]:1 [2]:1 [3]:1 [4]:1 [5]:1 [6]:1 [7]:1 [8]:1 [9]:1 [10]:1 [11]:1 [12]:1 [13]:1 [14]:1 [15]:1 [16]:1 [17]:1 [18]:1 [19]:1 [20]:1 [21]:1 [22]:1 [23]:1 [24]:1 [25]:1 [26]:1 [27]:1 [28]:1 [29]:1 [30]:1 [31]:1 [32]:1 [33]:1 [34]:1 [35]:1 [36]:1 [37]:1 [38]:1 [39]:1 +pretty print: +0 0 0 0 0 0 x x 0 0 x x x 0 0 0 0 0 0 0 +0 0 0 0 0 0 x x x 0 x x x 0 0 0 0 0 0 0 +0 0 0 0 0 0 0 0 0 0 0 0 0 x x x x x 0 0 +0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 0 0 +0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 0 0 x 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 +x x 0 0 0 x 0 x x 0 x x x 0 0 0 0 0 0 0 +x x 0 0 0 0 x 0 x x 0 0 0 0 0 0 0 0 0 x +0 x 0 0 0 0 x x 0 x 0 0 0 0 0 0 0 0 0 0 +0 0 0 0 0 0 0 x x 0 0 0 0 0 0 0 0 0 0 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 x x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 0 0 0 0 0 0 0 0 x x 0 0 x x 0 x 0 0 +0 0 0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 +csrColInd: +[0]:6 [1]:7 [2]:10 [3]:11 [4]:12 [5]:21 [6]:22 [7]:33 [8]:34 [9]:36 [10]:38 [11]:6 [12]:7 [13]:8 [14]:10 [15]:11 [16]:12 [17]:21 [18]:22 [19]:34 [20]:35 [21]:38 [22]:13 [23]:14 [24]:15 [25]:16 [26]:17 [27]:35 [28]:38 [29]:16 [30]:35 [31]:36 [32]:38 [33]:5 [34]:33 [35]:34 [36]:38 [37]:4 [38]:6 [39]:34 +csrRowPtr: +[0]:0 [1]:11 [2]:22 [3]:29 [4]:33 [5]:37 [6]:41 [7]:51 [8]:64 [9]:71 [10]:76 [11]:83 [12]:92 [13]:99 [14]:107 [15]:116 [16]:121 [17]:125 [18]:134 [19]:143 [20]:147 [21]:152 [22]:165 [23]:175 [24]:179 [25]:185 [26]:192 [27]:201 [28]:207 [29]:213 [30]:219 [31]:224 [32]:231 [33]:238 [34]:241 [35]:256 [36]:285 [37]:289 [38]:307 [39]:340 +csrVal: +[0]:1 [1]:1 [2]:1 [3]:1 [4]:1 [5]:1 [6]:1 [7]:1 [8]:1 [9]:1 [10]:1 [11]:1 [12]:1 [13]:1 [14]:1 [15]:1 [16]:1 [17]:1 [18]:1 [19]:1 [20]:1 [21]:1 [22]:1 [23]:1 [24]:1 [25]:1 [26]:1 [27]:1 [28]:1 [29]:1 [30]:1 [31]:1 [32]:1 [33]:1 [34]:1 [35]:1 [36]:1 [37]:1 [38]:1 [39]:1 +pretty print: +0 0 0 0 0 0 x x 0 0 x x x 0 0 0 0 0 0 0 +0 0 0 0 0 0 x x x 0 x x x 0 0 0 0 0 0 0 +0 0 0 0 0 0 0 0 0 0 0 0 0 x x x x x 0 0 +0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 0 0 +0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 0 0 x 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 +x x 0 0 0 x 0 x x 0 x x x 0 0 0 0 0 0 0 +x x 0 0 0 0 x 0 x x 0 0 0 0 0 0 0 0 0 x +0 x 0 0 0 0 x x 0 x 0 0 0 0 0 0 0 0 0 0 +0 0 0 0 0 0 0 x x 0 0 0 0 0 0 0 0 0 0 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 x x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 0 0 0 0 0 0 0 0 x x 0 0 x x 0 x 0 0 +0 0 0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 +cscRowInd: +[0]:6 [1]:7 [2]:10 [3]:11 [4]:12 [5]:21 [6]:22 [7]:33 [8]:34 [9]:36 [10]:38 [11]:6 [12]:7 [13]:8 [14]:10 [15]:11 [16]:12 [17]:21 [18]:22 [19]:34 [20]:35 [21]:38 [22]:13 [23]:14 [24]:15 [25]:16 [26]:17 [27]:35 [28]:38 [29]:16 [30]:35 [31]:36 [32]:38 [33]:5 [34]:33 [35]:34 [36]:38 [37]:4 [38]:6 [39]:34 +cscColPtr: +[0]:0 [1]:11 [2]:22 [3]:29 [4]:33 [5]:37 [6]:41 [7]:51 [8]:64 [9]:71 [10]:76 [11]:83 [12]:92 [13]:99 [14]:107 [15]:116 [16]:121 [17]:125 [18]:134 [19]:143 [20]:147 [21]:152 [22]:165 [23]:175 [24]:179 [25]:185 [26]:192 [27]:201 [28]:207 [29]:213 [30]:219 [31]:224 [32]:231 [33]:238 [34]:241 [35]:256 [36]:285 [37]:289 [38]:307 [39]:340 +cscVal: +[0]:1 [1]:1 [2]:1 [3]:1 [4]:1 [5]:1 [6]:1 [7]:1 [8]:1 [9]:1 [10]:1 [11]:1 [12]:1 [13]:1 [14]:1 [15]:1 [16]:1 [17]:1 [18]:1 [19]:1 [20]:1 [21]:1 [22]:1 [23]:1 [24]:1 [25]:1 [26]:1 [27]:1 [28]:1 [29]:1 [30]:1 [31]:1 [32]:1 [33]:1 [34]:1 [35]:1 [36]:1 [37]:1 [38]:1 [39]:1 +pretty print: +0 0 0 0 0 0 x x 0 0 x x x 0 0 0 0 0 0 0 +0 0 0 0 0 0 x x x 0 x x x 0 0 0 0 0 0 0 +0 0 0 0 0 0 0 0 0 0 0 0 0 x x x x x 0 0 +0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 0 0 +0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 0 0 x 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 +x x 0 0 0 x 0 x x 0 x x x 0 0 0 0 0 0 0 +x x 0 0 0 0 x 0 x x 0 0 0 0 0 0 0 0 0 x +0 x 0 0 0 0 x x 0 x 0 0 0 0 0 0 0 0 0 0 +0 0 0 0 0 0 0 x x 0 0 0 0 0 0 0 0 0 0 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 x x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 0 0 0 0 0 0 0 0 x x 0 0 x x 0 x 0 0 +0 0 0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 +Allocate 40 +Do not allocate 0 0 +Do not allocate 0 0 +===End mxm=== +csrColInd: +[0]:0 [1]:1 [2]:2 [3]:3 [4]:4 [5]:5 [6]:6 [7]:7 [8]:8 [9]:9 [10]:10 [11]:11 [12]:12 [13]:13 [14]:14 [15]:15 [16]:16 [17]:17 [18]:18 [19]:19 [20]:20 [21]:21 [22]:22 [23]:23 [24]:24 [25]:25 [26]:26 [27]:27 [28]:28 [29]:29 [30]:30 [31]:31 [32]:32 [33]:33 [34]:34 [35]:35 [36]:36 [37]:37 [38]:38 [39]:0 +csrRowPtr: +[0]:0 [1]:39 [2]:77 [3]:113 [4]:150 [5]:187 [6]:225 [7]:264 [8]:303 [9]:341 [10]:379 [11]:418 [12]:457 [13]:495 [14]:532 [15]:569 [16]:606 [17]:643 [18]:680 [19]:717 [20]:753 [21]:790 [22]:829 [23]:868 [24]:904 [25]:940 [26]:977 [27]:1014 [28]:1051 [29]:1087 [30]:1124 [31]:1161 [32]:1198 [33]:1235 [34]:1249 [35]:1288 [36]:1324 [37]:1343 [38]:1372 [39]:1411 +csrVal: +[0]:11 [1]:9 [2]:1 [3]:2 [4]:3 [5]:3 [6]:6 [7]:5 [8]:4 [9]:3 [10]:3 [11]:3 [12]:3 [13]:1 [14]:1 [15]:1 [16]:1 [17]:1 [18]:3 [19]:2 [20]:2 [21]:3 [22]:3 [23]:2 [24]:1 [25]:2 [26]:2 [27]:2 [28]:1 [29]:3 [30]:2 [31]:3 [32]:3 [33]:1 [34]:8 [35]:6 [36]:2 [37]:5 [38]:7 [39]:9 +pretty print: +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +csrColInd: +[0]:0 [1]:1 [2]:2 [3]:3 [4]:4 [5]:5 [6]:6 [7]:7 [8]:8 [9]:9 [10]:10 [11]:11 [12]:12 [13]:13 [14]:14 [15]:15 [16]:16 [17]:17 [18]:18 [19]:19 [20]:20 [21]:21 [22]:22 [23]:23 [24]:24 [25]:25 [26]:26 [27]:27 [28]:28 [29]:29 [30]:30 [31]:31 [32]:32 [33]:33 [34]:34 [35]:35 [36]:36 [37]:37 [38]:38 [39]:0 +csrRowPtr: +[0]:0 [1]:39 [2]:77 [3]:113 [4]:150 [5]:187 [6]:225 [7]:264 [8]:303 [9]:341 [10]:379 [11]:418 [12]:457 [13]:495 [14]:532 [15]:569 [16]:606 [17]:643 [18]:680 [19]:717 [20]:753 [21]:790 [22]:829 [23]:868 [24]:904 [25]:940 [26]:977 [27]:1014 [28]:1051 [29]:1087 [30]:1124 [31]:1161 [32]:1198 [33]:1235 [34]:1249 [35]:1288 [36]:1324 [37]:1343 [38]:1372 [39]:1411 +csrVal: +[0]:11 [1]:9 [2]:1 [3]:2 [4]:3 [5]:3 [6]:6 [7]:5 [8]:4 [9]:3 [10]:3 [11]:3 [12]:3 [13]:1 [14]:1 [15]:1 [16]:1 [17]:1 [18]:3 [19]:2 [20]:2 [21]:3 [22]:3 [23]:2 [24]:1 [25]:2 [26]:2 [27]:2 [28]:1 [29]:3 [30]:2 [31]:3 [32]:3 [33]:1 [34]:8 [35]:6 [36]:2 [37]:5 [38]:7 [39]:9 +pretty print: +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +Allocate 40 +Allocate 408 +Allocate 408 +Allocate 40 +Allocate 408 +Allocate 408 +===Begin mxm=== +csrColInd: +[0]:6 [1]:7 [2]:10 [3]:11 [4]:12 [5]:21 [6]:22 [7]:33 [8]:34 [9]:36 [10]:38 [11]:6 [12]:7 [13]:8 [14]:10 [15]:11 [16]:12 [17]:21 [18]:22 [19]:34 [20]:35 [21]:38 [22]:13 [23]:14 [24]:15 [25]:16 [26]:17 [27]:35 [28]:38 [29]:16 [30]:35 [31]:36 [32]:38 [33]:5 [34]:33 [35]:34 [36]:38 [37]:4 [38]:6 [39]:34 +csrRowPtr: +[0]:0 [1]:11 [2]:22 [3]:29 [4]:33 [5]:37 [6]:41 [7]:51 [8]:64 [9]:71 [10]:76 [11]:83 [12]:92 [13]:99 [14]:107 [15]:116 [16]:121 [17]:125 [18]:134 [19]:143 [20]:147 [21]:152 [22]:165 [23]:175 [24]:179 [25]:185 [26]:192 [27]:201 [28]:207 [29]:213 [30]:219 [31]:224 [32]:231 [33]:238 [34]:241 [35]:256 [36]:285 [37]:289 [38]:307 [39]:340 +csrVal: +[0]:1 [1]:1 [2]:1 [3]:1 [4]:1 [5]:1 [6]:1 [7]:1 [8]:1 [9]:1 [10]:1 [11]:1 [12]:1 [13]:1 [14]:1 [15]:1 [16]:1 [17]:1 [18]:1 [19]:1 [20]:1 [21]:1 [22]:1 [23]:1 [24]:1 [25]:1 [26]:1 [27]:1 [28]:1 [29]:1 [30]:1 [31]:1 [32]:1 [33]:1 [34]:1 [35]:1 [36]:1 [37]:1 [38]:1 [39]:1 +pretty print: +0 0 0 0 0 0 x x 0 0 x x x 0 0 0 0 0 0 0 +0 0 0 0 0 0 x x x 0 x x x 0 0 0 0 0 0 0 +0 0 0 0 0 0 0 0 0 0 0 0 0 x x x x x 0 0 +0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 0 0 +0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 0 0 x 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 +x x 0 0 0 x 0 x x 0 x x x 0 0 0 0 0 0 0 +x x 0 0 0 0 x 0 x x 0 0 0 0 0 0 0 0 0 x +0 x 0 0 0 0 x x 0 x 0 0 0 0 0 0 0 0 0 0 +0 0 0 0 0 0 0 x x 0 0 0 0 0 0 0 0 0 0 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 x x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 0 0 0 0 0 0 0 0 x x 0 0 x x 0 x 0 0 +0 0 0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 +csrColInd: +[0]:6 [1]:7 [2]:10 [3]:11 [4]:12 [5]:21 [6]:22 [7]:33 [8]:34 [9]:36 [10]:38 [11]:6 [12]:7 [13]:8 [14]:10 [15]:11 [16]:12 [17]:21 [18]:22 [19]:34 [20]:35 [21]:38 [22]:13 [23]:14 [24]:15 [25]:16 [26]:17 [27]:35 [28]:38 [29]:16 [30]:35 [31]:36 [32]:38 [33]:5 [34]:33 [35]:34 [36]:38 [37]:4 [38]:6 [39]:34 +csrRowPtr: +[0]:0 [1]:11 [2]:22 [3]:29 [4]:33 [5]:37 [6]:41 [7]:51 [8]:64 [9]:71 [10]:76 [11]:83 [12]:92 [13]:99 [14]:107 [15]:116 [16]:121 [17]:125 [18]:134 [19]:143 [20]:147 [21]:152 [22]:165 [23]:175 [24]:179 [25]:185 [26]:192 [27]:201 [28]:207 [29]:213 [30]:219 [31]:224 [32]:231 [33]:238 [34]:241 [35]:256 [36]:285 [37]:289 [38]:307 [39]:340 +csrVal: +[0]:1 [1]:1 [2]:1 [3]:1 [4]:1 [5]:1 [6]:1 [7]:1 [8]:1 [9]:1 [10]:1 [11]:1 [12]:1 [13]:1 [14]:1 [15]:1 [16]:1 [17]:1 [18]:1 [19]:1 [20]:1 [21]:1 [22]:1 [23]:1 [24]:1 [25]:1 [26]:1 [27]:1 [28]:1 [29]:1 [30]:1 [31]:1 [32]:1 [33]:1 [34]:1 [35]:1 [36]:1 [37]:1 [38]:1 [39]:1 +pretty print: +0 0 0 0 0 0 x x 0 0 x x x 0 0 0 0 0 0 0 +0 0 0 0 0 0 x x x 0 x x x 0 0 0 0 0 0 0 +0 0 0 0 0 0 0 0 0 0 0 0 0 x x x x x 0 0 +0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 0 0 +0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 0 0 x 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 +x x 0 0 0 x 0 x x 0 x x x 0 0 0 0 0 0 0 +x x 0 0 0 0 x 0 x x 0 0 0 0 0 0 0 0 0 x +0 x 0 0 0 0 x x 0 x 0 0 0 0 0 0 0 0 0 0 +0 0 0 0 0 0 0 x x 0 0 0 0 0 0 0 0 0 0 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 +x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 x x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 +0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 +0 0 0 0 0 0 0 0 0 0 x x 0 0 x x 0 x 0 0 +0 0 0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 +Allocate 40 +Do not allocate 0 0 +Do not allocate 0 0 +===End mxm=== +csrColInd: +[0]:0 [1]:1 [2]:2 [3]:3 [4]:4 [5]:5 [6]:6 [7]:7 [8]:8 [9]:9 [10]:10 [11]:11 [12]:12 [13]:13 [14]:14 [15]:15 [16]:16 [17]:17 [18]:18 [19]:19 [20]:20 [21]:21 [22]:22 [23]:23 [24]:24 [25]:25 [26]:26 [27]:27 [28]:28 [29]:29 [30]:30 [31]:31 [32]:32 [33]:33 [34]:34 [35]:35 [36]:36 [37]:37 [38]:38 [39]:0 +csrRowPtr: +[0]:0 [1]:39 [2]:77 [3]:113 [4]:150 [5]:187 [6]:225 [7]:264 [8]:303 [9]:341 [10]:379 [11]:418 [12]:457 [13]:495 [14]:532 [15]:569 [16]:606 [17]:643 [18]:680 [19]:717 [20]:753 [21]:790 [22]:829 [23]:868 [24]:904 [25]:940 [26]:977 [27]:1014 [28]:1051 [29]:1087 [30]:1124 [31]:1161 [32]:1198 [33]:1235 [34]:1249 [35]:1288 [36]:1324 [37]:1343 [38]:1372 [39]:1411 +csrVal: +[0]:11 [1]:9 [2]:1 [3]:2 [4]:3 [5]:3 [6]:6 [7]:5 [8]:4 [9]:3 [10]:3 [11]:3 [12]:3 [13]:1 [14]:1 [15]:1 [16]:1 [17]:1 [18]:3 [19]:2 [20]:2 [21]:3 [22]:3 [23]:2 [24]:1 [25]:2 [26]:2 [27]:2 [28]:1 [29]:3 [30]:2 [31]:3 [32]:3 [33]:1 [34]:8 [35]:6 [36]:2 [37]:5 [38]:7 [39]:9 +pretty print: +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x +x x x x x x x x x x x x x x x x x x x x diff --git a/ext/GALATIC b/ext/GALATIC new file mode 160000 index 0000000..f7537d4 --- /dev/null +++ b/ext/GALATIC @@ -0,0 +1 @@ +Subproject commit f7537d409b8fd2449a1c2d9a5dda7cd51371e2ea diff --git a/ext/cub b/ext/cub deleted file mode 160000 index d622848..0000000 --- a/ext/cub +++ /dev/null @@ -1 +0,0 @@ -Subproject commit d622848f9fb62f13e5e064e1deb43b6bcbb12bad diff --git a/graphblas/backend/cuda/operations.hpp b/graphblas/backend/cuda/operations.hpp index fdcb1ba..95b40bd 100644 --- a/graphblas/backend/cuda/operations.hpp +++ b/graphblas/backend/cuda/operations.hpp @@ -43,8 +43,13 @@ Info mxm(Matrix* C, desc)); } else if (typeid(c) == typeid(float) && typeid(a) == typeid(float) && typeid(b) == typeid(float)) { + // GALATIC_spgemm(&C->sparse_, + // op, + // &A->sparse_, + // &B->sparse_, + // desc); CHECK(cusparse_spgemm2(&C->sparse_, mask, accum, op, &A->sparse_, - &B->sparse_, desc)); + &B->sparse_, desc)); } else { std::cout << "Error: Unmasked SpGEMM not implemented yet!\n"; return GrB_NOT_IMPLEMENTED; diff --git a/graphblas/backend/cuda/reduce.hpp b/graphblas/backend/cuda/reduce.hpp index 3adbca4..5c90586 100644 --- a/graphblas/backend/cuda/reduce.hpp +++ b/graphblas/backend/cuda/reduce.hpp @@ -1,7 +1,7 @@ #ifndef GRAPHBLAS_BACKEND_CUDA_REDUCE_HPP_ #define GRAPHBLAS_BACKEND_CUDA_REDUCE_HPP_ -#include +#include #include diff --git a/graphblas/backend/cuda/spgemm.hpp b/graphblas/backend/cuda/spgemm.hpp index 20aae02..9df7ad9 100644 --- a/graphblas/backend/cuda/spgemm.hpp +++ b/graphblas/backend/cuda/spgemm.hpp @@ -1,14 +1,24 @@ #ifndef GRAPHBLAS_BACKEND_CUDA_SPGEMM_HPP_ #define GRAPHBLAS_BACKEND_CUDA_SPGEMM_HPP_ + +#include "../../../ext/GALATIC/include/dCSR.cuh" +#include "../../../ext/GALATIC/include/SemiRingInterface.h" +#include "../../../ext/GALATIC/source/device/Multiply.cuh" + #include "graphblas/backend/cuda/sparse_matrix.hpp" + #include #include #include #include + + + + namespace graphblas { namespace backend { template @@ -17,6 +27,37 @@ class SparseMatrix; template class DenseMatrix; +namespace { +template +void matrixToGalatic(const SparseMatrix *input , dCSR& output) { + output.col_ids = reinterpret_cast(input->d_csrColInd_); + output.data = input->d_csrVal_; + output.row_offsets = reinterpret_cast(input->d_csrRowPtr_); + output.rows = input->nrows_; + output.cols = input->ncols_; + output.nnz = input->nvals_; +} +// for copying results back after matrix multiplicaiton +template +void galaticToSparse(SparseMatrix *output , const dCSR& input) { + output->d_csrColInd_ = reinterpret_cast(input.col_ids); + output->d_csrVal_ = input.data; + output->d_csrRowPtr_ = reinterpret_cast(input.row_offsets); + output->nvals_ = input.nnz; + output->ncapacity_ = input.nnz; +} +// to prevent double freeing of resources when destructor runs +// as we do a shallow copy in matrixToGalatic and galaticToSparse. +template +void nullizeGalaticMatrix(dCSR& m) { + m.data = nullptr; + m.col_ids = nullptr; + m.row_offsets = nullptr; +} +} // End anonymous Namespace + + + template Info spgemmMasked(SparseMatrix* C, @@ -109,6 +150,139 @@ Info spgemmMasked(SparseMatrix* C, return GrB_SUCCESS; } +// A shim between graphblast's and GALATIC's semiring interfaces +template + struct GalaticSemiring : SemiRing { + NativeSR nativeSemiring; + + __device__ c multiply(const a& left, const b& right) const + { return nativeSemiring.mul_op(left, right); } + __device__ c add(const c& left,const c& right) const + { return nativeSemiring.add_op(left, right); } + __device__ static c AdditiveIdentity() + { return NativeSR::identity(); } +}; +template +Info GALATIC_spgemm(SparseMatrix* C, + SemiringT op, + const SparseMatrix* A, + const SparseMatrix* B, + Descriptor* desc) { + + Index A_nrows, A_ncols, A_nvals; + Index B_nrows, B_ncols, B_nvals; + Index C_nrows, C_ncols, C_nvals; + + A_nrows = A->nrows_; + A_ncols = A->ncols_; + A_nvals = A->nvals_; + B_nrows = B->nrows_; + B_ncols = B->ncols_; + B_nvals = B->nvals_; + C_nrows = C->nrows_; + C_ncols = C->ncols_; + + // Dimension compatibility check + if ((A_ncols != B_nrows) || (C_ncols != B_ncols) || (C_nrows != A_nrows)) { + std::cout << "Dim mismatch mxm" << std::endl; + std::cout << A_ncols << " " << B_nrows << std::endl; + std::cout << C_ncols << " " << B_ncols << std::endl; + std::cout << C_nrows << " " << A_nrows << std::endl; + return GrB_DIMENSION_MISMATCH; + } + + + // cusparseSetMatType(descr, CUSPARSE_MATRIX_TYPE_GENERAL); + // cusparseSetMatIndexBase(descr, CUSPARSE_INDEX_BASE_ZERO); + // cusparseStatus_t status; + + + if (C->d_csrColInd_ != NULL) { + CUDA_CALL(cudaFree(C->d_csrColInd_)); + CUDA_CALL(cudaFree(C->d_csrVal_)); + C->d_csrColInd_ = NULL; + C->d_csrVal_ = NULL; + } + + if (C->d_csrRowPtr_ != NULL) { + CUDA_CALL(cudaFree( C->d_csrRowPtr_)); + C->d_csrRowPtr_ = NULL; + } + + if (C->h_csrColInd_ != NULL) { + free(C->h_csrColInd_); + free(C->h_csrVal_); + C->h_csrColInd_ = NULL; + C->h_csrVal_ = NULL; + } + + + if (C->h_csrRowPtr_ == NULL) + C->h_csrRowPtr_ = reinterpret_cast(malloc((A_nrows+1)* + sizeof(Index))); + + + + dCSR outMatrixGPU; + dCSR leftInputMatrixGPU; + dCSR rightInputMatrixGPU; + + matrixToGalatic(A, leftInputMatrixGPU); + matrixToGalatic(B, rightInputMatrixGPU); + + + const int Threads = 128; + const int BlocksPerMP = 1; + const int NNZPerThread = 2; + const int InputElementsPerThreads = 2; + const int RetainElementsPerThreads = 1; + const int MaxChunksToMerge = 16; + const int MaxChunksGeneralizedMerge = 256; // MAX: 865 + const int MergePathOptions = 8; + + + GPUMatrixMatrixMultiplyTraits DefaultTraits(Threads, BlocksPerMP, NNZPerThread, + InputElementsPerThreads, RetainElementsPerThreads, + MaxChunksToMerge, MaxChunksGeneralizedMerge, MergePathOptions ); + + const bool Debug_Mode = false; + + + // GALATIC has its own semiring interface; + // GalaticSemiring is a shim here for conversion of SemiringT, see above + + GalaticSemiring sr; + ExecutionStats stats; + + sr.nativeSemiring = op; + + ACSpGEMM::Multiply>(leftInputMatrixGPU, rightInputMatrixGPU, + outMatrixGPU, DefaultTraits, stats, + Debug_Mode, sr); + + galaticToSparse(C , outMatrixGPU); + + // prevent allocations being freed twice when destructors are ran, + // as we are doing shallow copies: + // + // A, B -> leftInputMatrixGPU, rightInputMatrixGPU + // outMatrixGPU -> C. + nullizeGalaticMatrix(outMatrixGPU); + nullizeGalaticMatrix(leftInputMatrixGPU); + nullizeGalaticMatrix(rightInputMatrixGPU); + + C->h_csrColInd_ = reinterpret_cast(malloc(C->ncapacity_*sizeof(Index))); + C->h_csrVal_ = reinterpret_cast(malloc(C->ncapacity_*sizeof(c))); + + + C->need_update_ = true; // Set flag that we need to copy data from GPU + C->csr_initialized_ = true; + C->csc_initialized_ = false; + return GrB_SUCCESS; +} + + + template Info cusparse_spgemm(SparseMatrix* C, diff --git a/graphblas/backend/cuda/spmspv_inner.hpp b/graphblas/backend/cuda/spmspv_inner.hpp index 38bac9e..2704a6b 100644 --- a/graphblas/backend/cuda/spmspv_inner.hpp +++ b/graphblas/backend/cuda/spmspv_inner.hpp @@ -5,7 +5,7 @@ #include #include -#include +#include #include #include diff --git a/graphblas/backend/cuda/spmv.hpp b/graphblas/backend/cuda/spmv.hpp index dd61f22..2d2077b 100644 --- a/graphblas/backend/cuda/spmv.hpp +++ b/graphblas/backend/cuda/spmv.hpp @@ -5,7 +5,7 @@ #include #include -#include +#include #include #include diff --git a/graphblas/stddef.hpp b/graphblas/stddef.hpp index 1a12f21..52185ff 100644 --- a/graphblas/stddef.hpp +++ b/graphblas/stddef.hpp @@ -73,20 +73,22 @@ struct less_equal { return lhs <= rhs; } }; - +namespace fixme { template struct first { inline GRB_HOST_DEVICE T_out operator()(T_in1 lhs, T_in2 rhs) { return lhs; } }; - +} +namespace fixme { template struct second { inline GRB_HOST_DEVICE T_out operator()(T_in1 lhs, T_in2 rhs) { return rhs; } }; +} template struct minimum { @@ -180,13 +182,13 @@ struct SR_NAME \ typedef T_out result_type; \ typedef T_out T_out_type; \ \ - inline T_out identity() const \ + static inline __host__ __device__ T_out identity() \ { return ADD_MONOID().identity(); } \ \ - inline __host__ __device__ T_out add_op(T_out lhs, T_out rhs) \ + inline __host__ __device__ T_out add_op(const T_out& lhs, const T_out& rhs) const \ { return ADD_MONOID()(lhs, rhs); } \ \ - inline __host__ __device__ T_out mul_op(T_in1 lhs, T_in2 rhs) \ + inline __host__ __device__ T_out mul_op(const T_in1& lhs, const T_in2& rhs) const \ { return MULT_BINARYOP()(lhs, rhs); } \ }; From 179666b3f706032d68f53adee829708bffdc7356 Mon Sep 17 00:00:00 2001 From: Richard Lettich Date: Tue, 22 Jun 2021 10:20:53 -0700 Subject: [PATCH 02/12] formatting --- graphblas/backend/cuda/spgemm.hpp | 139 ++++++++++++++++-------------- 1 file changed, 74 insertions(+), 65 deletions(-) diff --git a/graphblas/backend/cuda/spgemm.hpp b/graphblas/backend/cuda/spgemm.hpp index 9df7ad9..a4d58eb 100644 --- a/graphblas/backend/cuda/spgemm.hpp +++ b/graphblas/backend/cuda/spgemm.hpp @@ -27,37 +27,6 @@ class SparseMatrix; template class DenseMatrix; -namespace { -template -void matrixToGalatic(const SparseMatrix *input , dCSR& output) { - output.col_ids = reinterpret_cast(input->d_csrColInd_); - output.data = input->d_csrVal_; - output.row_offsets = reinterpret_cast(input->d_csrRowPtr_); - output.rows = input->nrows_; - output.cols = input->ncols_; - output.nnz = input->nvals_; -} -// for copying results back after matrix multiplicaiton -template -void galaticToSparse(SparseMatrix *output , const dCSR& input) { - output->d_csrColInd_ = reinterpret_cast(input.col_ids); - output->d_csrVal_ = input.data; - output->d_csrRowPtr_ = reinterpret_cast(input.row_offsets); - output->nvals_ = input.nnz; - output->ncapacity_ = input.nnz; -} -// to prevent double freeing of resources when destructor runs -// as we do a shallow copy in matrixToGalatic and galaticToSparse. -template -void nullizeGalaticMatrix(dCSR& m) { - m.data = nullptr; - m.col_ids = nullptr; - m.row_offsets = nullptr; -} -} // End anonymous Namespace - - - template Info spgemmMasked(SparseMatrix* C, @@ -149,21 +118,52 @@ Info spgemmMasked(SparseMatrix* C, C->csc_initialized_ = false; return GrB_SUCCESS; } +// Shallow copy graphblast sparsematrix -> Galatic dCSR format +template +static void matrixToGalatic(const SparseMatrix *input , dCSR& output) { + output.col_ids = reinterpret_cast(input->d_csrColInd_); + output.data = input->d_csrVal_; + output.row_offsets = reinterpret_cast(input->d_csrRowPtr_); + output.rows = input->nrows_; + output.cols = input->ncols_; + output.nnz = input->nvals_; +} -// A shim between graphblast's and GALATIC's semiring interfaces +// Shallow copy Galatic dCSR format -> graphblast sparsematrix +template +static void galaticToSparse(SparseMatrix *output , const dCSR& input) { + output->d_csrColInd_ = reinterpret_cast(input.col_ids); + output->d_csrVal_ = input.data; + output->d_csrRowPtr_ = reinterpret_cast(input.row_offsets); + output->nvals_ = input.nnz; + output->ncapacity_ = input.nnz; +} + +// Nullize pointers in Galatic's sparse matrices; +// Galatic's destructors check for null. Doing this will prevent double +// freeing when shallowcopying with matrixToGalatic & galaticToSparse +template +static void nullizeGalaticMatrix(dCSR& m) { + m.data = nullptr; + m.col_ids = nullptr; + m.row_offsets = nullptr; +} + +// A generic shim between graphblast's and GALATIC's semiring interfaces template - struct GalaticSemiring : SemiRing { - NativeSR nativeSemiring; - - __device__ c multiply(const a& left, const b& right) const - { return nativeSemiring.mul_op(left, right); } - __device__ c add(const c& left,const c& right) const - { return nativeSemiring.add_op(left, right); } - __device__ static c AdditiveIdentity() - { return NativeSR::identity(); } +static struct GalaticSemiring : SemiRing { + NativeSR nativeSemiring; + + __device__ c multiply(const a& left, const b& right) const + { return nativeSemiring.mul_op(left, right); } + __device__ c add(const c& left,const c& right) const + { return nativeSemiring.add_op(left, right); } + __device__ static c AdditiveIdentity() + { return NativeSR::identity(); } }; + template -Info GALATIC_spgemm(SparseMatrix* C, +Info GALATIC_spgemm(SparseMatrix* C, SemiringT op, const SparseMatrix* A, const SparseMatrix* B, @@ -191,10 +191,9 @@ Info GALATIC_spgemm(SparseMatrix* C, return GrB_DIMENSION_MISMATCH; } - - // cusparseSetMatType(descr, CUSPARSE_MATRIX_TYPE_GENERAL); - // cusparseSetMatIndexBase(descr, CUSPARSE_INDEX_BASE_ZERO); - // cusparseStatus_t status; + //fixme, not sure if this is nessecary or sufficent + cusparseSetMatType(descr, CUSPARSE_MATRIX_TYPE_GENERAL); + cusparseSetMatIndexBase(descr, CUSPARSE_INDEX_BASE_ZERO); if (C->d_csrColInd_ != NULL) { @@ -205,7 +204,7 @@ Info GALATIC_spgemm(SparseMatrix* C, } if (C->d_csrRowPtr_ != NULL) { - CUDA_CALL(cudaFree( C->d_csrRowPtr_)); + CUDA_CALL(cudaFree(C->d_csrRowPtr_)); C->d_csrRowPtr_ = NULL; } @@ -216,17 +215,11 @@ Info GALATIC_spgemm(SparseMatrix* C, C->h_csrVal_ = NULL; } - - if (C->h_csrRowPtr_ == NULL) - C->h_csrRowPtr_ = reinterpret_cast(malloc((A_nrows+1)* - sizeof(Index))); - - - dCSR outMatrixGPU; dCSR leftInputMatrixGPU; dCSR rightInputMatrixGPU; + //shallow copy input matrices to galatic format matrixToGalatic(A, leftInputMatrixGPU); matrixToGalatic(B, rightInputMatrixGPU); @@ -241,25 +234,37 @@ Info GALATIC_spgemm(SparseMatrix* C, const int MergePathOptions = 8; - GPUMatrixMatrixMultiplyTraits DefaultTraits(Threads, BlocksPerMP, NNZPerThread, - InputElementsPerThreads, RetainElementsPerThreads, - MaxChunksToMerge, MaxChunksGeneralizedMerge, MergePathOptions ); + GPUMatrixMatrixMultiplyTraits DefaultTraits( + Threads, BlocksPerMP, NNZPerThread, InputElementsPerThreads, + RetainElementsPerThreads, MaxChunksToMerge,MaxChunksGeneralizedMerge, + MergePathOptions + ); const bool Debug_Mode = false; // GALATIC has its own semiring interface; - // GalaticSemiring is a shim here for conversion of SemiringT, see above - - GalaticSemiring sr; - ExecutionStats stats; - + // GalaticSemiring is a shim here for conversion of graphblast-style + // SemiringT type. GalaticSemiring definition is above this function + GalaticSemiring semiring_shim; sr.nativeSemiring = op; - ACSpGEMM::Multiply>(leftInputMatrixGPU, rightInputMatrixGPU, - outMatrixGPU, DefaultTraits, stats, - Debug_Mode, sr); + ExecutionStats stats; + try { + ACSpGEMM::Multiply>( + leftInputMatrixGPU, rightInputMatrixGPU, outMatrixGPU, + DefaultTraits, stats, Debug_Mode, semiring_shim + ); + } catch(std::exception& e) { + std::cout + << "Exception occured in GALATIC SpGEMM, called from GALATIC_spgemm\n" + << "Exception:\n" + << e + << std::endl; + return GrB_OUT_OF_MEMORY; //most likely issue, fixme + } + // shallow copy to native format. galaticToSparse(C , outMatrixGPU); // prevent allocations being freed twice when destructors are ran, @@ -271,6 +276,10 @@ Info GALATIC_spgemm(SparseMatrix* C, nullizeGalaticMatrix(leftInputMatrixGPU); nullizeGalaticMatrix(rightInputMatrixGPU); + + if (C->h_csrRowPtr_ == NULL) + C->h_csrRowPtr_ = reinterpret_cast(malloc((A_nrows+1)* + sizeof(Index))); C->h_csrColInd_ = reinterpret_cast(malloc(C->ncapacity_*sizeof(Index))); C->h_csrVal_ = reinterpret_cast(malloc(C->ncapacity_*sizeof(c))); From 5080cc317c043698597c3a20930abd953dd33ccb Mon Sep 17 00:00:00 2001 From: Richard Lettich Date: Tue, 22 Jun 2021 23:49:52 +0000 Subject: [PATCH 03/12] Better integration, select threads/spgemm --- CMakeLists.txt | 6 +-- graphblas/backend/cuda/descriptor.hpp | 8 ++++ graphblas/backend/cuda/operations.hpp | 32 +++++++++++----- graphblas/backend/cuda/spgemm.hpp | 55 +++++++++++++++++++++------ graphblas/types.hpp | 1 + test/gspgemm.cu | 15 ++++++-- 6 files changed, 91 insertions(+), 26 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index dad5c8d..b027944 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1,7 +1,7 @@ # CMakeLists.txt for unittest # required cmakeversion -cmake_minimum_required(VERSION 3.20) +cmake_minimum_required(VERSION 2.8) # packages find_package(CUDA) find_package(Boost COMPONENTS program_options REQUIRED) @@ -23,12 +23,12 @@ set( CUDA_CUSPARSE_LIBRARY "$ENV{CUDA_HOME}/lib64/libcusparse.so" ) #FILE( GLOB_RECURSE PROJ_LIBRARIES ext/cublas1.1/*.cu ) FILE( GLOB_RECURSE PROJ_HEADERS graphblas/*.hpp) # nvcc flags -set(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS} -arch=sm_50 -lineinfo -O3 -use_fast_math -Xptxas=-v --expt-relaxed-constexpr ") +set(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS} -arch=sm_35 -lineinfo -O3 -use_fast_math -Xptxas=-v --expt-relaxed-constexpr ") #set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-fpermissive;-arch=sm_35;-lineinfo;-Xptxas=-v;-dlcm=ca;-maxrregcount=64) #set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-gencode arch=compute_20,code=sm_21) # needed for cudamalloc set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS}") -set(CMAKE_CXX_FLAGS "-fpermissive -g -m64 -std=c++14" ) +set(CMAKE_CXX_FLAGS "-fpermissive -g -std=c++14" ) #set(CMAKE_CXX_FLAGS "-fpermissive -pg -m64 -std=c++11" ) #set(CMAKE_CXX_FLAGS "-fpermissive -g -m64 -std=c++11 -H" ) set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${OpenMP_CXX_FLAGS}") diff --git a/graphblas/backend/cuda/descriptor.hpp b/graphblas/backend/cuda/descriptor.hpp index 9788143..0f1ae7d 100644 --- a/graphblas/backend/cuda/descriptor.hpp +++ b/graphblas/backend/cuda/descriptor.hpp @@ -280,6 +280,14 @@ Info Descriptor::loadArgs(const po::variables_map& vm) { std::cout << "Error: incorrect nthread selection!\n"; } + if(mode_ == "galatic") { + CHECK(set(GrB_MODE, GrB_GALATIC)); + } else if (mode_ == "cusparse2") { + CHECK(set(GrB_MODE, GrB_CUSPARSE2)); + } else { + std::cout << R"(Invalid mode: Options are "galatic" and "cusparse2")" << std::endl; + } + // TODO(@ctcyang): Enable device selection using ndevice_ // if( ndevice_!=0 ) diff --git a/graphblas/backend/cuda/operations.hpp b/graphblas/backend/cuda/operations.hpp index 95b40bd..80dfaf7 100644 --- a/graphblas/backend/cuda/operations.hpp +++ b/graphblas/backend/cuda/operations.hpp @@ -43,16 +43,30 @@ Info mxm(Matrix* C, desc)); } else if (typeid(c) == typeid(float) && typeid(a) == typeid(float) && typeid(b) == typeid(float)) { - // GALATIC_spgemm(&C->sparse_, - // op, - // &A->sparse_, - // &B->sparse_, - // desc); - CHECK(cusparse_spgemm2(&C->sparse_, mask, accum, op, &A->sparse_, - &B->sparse_, desc)); + + Desc_value s_mode; + CHECK(desc->get(GrB_MODE, &s_mode)); + + if (s_mode == GrB_CUSPARSE2) + CHECK(cusparse_spgemm2(&C->sparse_, mask, accum, op, &A->sparse_, + &B->sparse_, desc)); + else { + if (s_mode != GrB_GALATIC) { + std::cout << R"(Unknown mode (Options are: "cusspare2" and "galatic"; defaulting to galatic)" << std::endl; + } + CHECK(GALATIC_spgemm(&C->sparse_, + op, + &A->sparse_, + &B->sparse_, + desc)); + + } } else { - std::cout << "Error: Unmasked SpGEMM not implemented yet!\n"; - return GrB_NOT_IMPLEMENTED; + CHECK(GALATIC_spgemm(&C->sparse_, + op, + &A->sparse_, + &B->sparse_, + desc)); } } else { std::cout << "Error: SpMM and GEMM not implemented yet!\n"; diff --git a/graphblas/backend/cuda/spgemm.hpp b/graphblas/backend/cuda/spgemm.hpp index a4d58eb..56071fc 100644 --- a/graphblas/backend/cuda/spgemm.hpp +++ b/graphblas/backend/cuda/spgemm.hpp @@ -151,7 +151,7 @@ static void nullizeGalaticMatrix(dCSR& m) { // A generic shim between graphblast's and GALATIC's semiring interfaces template -static struct GalaticSemiring : SemiRing { +struct GalaticSemiring : SemiRing { NativeSR nativeSemiring; __device__ c multiply(const a& left, const b& right) const @@ -192,8 +192,8 @@ Info GALATIC_spgemm(SparseMatrix* C, } //fixme, not sure if this is nessecary or sufficent - cusparseSetMatType(descr, CUSPARSE_MATRIX_TYPE_GENERAL); - cusparseSetMatIndexBase(descr, CUSPARSE_INDEX_BASE_ZERO); +// cusparseSetMatType(desc, CUSPARSE_MATRIX_TYPE_GENERAL); +// cusparseSetMatIndexBase(descr, CUSPARSE_INDEX_BASE_ZERO); if (C->d_csrColInd_ != NULL) { @@ -222,6 +222,7 @@ Info GALATIC_spgemm(SparseMatrix* C, //shallow copy input matrices to galatic format matrixToGalatic(A, leftInputMatrixGPU); matrixToGalatic(B, rightInputMatrixGPU); + const int Threads = 128; @@ -247,21 +248,53 @@ Info GALATIC_spgemm(SparseMatrix* C, // GalaticSemiring is a shim here for conversion of graphblast-style // SemiringT type. GalaticSemiring definition is above this function GalaticSemiring semiring_shim; - sr.nativeSemiring = op; + semiring_shim.nativeSemiring = op; ExecutionStats stats; try { - ACSpGEMM::Multiply>( - leftInputMatrixGPU, rightInputMatrixGPU, outMatrixGPU, - DefaultTraits, stats, Debug_Mode, semiring_shim - ); + + + Desc_value nt_mode; + CHECK(desc->get(GrB_NT, &nt_mode)); + const int num_threads = static_cast(nt_mode); + + switch (num_threads) { + case 64: + ACSpGEMM::MultiplyImplementation, + 64, 4, 2, 8, 4, 16, 512, 8, 0, a, b, c, + GalaticSemiring> + (leftInputMatrixGPU, rightInputMatrixGPU, + outMatrixGPU, DefaultTraits, stats, semiring_shim); + break; + case 128: + ACSpGEMM::MultiplyImplementation, + 128, 4, 2, 4, 4, 16, 512, 8, 0, a, b, c, + GalaticSemiring> + ( leftInputMatrixGPU, rightInputMatrixGPU, + outMatrixGPU, DefaultTraits, stats, semiring_shim); + break; + case 512: + ACSpGEMM::MultiplyImplementation, + 512, 1, 1, 1, 2, 16, 512, 8, 0, a, b, c, + GalaticSemiring> + (leftInputMatrixGPU, rightInputMatrixGPU, + outMatrixGPU, DefaultTraits, stats, semiring_shim); + break; + default: // 256 + ACSpGEMM::MultiplyImplementation, + 256, 4, 2, 4, 4, 16, 512, 8, 0, a, b, c, + GalaticSemiring> + (leftInputMatrixGPU, rightInputMatrixGPU, + outMatrixGPU, DefaultTraits, stats, semiring_shim); + break; + } } catch(std::exception& e) { - std::cout + std::cerr << "Exception occured in GALATIC SpGEMM, called from GALATIC_spgemm\n" << "Exception:\n" - << e + << e.what() << std::endl; - return GrB_OUT_OF_MEMORY; //most likely issue, fixme + return GrB_OUT_OF_MEMORY; //the most likely issue, fixme } // shallow copy to native format. diff --git a/graphblas/types.hpp b/graphblas/types.hpp index cf5c6ba..35a8892 100644 --- a/graphblas/types.hpp +++ b/graphblas/types.hpp @@ -60,6 +60,7 @@ enum Desc_value {GrB_SCMP, // for GrB_MASK GrB_DEFAULT, GrB_CUSPARSE, // for SpMV, SpMM GrB_CUSPARSE2, + GrB_GALATIC, GrB_FIXEDROW, GrB_FIXEDCOL, GrB_MERGEPATH = 9, diff --git a/test/gspgemm.cu b/test/gspgemm.cu index 9e08388..df1b0b8 100644 --- a/test/gspgemm.cu +++ b/test/gspgemm.cu @@ -23,7 +23,9 @@ int main( int argc, char** argv ) graphblas::Index b_num_rows, b_num_cols, b_num_edges; char* dat_name; - // Load A + + + // Load A std::cout << "loading A" << std::endl; readMtx("../data/small/chesapeake.mtx", &a_row_indices, &a_col_indices, &a_values, &a_num_rows, &a_num_cols, &a_num_edges, 0, false, &dat_name); @@ -44,7 +46,13 @@ int main( int argc, char** argv ) // graphblas::Matrix c(a_num_rows, b_num_cols); graphblas::Descriptor desc; - desc.descriptor_.debug_ = true; + + po::variables_map vm; + parseArgs(argc, argv, &vm); + CHECK(desc.loadArgs(vm)); + + + desc.descriptor_.debug_ = true; graphblas::mxm( &c, GrB_NULL, @@ -64,9 +72,10 @@ int main( int argc, char** argv ) A.build(a.matrix_.sparse_.d_csrRowPtr_, a.matrix_.sparse_.d_csrColInd_, a.matrix_.sparse_.d_csrVal_, a.matrix_.sparse_.nvals_); B.build(b.matrix_.sparse_.d_csrRowPtr_, b.matrix_.sparse_.d_csrColInd_, b.matrix_.sparse_.d_csrVal_, b.matrix_.sparse_.nvals_); + desc.descriptor_.debug_ = true; - graphblas::mxm(&C, GrB_NULL, GrB_NULL, graphblas::PlusMultipliesSemiring(), + graphblas::mxm(&C, GrB_NULL, GrB_NULL, graphblas::PlusDividesSemiring(), &A, &B, &desc); // Multiply using CPU array initialization. From 4cb56c768e9b394ba6693ea98d43dfa473ee1ba3 Mon Sep 17 00:00:00 2001 From: Richard Lettich Date: Tue, 22 Jun 2021 23:51:42 +0000 Subject: [PATCH 04/12] slight cleanup --- graphblas/backend/cuda/spgemm.hpp | 29 +---------------------------- 1 file changed, 1 insertion(+), 28 deletions(-) diff --git a/graphblas/backend/cuda/spgemm.hpp b/graphblas/backend/cuda/spgemm.hpp index 56071fc..d1e3756 100644 --- a/graphblas/backend/cuda/spgemm.hpp +++ b/graphblas/backend/cuda/spgemm.hpp @@ -190,11 +190,6 @@ Info GALATIC_spgemm(SparseMatrix* C, std::cout << C_nrows << " " << A_nrows << std::endl; return GrB_DIMENSION_MISMATCH; } - - //fixme, not sure if this is nessecary or sufficent -// cusparseSetMatType(desc, CUSPARSE_MATRIX_TYPE_GENERAL); -// cusparseSetMatIndexBase(descr, CUSPARSE_INDEX_BASE_ZERO); - if (C->d_csrColInd_ != NULL) { CUDA_CALL(cudaFree(C->d_csrColInd_)); @@ -223,26 +218,6 @@ Info GALATIC_spgemm(SparseMatrix* C, matrixToGalatic(A, leftInputMatrixGPU); matrixToGalatic(B, rightInputMatrixGPU); - - - const int Threads = 128; - const int BlocksPerMP = 1; - const int NNZPerThread = 2; - const int InputElementsPerThreads = 2; - const int RetainElementsPerThreads = 1; - const int MaxChunksToMerge = 16; - const int MaxChunksGeneralizedMerge = 256; // MAX: 865 - const int MergePathOptions = 8; - - - GPUMatrixMatrixMultiplyTraits DefaultTraits( - Threads, BlocksPerMP, NNZPerThread, InputElementsPerThreads, - RetainElementsPerThreads, MaxChunksToMerge,MaxChunksGeneralizedMerge, - MergePathOptions - ); - - const bool Debug_Mode = false; - // GALATIC has its own semiring interface; // GalaticSemiring is a shim here for conversion of graphblast-style @@ -252,8 +227,6 @@ Info GALATIC_spgemm(SparseMatrix* C, ExecutionStats stats; try { - - Desc_value nt_mode; CHECK(desc->get(GrB_NT, &nt_mode)); const int num_threads = static_cast(nt_mode); @@ -273,7 +246,7 @@ Info GALATIC_spgemm(SparseMatrix* C, ( leftInputMatrixGPU, rightInputMatrixGPU, outMatrixGPU, DefaultTraits, stats, semiring_shim); break; - case 512: + case 512: // experimental ACSpGEMM::MultiplyImplementation, 512, 1, 1, 1, 2, 16, 512, 8, 0, a, b, c, GalaticSemiring> From ca71c72944213c7169effef866b5f6594fb633eb Mon Sep 17 00:00:00 2001 From: Richard Lettich Date: Wed, 23 Jun 2021 00:09:14 +0000 Subject: [PATCH 05/12] some formatting and minor adjustments --- graphblas/backend/cuda/spgemm.hpp | 8 +++++++- graphblas/stddef.hpp | 28 +++++++++++++--------------- test/gspgemm.cu | 20 ++++++++++---------- 3 files changed, 30 insertions(+), 26 deletions(-) diff --git a/graphblas/backend/cuda/spgemm.hpp b/graphblas/backend/cuda/spgemm.hpp index d1e3756..0a3b4ba 100644 --- a/graphblas/backend/cuda/spgemm.hpp +++ b/graphblas/backend/cuda/spgemm.hpp @@ -219,6 +219,10 @@ Info GALATIC_spgemm(SparseMatrix* C, matrixToGalatic(B, rightInputMatrixGPU); + + GPUMatrixMatrixMultiplyTraits DefaultTraits; + + // GALATIC has its own semiring interface; // GalaticSemiring is a shim here for conversion of graphblast-style // SemiringT type. GalaticSemiring definition is above this function @@ -227,6 +231,8 @@ Info GALATIC_spgemm(SparseMatrix* C, ExecutionStats stats; try { + + Desc_value nt_mode; CHECK(desc->get(GrB_NT, &nt_mode)); const int num_threads = static_cast(nt_mode); @@ -246,7 +252,7 @@ Info GALATIC_spgemm(SparseMatrix* C, ( leftInputMatrixGPU, rightInputMatrixGPU, outMatrixGPU, DefaultTraits, stats, semiring_shim); break; - case 512: // experimental + case 512: ACSpGEMM::MultiplyImplementation, 512, 1, 1, 1, 2, 16, 512, 8, 0, a, b, c, GalaticSemiring> diff --git a/graphblas/stddef.hpp b/graphblas/stddef.hpp index 52185ff..99ac3cc 100644 --- a/graphblas/stddef.hpp +++ b/graphblas/stddef.hpp @@ -175,21 +175,19 @@ REGISTER_MONOID(NotEqualToMonoid, not_equal_to, std::numeric_limits::max( } // namespace graphblas // Semiring generator macro provided by Scott McMillan -#define REGISTER_SEMIRING(SR_NAME, ADD_MONOID, MULT_BINARYOP) \ -template \ -struct SR_NAME \ -{ \ - typedef T_out result_type; \ - typedef T_out T_out_type; \ - \ - static inline __host__ __device__ T_out identity() \ - { return ADD_MONOID().identity(); } \ - \ - inline __host__ __device__ T_out add_op(const T_out& lhs, const T_out& rhs) const \ - { return ADD_MONOID()(lhs, rhs); } \ - \ - inline __host__ __device__ T_out mul_op(const T_in1& lhs, const T_in2& rhs) const \ - { return MULT_BINARYOP()(lhs, rhs); } \ +#define REGISTER_SEMIRING(SR_NAME, ADD_MONOID, MULT_BINARYOP) \ +template \ +struct SR_NAME \ +{ \ + typedef T_out result_type; \ + typedef T_out T_out_type; \ + \ + static inline GRB_HOST_DEVICE T_out identity() \ + { return ADD_MONOID().identity(); } \ + inline GRB_HOST_DEVICE T_out add_op(const T_out& lhs, const T_out& rhs) const \ + { return ADD_MONOID()(lhs, rhs); } \ + inline GRB_HOST_DEVICE T_out mul_op(const T_in1& lhs, const T_in2& rhs) const \ + { return MULT_BINARYOP()(lhs, rhs); } \ }; namespace graphblas { diff --git a/test/gspgemm.cu b/test/gspgemm.cu index df1b0b8..c7668a4 100644 --- a/test/gspgemm.cu +++ b/test/gspgemm.cu @@ -18,7 +18,7 @@ int main( int argc, char** argv ) std::vector a_row_indices, b_row_indices; std::vector a_col_indices, b_col_indices; - std::vector a_values, b_values; + std::vector a_values, b_values; graphblas::Index a_num_rows, a_num_cols, a_num_edges; graphblas::Index b_num_rows, b_num_cols, b_num_edges; char* dat_name; @@ -29,7 +29,7 @@ int main( int argc, char** argv ) std::cout << "loading A" << std::endl; readMtx("../data/small/chesapeake.mtx", &a_row_indices, &a_col_indices, &a_values, &a_num_rows, &a_num_cols, &a_num_edges, 0, false, &dat_name); - graphblas::Matrix a(a_num_rows, a_num_cols); + graphblas::Matrix a(a_num_rows, a_num_cols); a.build(&a_row_indices, &a_col_indices, &a_values, a_num_edges, GrB_NULL, dat_name); if(DEBUG) a.print(); @@ -38,13 +38,13 @@ int main( int argc, char** argv ) std::cout << "loading B" << std::endl; readMtx("../data/small/chesapeake.mtx", &b_row_indices, &b_col_indices, &b_values, &b_num_rows, &b_num_cols, &b_num_edges, 0, false, &dat_name); - graphblas::Matrix b(b_num_rows, b_num_cols); + graphblas::Matrix b(b_num_rows, b_num_cols); b.build(&b_row_indices, &b_col_indices, &b_values, b_num_edges, GrB_NULL, dat_name); if(DEBUG) b.print(); // - graphblas::Matrix c(a_num_rows, b_num_cols); + graphblas::Matrix c(a_num_rows, b_num_cols); graphblas::Descriptor desc; po::variables_map vm; @@ -53,11 +53,11 @@ int main( int argc, char** argv ) desc.descriptor_.debug_ = true; - graphblas::mxm( + graphblas::mxm( &c, GrB_NULL, GrB_NULL, - graphblas::PlusMultipliesSemiring(), + graphblas::PlusMultipliesSemiring(), &a, &b, &desc @@ -65,9 +65,9 @@ int main( int argc, char** argv ) if(DEBUG) c.print(); // Multiply using GPU array initialization. - graphblas::Matrix A(a_num_rows, a_num_cols); - graphblas::Matrix B(b_num_rows, b_num_cols); - graphblas::Matrix C(a_num_rows, b_num_cols); + graphblas::Matrix A(a_num_rows, a_num_cols); + graphblas::Matrix B(b_num_rows, b_num_cols); + graphblas::Matrix C(a_num_rows, b_num_cols); A.build(a.matrix_.sparse_.d_csrRowPtr_, a.matrix_.sparse_.d_csrColInd_, a.matrix_.sparse_.d_csrVal_, a.matrix_.sparse_.nvals_); B.build(b.matrix_.sparse_.d_csrRowPtr_, b.matrix_.sparse_.d_csrColInd_, b.matrix_.sparse_.d_csrVal_, b.matrix_.sparse_.nvals_); @@ -75,7 +75,7 @@ int main( int argc, char** argv ) desc.descriptor_.debug_ = true; - graphblas::mxm(&C, GrB_NULL, GrB_NULL, graphblas::PlusDividesSemiring(), + graphblas::mxm(&C, GrB_NULL, GrB_NULL, graphblas::CustomLessPlusSemiring(), &A, &B, &desc); // Multiply using CPU array initialization. From b9e075cc3104cf18a191cb9c63f78624815b4b54 Mon Sep 17 00:00:00 2001 From: Richard Lettich Date: Wed, 23 Jun 2021 00:28:33 +0000 Subject: [PATCH 06/12] edit readme --- README.md | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/README.md b/README.md index 46c81f1..f5eb211 100644 --- a/README.md +++ b/README.md @@ -9,11 +9,13 @@ GraphBLAST is a GPU implementation of [GraphBLAS](http://www.graphblas.org), an ## Prerequisites -This software has been tested on the following dependencies: +This software has been tested to build with the following dependencies: -* CUDA 9.1, 9.2 -* Boost 1.58 -* g++ 4.9.3, 5.4.0 +* CUDA 11.3 + * (Change: CUDA > 11 is now required) +* Boost 1.74 +* g++ 8.3.0 + * (Change: C++14 is required) Optional: From 2c7bdb99c4ff52c3d19decc89ab63781948e1370 Mon Sep 17 00:00:00 2001 From: Richard Lettich Date: Wed, 23 Jun 2021 00:40:16 +0000 Subject: [PATCH 07/12] removed submodule cub --- .gitmodules | 3 --- 1 file changed, 3 deletions(-) diff --git a/.gitmodules b/.gitmodules index e1d59bf..b003727 100644 --- a/.gitmodules +++ b/.gitmodules @@ -1,9 +1,6 @@ [submodule "ext/moderngpu"] path = ext/moderngpu url = https://ctcyang@github.com/ctcyang/moderngpu.git -[submodule "ext/cub"] - path = ext/cub - url = https://ctcyang@github.com/NVlabs/cub.git [submodule "ext/GALATIC"] path = ext/GALATIC url = git@github.com:richardlett/GALATIC.git From 99c74ba77ca969322c7621ccefbb438e460b9089 Mon Sep 17 00:00:00 2001 From: Richard Lettich Date: Wed, 23 Jun 2021 00:42:40 +0000 Subject: [PATCH 08/12] removed build directory --- build/outtt | 385 ---------------------------------------------------- 1 file changed, 385 deletions(-) delete mode 100644 build/outtt diff --git a/build/outtt b/build/outtt deleted file mode 100644 index 17b909b..0000000 --- a/build/outtt +++ /dev/null @@ -1,385 +0,0 @@ -loading A -Undirected due to mtx: 1 -Undirected due to cmd: 0 -Undirected: 1 -Remove self-loop: 1 -Reading ../data/small/.chesapeake.mtx.ud.nosl.bin -Allocate 40 -Allocate 408 -Allocate 408 -csrColInd: -[0]:6 [1]:7 [2]:10 [3]:11 [4]:12 [5]:21 [6]:22 [7]:33 [8]:34 [9]:36 [10]:38 [11]:6 [12]:7 [13]:8 [14]:10 [15]:11 [16]:12 [17]:21 [18]:22 [19]:34 [20]:35 [21]:38 [22]:13 [23]:14 [24]:15 [25]:16 [26]:17 [27]:35 [28]:38 [29]:16 [30]:35 [31]:36 [32]:38 [33]:5 [34]:33 [35]:34 [36]:38 [37]:4 [38]:6 [39]:34 -csrRowPtr: -[0]:0 [1]:11 [2]:22 [3]:29 [4]:33 [5]:37 [6]:41 [7]:51 [8]:64 [9]:71 [10]:76 [11]:83 [12]:92 [13]:99 [14]:107 [15]:116 [16]:121 [17]:125 [18]:134 [19]:143 [20]:147 [21]:152 [22]:165 [23]:175 [24]:179 [25]:185 [26]:192 [27]:201 [28]:207 [29]:213 [30]:219 [31]:224 [32]:231 [33]:238 [34]:241 [35]:256 [36]:285 [37]:289 [38]:307 [39]:340 -csrVal: -[0]:1 [1]:1 [2]:1 [3]:1 [4]:1 [5]:1 [6]:1 [7]:1 [8]:1 [9]:1 [10]:1 [11]:1 [12]:1 [13]:1 [14]:1 [15]:1 [16]:1 [17]:1 [18]:1 [19]:1 [20]:1 [21]:1 [22]:1 [23]:1 [24]:1 [25]:1 [26]:1 [27]:1 [28]:1 [29]:1 [30]:1 [31]:1 [32]:1 [33]:1 [34]:1 [35]:1 [36]:1 [37]:1 [38]:1 [39]:1 -pretty print: -0 0 0 0 0 0 x x 0 0 x x x 0 0 0 0 0 0 0 -0 0 0 0 0 0 x x x 0 x x x 0 0 0 0 0 0 0 -0 0 0 0 0 0 0 0 0 0 0 0 0 x x x x x 0 0 -0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 0 0 -0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 0 0 x 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 -x x 0 0 0 x 0 x x 0 x x x 0 0 0 0 0 0 0 -x x 0 0 0 0 x 0 x x 0 0 0 0 0 0 0 0 0 x -0 x 0 0 0 0 x x 0 x 0 0 0 0 0 0 0 0 0 0 -0 0 0 0 0 0 0 x x 0 0 0 0 0 0 0 0 0 0 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 x x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 0 0 0 0 0 0 0 0 x x 0 0 x x 0 x 0 0 -0 0 0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 -cscRowInd: -[0]:6 [1]:7 [2]:10 [3]:11 [4]:12 [5]:21 [6]:22 [7]:33 [8]:34 [9]:36 [10]:38 [11]:6 [12]:7 [13]:8 [14]:10 [15]:11 [16]:12 [17]:21 [18]:22 [19]:34 [20]:35 [21]:38 [22]:13 [23]:14 [24]:15 [25]:16 [26]:17 [27]:35 [28]:38 [29]:16 [30]:35 [31]:36 [32]:38 [33]:5 [34]:33 [35]:34 [36]:38 [37]:4 [38]:6 [39]:34 -cscColPtr: -[0]:0 [1]:11 [2]:22 [3]:29 [4]:33 [5]:37 [6]:41 [7]:51 [8]:64 [9]:71 [10]:76 [11]:83 [12]:92 [13]:99 [14]:107 [15]:116 [16]:121 [17]:125 [18]:134 [19]:143 [20]:147 [21]:152 [22]:165 [23]:175 [24]:179 [25]:185 [26]:192 [27]:201 [28]:207 [29]:213 [30]:219 [31]:224 [32]:231 [33]:238 [34]:241 [35]:256 [36]:285 [37]:289 [38]:307 [39]:340 -cscVal: -[0]:1 [1]:1 [2]:1 [3]:1 [4]:1 [5]:1 [6]:1 [7]:1 [8]:1 [9]:1 [10]:1 [11]:1 [12]:1 [13]:1 [14]:1 [15]:1 [16]:1 [17]:1 [18]:1 [19]:1 [20]:1 [21]:1 [22]:1 [23]:1 [24]:1 [25]:1 [26]:1 [27]:1 [28]:1 [29]:1 [30]:1 [31]:1 [32]:1 [33]:1 [34]:1 [35]:1 [36]:1 [37]:1 [38]:1 [39]:1 -pretty print: -0 0 0 0 0 0 x x 0 0 x x x 0 0 0 0 0 0 0 -0 0 0 0 0 0 x x x 0 x x x 0 0 0 0 0 0 0 -0 0 0 0 0 0 0 0 0 0 0 0 0 x x x x x 0 0 -0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 0 0 -0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 0 0 x 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 -x x 0 0 0 x 0 x x 0 x x x 0 0 0 0 0 0 0 -x x 0 0 0 0 x 0 x x 0 0 0 0 0 0 0 0 0 x -0 x 0 0 0 0 x x 0 x 0 0 0 0 0 0 0 0 0 0 -0 0 0 0 0 0 0 x x 0 0 0 0 0 0 0 0 0 0 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 x x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 0 0 0 0 0 0 0 0 x x 0 0 x x 0 x 0 0 -0 0 0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 -loading B -Undirected due to mtx: 1 -Undirected due to cmd: 0 -Undirected: 1 -Remove self-loop: 1 -Reading ../data/small/.chesapeake.mtx.ud.nosl.bin -Allocate 40 -Allocate 408 -Allocate 408 -csrColInd: -[0]:6 [1]:7 [2]:10 [3]:11 [4]:12 [5]:21 [6]:22 [7]:33 [8]:34 [9]:36 [10]:38 [11]:6 [12]:7 [13]:8 [14]:10 [15]:11 [16]:12 [17]:21 [18]:22 [19]:34 [20]:35 [21]:38 [22]:13 [23]:14 [24]:15 [25]:16 [26]:17 [27]:35 [28]:38 [29]:16 [30]:35 [31]:36 [32]:38 [33]:5 [34]:33 [35]:34 [36]:38 [37]:4 [38]:6 [39]:34 -csrRowPtr: -[0]:0 [1]:11 [2]:22 [3]:29 [4]:33 [5]:37 [6]:41 [7]:51 [8]:64 [9]:71 [10]:76 [11]:83 [12]:92 [13]:99 [14]:107 [15]:116 [16]:121 [17]:125 [18]:134 [19]:143 [20]:147 [21]:152 [22]:165 [23]:175 [24]:179 [25]:185 [26]:192 [27]:201 [28]:207 [29]:213 [30]:219 [31]:224 [32]:231 [33]:238 [34]:241 [35]:256 [36]:285 [37]:289 [38]:307 [39]:340 -csrVal: -[0]:1 [1]:1 [2]:1 [3]:1 [4]:1 [5]:1 [6]:1 [7]:1 [8]:1 [9]:1 [10]:1 [11]:1 [12]:1 [13]:1 [14]:1 [15]:1 [16]:1 [17]:1 [18]:1 [19]:1 [20]:1 [21]:1 [22]:1 [23]:1 [24]:1 [25]:1 [26]:1 [27]:1 [28]:1 [29]:1 [30]:1 [31]:1 [32]:1 [33]:1 [34]:1 [35]:1 [36]:1 [37]:1 [38]:1 [39]:1 -pretty print: -0 0 0 0 0 0 x x 0 0 x x x 0 0 0 0 0 0 0 -0 0 0 0 0 0 x x x 0 x x x 0 0 0 0 0 0 0 -0 0 0 0 0 0 0 0 0 0 0 0 0 x x x x x 0 0 -0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 0 0 -0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 0 0 x 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 -x x 0 0 0 x 0 x x 0 x x x 0 0 0 0 0 0 0 -x x 0 0 0 0 x 0 x x 0 0 0 0 0 0 0 0 0 x -0 x 0 0 0 0 x x 0 x 0 0 0 0 0 0 0 0 0 0 -0 0 0 0 0 0 0 x x 0 0 0 0 0 0 0 0 0 0 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 x x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 0 0 0 0 0 0 0 0 x x 0 0 x x 0 x 0 0 -0 0 0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 -cscRowInd: -[0]:6 [1]:7 [2]:10 [3]:11 [4]:12 [5]:21 [6]:22 [7]:33 [8]:34 [9]:36 [10]:38 [11]:6 [12]:7 [13]:8 [14]:10 [15]:11 [16]:12 [17]:21 [18]:22 [19]:34 [20]:35 [21]:38 [22]:13 [23]:14 [24]:15 [25]:16 [26]:17 [27]:35 [28]:38 [29]:16 [30]:35 [31]:36 [32]:38 [33]:5 [34]:33 [35]:34 [36]:38 [37]:4 [38]:6 [39]:34 -cscColPtr: -[0]:0 [1]:11 [2]:22 [3]:29 [4]:33 [5]:37 [6]:41 [7]:51 [8]:64 [9]:71 [10]:76 [11]:83 [12]:92 [13]:99 [14]:107 [15]:116 [16]:121 [17]:125 [18]:134 [19]:143 [20]:147 [21]:152 [22]:165 [23]:175 [24]:179 [25]:185 [26]:192 [27]:201 [28]:207 [29]:213 [30]:219 [31]:224 [32]:231 [33]:238 [34]:241 [35]:256 [36]:285 [37]:289 [38]:307 [39]:340 -cscVal: -[0]:1 [1]:1 [2]:1 [3]:1 [4]:1 [5]:1 [6]:1 [7]:1 [8]:1 [9]:1 [10]:1 [11]:1 [12]:1 [13]:1 [14]:1 [15]:1 [16]:1 [17]:1 [18]:1 [19]:1 [20]:1 [21]:1 [22]:1 [23]:1 [24]:1 [25]:1 [26]:1 [27]:1 [28]:1 [29]:1 [30]:1 [31]:1 [32]:1 [33]:1 [34]:1 [35]:1 [36]:1 [37]:1 [38]:1 [39]:1 -pretty print: -0 0 0 0 0 0 x x 0 0 x x x 0 0 0 0 0 0 0 -0 0 0 0 0 0 x x x 0 x x x 0 0 0 0 0 0 0 -0 0 0 0 0 0 0 0 0 0 0 0 0 x x x x x 0 0 -0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 0 0 -0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 0 0 x 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 -x x 0 0 0 x 0 x x 0 x x x 0 0 0 0 0 0 0 -x x 0 0 0 0 x 0 x x 0 0 0 0 0 0 0 0 0 x -0 x 0 0 0 0 x x 0 x 0 0 0 0 0 0 0 0 0 0 -0 0 0 0 0 0 0 x x 0 0 0 0 0 0 0 0 0 0 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 x x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 0 0 0 0 0 0 0 0 x x 0 0 x x 0 x 0 0 -0 0 0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 -===Begin mxm=== -csrColInd: -[0]:6 [1]:7 [2]:10 [3]:11 [4]:12 [5]:21 [6]:22 [7]:33 [8]:34 [9]:36 [10]:38 [11]:6 [12]:7 [13]:8 [14]:10 [15]:11 [16]:12 [17]:21 [18]:22 [19]:34 [20]:35 [21]:38 [22]:13 [23]:14 [24]:15 [25]:16 [26]:17 [27]:35 [28]:38 [29]:16 [30]:35 [31]:36 [32]:38 [33]:5 [34]:33 [35]:34 [36]:38 [37]:4 [38]:6 [39]:34 -csrRowPtr: -[0]:0 [1]:11 [2]:22 [3]:29 [4]:33 [5]:37 [6]:41 [7]:51 [8]:64 [9]:71 [10]:76 [11]:83 [12]:92 [13]:99 [14]:107 [15]:116 [16]:121 [17]:125 [18]:134 [19]:143 [20]:147 [21]:152 [22]:165 [23]:175 [24]:179 [25]:185 [26]:192 [27]:201 [28]:207 [29]:213 [30]:219 [31]:224 [32]:231 [33]:238 [34]:241 [35]:256 [36]:285 [37]:289 [38]:307 [39]:340 -csrVal: -[0]:1 [1]:1 [2]:1 [3]:1 [4]:1 [5]:1 [6]:1 [7]:1 [8]:1 [9]:1 [10]:1 [11]:1 [12]:1 [13]:1 [14]:1 [15]:1 [16]:1 [17]:1 [18]:1 [19]:1 [20]:1 [21]:1 [22]:1 [23]:1 [24]:1 [25]:1 [26]:1 [27]:1 [28]:1 [29]:1 [30]:1 [31]:1 [32]:1 [33]:1 [34]:1 [35]:1 [36]:1 [37]:1 [38]:1 [39]:1 -pretty print: -0 0 0 0 0 0 x x 0 0 x x x 0 0 0 0 0 0 0 -0 0 0 0 0 0 x x x 0 x x x 0 0 0 0 0 0 0 -0 0 0 0 0 0 0 0 0 0 0 0 0 x x x x x 0 0 -0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 0 0 -0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 0 0 x 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 -x x 0 0 0 x 0 x x 0 x x x 0 0 0 0 0 0 0 -x x 0 0 0 0 x 0 x x 0 0 0 0 0 0 0 0 0 x -0 x 0 0 0 0 x x 0 x 0 0 0 0 0 0 0 0 0 0 -0 0 0 0 0 0 0 x x 0 0 0 0 0 0 0 0 0 0 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 x x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 0 0 0 0 0 0 0 0 x x 0 0 x x 0 x 0 0 -0 0 0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 -cscRowInd: -[0]:6 [1]:7 [2]:10 [3]:11 [4]:12 [5]:21 [6]:22 [7]:33 [8]:34 [9]:36 [10]:38 [11]:6 [12]:7 [13]:8 [14]:10 [15]:11 [16]:12 [17]:21 [18]:22 [19]:34 [20]:35 [21]:38 [22]:13 [23]:14 [24]:15 [25]:16 [26]:17 [27]:35 [28]:38 [29]:16 [30]:35 [31]:36 [32]:38 [33]:5 [34]:33 [35]:34 [36]:38 [37]:4 [38]:6 [39]:34 -cscColPtr: -[0]:0 [1]:11 [2]:22 [3]:29 [4]:33 [5]:37 [6]:41 [7]:51 [8]:64 [9]:71 [10]:76 [11]:83 [12]:92 [13]:99 [14]:107 [15]:116 [16]:121 [17]:125 [18]:134 [19]:143 [20]:147 [21]:152 [22]:165 [23]:175 [24]:179 [25]:185 [26]:192 [27]:201 [28]:207 [29]:213 [30]:219 [31]:224 [32]:231 [33]:238 [34]:241 [35]:256 [36]:285 [37]:289 [38]:307 [39]:340 -cscVal: -[0]:1 [1]:1 [2]:1 [3]:1 [4]:1 [5]:1 [6]:1 [7]:1 [8]:1 [9]:1 [10]:1 [11]:1 [12]:1 [13]:1 [14]:1 [15]:1 [16]:1 [17]:1 [18]:1 [19]:1 [20]:1 [21]:1 [22]:1 [23]:1 [24]:1 [25]:1 [26]:1 [27]:1 [28]:1 [29]:1 [30]:1 [31]:1 [32]:1 [33]:1 [34]:1 [35]:1 [36]:1 [37]:1 [38]:1 [39]:1 -pretty print: -0 0 0 0 0 0 x x 0 0 x x x 0 0 0 0 0 0 0 -0 0 0 0 0 0 x x x 0 x x x 0 0 0 0 0 0 0 -0 0 0 0 0 0 0 0 0 0 0 0 0 x x x x x 0 0 -0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 0 0 -0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 0 0 x 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 -x x 0 0 0 x 0 x x 0 x x x 0 0 0 0 0 0 0 -x x 0 0 0 0 x 0 x x 0 0 0 0 0 0 0 0 0 x -0 x 0 0 0 0 x x 0 x 0 0 0 0 0 0 0 0 0 0 -0 0 0 0 0 0 0 x x 0 0 0 0 0 0 0 0 0 0 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 x x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 0 0 0 0 0 0 0 0 x x 0 0 x x 0 x 0 0 -0 0 0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 -csrColInd: -[0]:6 [1]:7 [2]:10 [3]:11 [4]:12 [5]:21 [6]:22 [7]:33 [8]:34 [9]:36 [10]:38 [11]:6 [12]:7 [13]:8 [14]:10 [15]:11 [16]:12 [17]:21 [18]:22 [19]:34 [20]:35 [21]:38 [22]:13 [23]:14 [24]:15 [25]:16 [26]:17 [27]:35 [28]:38 [29]:16 [30]:35 [31]:36 [32]:38 [33]:5 [34]:33 [35]:34 [36]:38 [37]:4 [38]:6 [39]:34 -csrRowPtr: -[0]:0 [1]:11 [2]:22 [3]:29 [4]:33 [5]:37 [6]:41 [7]:51 [8]:64 [9]:71 [10]:76 [11]:83 [12]:92 [13]:99 [14]:107 [15]:116 [16]:121 [17]:125 [18]:134 [19]:143 [20]:147 [21]:152 [22]:165 [23]:175 [24]:179 [25]:185 [26]:192 [27]:201 [28]:207 [29]:213 [30]:219 [31]:224 [32]:231 [33]:238 [34]:241 [35]:256 [36]:285 [37]:289 [38]:307 [39]:340 -csrVal: -[0]:1 [1]:1 [2]:1 [3]:1 [4]:1 [5]:1 [6]:1 [7]:1 [8]:1 [9]:1 [10]:1 [11]:1 [12]:1 [13]:1 [14]:1 [15]:1 [16]:1 [17]:1 [18]:1 [19]:1 [20]:1 [21]:1 [22]:1 [23]:1 [24]:1 [25]:1 [26]:1 [27]:1 [28]:1 [29]:1 [30]:1 [31]:1 [32]:1 [33]:1 [34]:1 [35]:1 [36]:1 [37]:1 [38]:1 [39]:1 -pretty print: -0 0 0 0 0 0 x x 0 0 x x x 0 0 0 0 0 0 0 -0 0 0 0 0 0 x x x 0 x x x 0 0 0 0 0 0 0 -0 0 0 0 0 0 0 0 0 0 0 0 0 x x x x x 0 0 -0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 0 0 -0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 0 0 x 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 -x x 0 0 0 x 0 x x 0 x x x 0 0 0 0 0 0 0 -x x 0 0 0 0 x 0 x x 0 0 0 0 0 0 0 0 0 x -0 x 0 0 0 0 x x 0 x 0 0 0 0 0 0 0 0 0 0 -0 0 0 0 0 0 0 x x 0 0 0 0 0 0 0 0 0 0 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 x x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 0 0 0 0 0 0 0 0 x x 0 0 x x 0 x 0 0 -0 0 0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 -cscRowInd: -[0]:6 [1]:7 [2]:10 [3]:11 [4]:12 [5]:21 [6]:22 [7]:33 [8]:34 [9]:36 [10]:38 [11]:6 [12]:7 [13]:8 [14]:10 [15]:11 [16]:12 [17]:21 [18]:22 [19]:34 [20]:35 [21]:38 [22]:13 [23]:14 [24]:15 [25]:16 [26]:17 [27]:35 [28]:38 [29]:16 [30]:35 [31]:36 [32]:38 [33]:5 [34]:33 [35]:34 [36]:38 [37]:4 [38]:6 [39]:34 -cscColPtr: -[0]:0 [1]:11 [2]:22 [3]:29 [4]:33 [5]:37 [6]:41 [7]:51 [8]:64 [9]:71 [10]:76 [11]:83 [12]:92 [13]:99 [14]:107 [15]:116 [16]:121 [17]:125 [18]:134 [19]:143 [20]:147 [21]:152 [22]:165 [23]:175 [24]:179 [25]:185 [26]:192 [27]:201 [28]:207 [29]:213 [30]:219 [31]:224 [32]:231 [33]:238 [34]:241 [35]:256 [36]:285 [37]:289 [38]:307 [39]:340 -cscVal: -[0]:1 [1]:1 [2]:1 [3]:1 [4]:1 [5]:1 [6]:1 [7]:1 [8]:1 [9]:1 [10]:1 [11]:1 [12]:1 [13]:1 [14]:1 [15]:1 [16]:1 [17]:1 [18]:1 [19]:1 [20]:1 [21]:1 [22]:1 [23]:1 [24]:1 [25]:1 [26]:1 [27]:1 [28]:1 [29]:1 [30]:1 [31]:1 [32]:1 [33]:1 [34]:1 [35]:1 [36]:1 [37]:1 [38]:1 [39]:1 -pretty print: -0 0 0 0 0 0 x x 0 0 x x x 0 0 0 0 0 0 0 -0 0 0 0 0 0 x x x 0 x x x 0 0 0 0 0 0 0 -0 0 0 0 0 0 0 0 0 0 0 0 0 x x x x x 0 0 -0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 0 0 -0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 0 0 x 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 -x x 0 0 0 x 0 x x 0 x x x 0 0 0 0 0 0 0 -x x 0 0 0 0 x 0 x x 0 0 0 0 0 0 0 0 0 x -0 x 0 0 0 0 x x 0 x 0 0 0 0 0 0 0 0 0 0 -0 0 0 0 0 0 0 x x 0 0 0 0 0 0 0 0 0 0 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 x x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 0 0 0 0 0 0 0 0 x x 0 0 x x 0 x 0 0 -0 0 0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 -Allocate 40 -Do not allocate 0 0 -Do not allocate 0 0 -===End mxm=== -csrColInd: -[0]:0 [1]:1 [2]:2 [3]:3 [4]:4 [5]:5 [6]:6 [7]:7 [8]:8 [9]:9 [10]:10 [11]:11 [12]:12 [13]:13 [14]:14 [15]:15 [16]:16 [17]:17 [18]:18 [19]:19 [20]:20 [21]:21 [22]:22 [23]:23 [24]:24 [25]:25 [26]:26 [27]:27 [28]:28 [29]:29 [30]:30 [31]:31 [32]:32 [33]:33 [34]:34 [35]:35 [36]:36 [37]:37 [38]:38 [39]:0 -csrRowPtr: -[0]:0 [1]:39 [2]:77 [3]:113 [4]:150 [5]:187 [6]:225 [7]:264 [8]:303 [9]:341 [10]:379 [11]:418 [12]:457 [13]:495 [14]:532 [15]:569 [16]:606 [17]:643 [18]:680 [19]:717 [20]:753 [21]:790 [22]:829 [23]:868 [24]:904 [25]:940 [26]:977 [27]:1014 [28]:1051 [29]:1087 [30]:1124 [31]:1161 [32]:1198 [33]:1235 [34]:1249 [35]:1288 [36]:1324 [37]:1343 [38]:1372 [39]:1411 -csrVal: -[0]:11 [1]:9 [2]:1 [3]:2 [4]:3 [5]:3 [6]:6 [7]:5 [8]:4 [9]:3 [10]:3 [11]:3 [12]:3 [13]:1 [14]:1 [15]:1 [16]:1 [17]:1 [18]:3 [19]:2 [20]:2 [21]:3 [22]:3 [23]:2 [24]:1 [25]:2 [26]:2 [27]:2 [28]:1 [29]:3 [30]:2 [31]:3 [32]:3 [33]:1 [34]:8 [35]:6 [36]:2 [37]:5 [38]:7 [39]:9 -pretty print: -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -csrColInd: -[0]:0 [1]:1 [2]:2 [3]:3 [4]:4 [5]:5 [6]:6 [7]:7 [8]:8 [9]:9 [10]:10 [11]:11 [12]:12 [13]:13 [14]:14 [15]:15 [16]:16 [17]:17 [18]:18 [19]:19 [20]:20 [21]:21 [22]:22 [23]:23 [24]:24 [25]:25 [26]:26 [27]:27 [28]:28 [29]:29 [30]:30 [31]:31 [32]:32 [33]:33 [34]:34 [35]:35 [36]:36 [37]:37 [38]:38 [39]:0 -csrRowPtr: -[0]:0 [1]:39 [2]:77 [3]:113 [4]:150 [5]:187 [6]:225 [7]:264 [8]:303 [9]:341 [10]:379 [11]:418 [12]:457 [13]:495 [14]:532 [15]:569 [16]:606 [17]:643 [18]:680 [19]:717 [20]:753 [21]:790 [22]:829 [23]:868 [24]:904 [25]:940 [26]:977 [27]:1014 [28]:1051 [29]:1087 [30]:1124 [31]:1161 [32]:1198 [33]:1235 [34]:1249 [35]:1288 [36]:1324 [37]:1343 [38]:1372 [39]:1411 -csrVal: -[0]:11 [1]:9 [2]:1 [3]:2 [4]:3 [5]:3 [6]:6 [7]:5 [8]:4 [9]:3 [10]:3 [11]:3 [12]:3 [13]:1 [14]:1 [15]:1 [16]:1 [17]:1 [18]:3 [19]:2 [20]:2 [21]:3 [22]:3 [23]:2 [24]:1 [25]:2 [26]:2 [27]:2 [28]:1 [29]:3 [30]:2 [31]:3 [32]:3 [33]:1 [34]:8 [35]:6 [36]:2 [37]:5 [38]:7 [39]:9 -pretty print: -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -Allocate 40 -Allocate 408 -Allocate 408 -Allocate 40 -Allocate 408 -Allocate 408 -===Begin mxm=== -csrColInd: -[0]:6 [1]:7 [2]:10 [3]:11 [4]:12 [5]:21 [6]:22 [7]:33 [8]:34 [9]:36 [10]:38 [11]:6 [12]:7 [13]:8 [14]:10 [15]:11 [16]:12 [17]:21 [18]:22 [19]:34 [20]:35 [21]:38 [22]:13 [23]:14 [24]:15 [25]:16 [26]:17 [27]:35 [28]:38 [29]:16 [30]:35 [31]:36 [32]:38 [33]:5 [34]:33 [35]:34 [36]:38 [37]:4 [38]:6 [39]:34 -csrRowPtr: -[0]:0 [1]:11 [2]:22 [3]:29 [4]:33 [5]:37 [6]:41 [7]:51 [8]:64 [9]:71 [10]:76 [11]:83 [12]:92 [13]:99 [14]:107 [15]:116 [16]:121 [17]:125 [18]:134 [19]:143 [20]:147 [21]:152 [22]:165 [23]:175 [24]:179 [25]:185 [26]:192 [27]:201 [28]:207 [29]:213 [30]:219 [31]:224 [32]:231 [33]:238 [34]:241 [35]:256 [36]:285 [37]:289 [38]:307 [39]:340 -csrVal: -[0]:1 [1]:1 [2]:1 [3]:1 [4]:1 [5]:1 [6]:1 [7]:1 [8]:1 [9]:1 [10]:1 [11]:1 [12]:1 [13]:1 [14]:1 [15]:1 [16]:1 [17]:1 [18]:1 [19]:1 [20]:1 [21]:1 [22]:1 [23]:1 [24]:1 [25]:1 [26]:1 [27]:1 [28]:1 [29]:1 [30]:1 [31]:1 [32]:1 [33]:1 [34]:1 [35]:1 [36]:1 [37]:1 [38]:1 [39]:1 -pretty print: -0 0 0 0 0 0 x x 0 0 x x x 0 0 0 0 0 0 0 -0 0 0 0 0 0 x x x 0 x x x 0 0 0 0 0 0 0 -0 0 0 0 0 0 0 0 0 0 0 0 0 x x x x x 0 0 -0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 0 0 -0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 0 0 x 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 -x x 0 0 0 x 0 x x 0 x x x 0 0 0 0 0 0 0 -x x 0 0 0 0 x 0 x x 0 0 0 0 0 0 0 0 0 x -0 x 0 0 0 0 x x 0 x 0 0 0 0 0 0 0 0 0 0 -0 0 0 0 0 0 0 x x 0 0 0 0 0 0 0 0 0 0 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 x x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 0 0 0 0 0 0 0 0 x x 0 0 x x 0 x 0 0 -0 0 0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 -csrColInd: -[0]:6 [1]:7 [2]:10 [3]:11 [4]:12 [5]:21 [6]:22 [7]:33 [8]:34 [9]:36 [10]:38 [11]:6 [12]:7 [13]:8 [14]:10 [15]:11 [16]:12 [17]:21 [18]:22 [19]:34 [20]:35 [21]:38 [22]:13 [23]:14 [24]:15 [25]:16 [26]:17 [27]:35 [28]:38 [29]:16 [30]:35 [31]:36 [32]:38 [33]:5 [34]:33 [35]:34 [36]:38 [37]:4 [38]:6 [39]:34 -csrRowPtr: -[0]:0 [1]:11 [2]:22 [3]:29 [4]:33 [5]:37 [6]:41 [7]:51 [8]:64 [9]:71 [10]:76 [11]:83 [12]:92 [13]:99 [14]:107 [15]:116 [16]:121 [17]:125 [18]:134 [19]:143 [20]:147 [21]:152 [22]:165 [23]:175 [24]:179 [25]:185 [26]:192 [27]:201 [28]:207 [29]:213 [30]:219 [31]:224 [32]:231 [33]:238 [34]:241 [35]:256 [36]:285 [37]:289 [38]:307 [39]:340 -csrVal: -[0]:1 [1]:1 [2]:1 [3]:1 [4]:1 [5]:1 [6]:1 [7]:1 [8]:1 [9]:1 [10]:1 [11]:1 [12]:1 [13]:1 [14]:1 [15]:1 [16]:1 [17]:1 [18]:1 [19]:1 [20]:1 [21]:1 [22]:1 [23]:1 [24]:1 [25]:1 [26]:1 [27]:1 [28]:1 [29]:1 [30]:1 [31]:1 [32]:1 [33]:1 [34]:1 [35]:1 [36]:1 [37]:1 [38]:1 [39]:1 -pretty print: -0 0 0 0 0 0 x x 0 0 x x x 0 0 0 0 0 0 0 -0 0 0 0 0 0 x x x 0 x x x 0 0 0 0 0 0 0 -0 0 0 0 0 0 0 0 0 0 0 0 0 x x x x x 0 0 -0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 0 0 -0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 0 0 x 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 -x x 0 0 0 x 0 x x 0 x x x 0 0 0 0 0 0 0 -x x 0 0 0 0 x 0 x x 0 0 0 0 0 0 0 0 0 x -0 x 0 0 0 0 x x 0 x 0 0 0 0 0 0 0 0 0 0 -0 0 0 0 0 0 0 x x 0 0 0 0 0 0 0 0 0 0 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 x 0 -x x 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 x x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 -0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 x 0 -0 0 0 0 0 0 0 0 0 0 x x 0 0 x x 0 x 0 0 -0 0 0 0 0 0 0 x 0 0 0 0 0 0 0 0 0 0 0 0 -Allocate 40 -Do not allocate 0 0 -Do not allocate 0 0 -===End mxm=== -csrColInd: -[0]:0 [1]:1 [2]:2 [3]:3 [4]:4 [5]:5 [6]:6 [7]:7 [8]:8 [9]:9 [10]:10 [11]:11 [12]:12 [13]:13 [14]:14 [15]:15 [16]:16 [17]:17 [18]:18 [19]:19 [20]:20 [21]:21 [22]:22 [23]:23 [24]:24 [25]:25 [26]:26 [27]:27 [28]:28 [29]:29 [30]:30 [31]:31 [32]:32 [33]:33 [34]:34 [35]:35 [36]:36 [37]:37 [38]:38 [39]:0 -csrRowPtr: -[0]:0 [1]:39 [2]:77 [3]:113 [4]:150 [5]:187 [6]:225 [7]:264 [8]:303 [9]:341 [10]:379 [11]:418 [12]:457 [13]:495 [14]:532 [15]:569 [16]:606 [17]:643 [18]:680 [19]:717 [20]:753 [21]:790 [22]:829 [23]:868 [24]:904 [25]:940 [26]:977 [27]:1014 [28]:1051 [29]:1087 [30]:1124 [31]:1161 [32]:1198 [33]:1235 [34]:1249 [35]:1288 [36]:1324 [37]:1343 [38]:1372 [39]:1411 -csrVal: -[0]:11 [1]:9 [2]:1 [3]:2 [4]:3 [5]:3 [6]:6 [7]:5 [8]:4 [9]:3 [10]:3 [11]:3 [12]:3 [13]:1 [14]:1 [15]:1 [16]:1 [17]:1 [18]:3 [19]:2 [20]:2 [21]:3 [22]:3 [23]:2 [24]:1 [25]:2 [26]:2 [27]:2 [28]:1 [29]:3 [30]:2 [31]:3 [32]:3 [33]:1 [34]:8 [35]:6 [36]:2 [37]:5 [38]:7 [39]:9 -pretty print: -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x -x x x x x x x x x x x x x x x x x x x x From 04138ae8fefff6eb01a32376ff46c14771ce1d22 Mon Sep 17 00:00:00 2001 From: Richard Lettich Date: Tue, 22 Jun 2021 17:55:36 -0700 Subject: [PATCH 09/12] update submodule --- ext/GALATIC | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ext/GALATIC b/ext/GALATIC index f7537d4..fca085e 160000 --- a/ext/GALATIC +++ b/ext/GALATIC @@ -1 +1 @@ -Subproject commit f7537d409b8fd2449a1c2d9a5dda7cd51371e2ea +Subproject commit fca085e2bcf116129e6c9f974a16a3a0233449c0 From 50fa929a8610dbe58162fa69d2b0f02b19cc206d Mon Sep 17 00:00:00 2001 From: Richard Lettich Date: Wed, 23 Jun 2021 01:03:31 +0000 Subject: [PATCH 10/12] renamed first and second to something that won't conflict with thrust --- graphblas/stddef.hpp | 9 +++------ 1 file changed, 3 insertions(+), 6 deletions(-) diff --git a/graphblas/stddef.hpp b/graphblas/stddef.hpp index 99ac3cc..815caa5 100644 --- a/graphblas/stddef.hpp +++ b/graphblas/stddef.hpp @@ -73,22 +73,19 @@ struct less_equal { return lhs <= rhs; } }; -namespace fixme { template -struct first { +struct left_arg { inline GRB_HOST_DEVICE T_out operator()(T_in1 lhs, T_in2 rhs) { return lhs; } }; -} -namespace fixme { template -struct second { +struct right_arg { inline GRB_HOST_DEVICE T_out operator()(T_in1 lhs, T_in2 rhs) { return rhs; } }; -} + template struct minimum { From ce4a855cc0fcf06f1dedaf7b0f2a6fe1006f3494 Mon Sep 17 00:00:00 2001 From: Richard Lettich Date: Tue, 29 Jun 2021 20:10:27 +0000 Subject: [PATCH 11/12] fixed formatting issues --- .gitmodules | 2 +- CMakeLists.txt | 4 ++-- ext/GALATIC | 2 +- graphblas/backend/cuda/operations.hpp | 2 +- graphblas/backend/cuda/spgemm.hpp | 19 +------------------ 5 files changed, 6 insertions(+), 23 deletions(-) diff --git a/.gitmodules b/.gitmodules index b003727..455917c 100644 --- a/.gitmodules +++ b/.gitmodules @@ -1,6 +1,6 @@ [submodule "ext/moderngpu"] path = ext/moderngpu - url = https://ctcyang@github.com/ctcyang/moderngpu.git + url = git@github.com:ctcyang/moderngpu.git [submodule "ext/GALATIC"] path = ext/GALATIC url = git@github.com:richardlett/GALATIC.git diff --git a/CMakeLists.txt b/CMakeLists.txt index b027944..4653712 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -13,13 +13,13 @@ set( PROJ_PATH ${CMAKE_SOURCE_DIR}) set( PROJ_OUT_PATH ${CMAKE_BINARY_DIR}) set( PROJ_HEADERS "" ) set( PROJ_LIBRARIES "" ) -set( PROJ_INCLUDES "./" "ext/moderngpu/include") +set( PROJ_INCLUDES "./" "ext/moderngpu/include" "ext") set( mgpu_SRC_FILES "ext/moderngpu/src/mgpucontext.cu" "ext/moderngpu/src/mgpuutil.cpp") set( CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/../bin ) #set( CUDA_CURAND_LIBRARY "$ENV{CUDA_HOME}/lib64/libcurand.so" ) #set( CUDA_CUBLAS_LIBRARY "$ENV{CUDA_HOME}/lib64/libcublas.so" ) set( CUDA_CUSPARSE_LIBRARY "$ENV{CUDA_HOME}/lib64/libcusparse.so" ) -#FILE( GLOB_RECURSE PROJ_SOURCES graphblas/*.cu ../graphblas/*.cpp ) +#FILE( G LOB_RECURSE PROJ_SOURCES graphblas/*.cu ../graphblas/*.cpp ) #FILE( GLOB_RECURSE PROJ_LIBRARIES ext/cublas1.1/*.cu ) FILE( GLOB_RECURSE PROJ_HEADERS graphblas/*.hpp) # nvcc flags diff --git a/ext/GALATIC b/ext/GALATIC index fca085e..e82d65a 160000 --- a/ext/GALATIC +++ b/ext/GALATIC @@ -1 +1 @@ -Subproject commit fca085e2bcf116129e6c9f974a16a3a0233449c0 +Subproject commit e82d65a99006f60e98330daa0424319c94b62bd7 diff --git a/graphblas/backend/cuda/operations.hpp b/graphblas/backend/cuda/operations.hpp index 80dfaf7..951f190 100644 --- a/graphblas/backend/cuda/operations.hpp +++ b/graphblas/backend/cuda/operations.hpp @@ -52,7 +52,7 @@ Info mxm(Matrix* C, &B->sparse_, desc)); else { if (s_mode != GrB_GALATIC) { - std::cout << R"(Unknown mode (Options are: "cusspare2" and "galatic"; defaulting to galatic)" << std::endl; + std::cout << R"(Unknown mode (Options are: "cuspare2" and "galatic"; defaulting to galatic)" << std::endl; } CHECK(GALATIC_spgemm(&C->sparse_, op, diff --git a/graphblas/backend/cuda/spgemm.hpp b/graphblas/backend/cuda/spgemm.hpp index 0a3b4ba..c811356 100644 --- a/graphblas/backend/cuda/spgemm.hpp +++ b/graphblas/backend/cuda/spgemm.hpp @@ -1,24 +1,16 @@ #ifndef GRAPHBLAS_BACKEND_CUDA_SPGEMM_HPP_ #define GRAPHBLAS_BACKEND_CUDA_SPGEMM_HPP_ - -#include "../../../ext/GALATIC/include/dCSR.cuh" -#include "../../../ext/GALATIC/include/SemiRingInterface.h" -#include "../../../ext/GALATIC/source/device/Multiply.cuh" +#include #include "graphblas/backend/cuda/sparse_matrix.hpp" - #include #include #include #include - - - - namespace graphblas { namespace backend { template @@ -218,11 +210,8 @@ Info GALATIC_spgemm(SparseMatrix* C, matrixToGalatic(A, leftInputMatrixGPU); matrixToGalatic(B, rightInputMatrixGPU); - - GPUMatrixMatrixMultiplyTraits DefaultTraits; - // GALATIC has its own semiring interface; // GalaticSemiring is a shim here for conversion of graphblast-style // SemiringT type. GalaticSemiring definition is above this function @@ -231,8 +220,6 @@ Info GALATIC_spgemm(SparseMatrix* C, ExecutionStats stats; try { - - Desc_value nt_mode; CHECK(desc->get(GrB_NT, &nt_mode)); const int num_threads = static_cast(nt_mode); @@ -288,22 +275,18 @@ Info GALATIC_spgemm(SparseMatrix* C, nullizeGalaticMatrix(leftInputMatrixGPU); nullizeGalaticMatrix(rightInputMatrixGPU); - if (C->h_csrRowPtr_ == NULL) C->h_csrRowPtr_ = reinterpret_cast(malloc((A_nrows+1)* sizeof(Index))); C->h_csrColInd_ = reinterpret_cast(malloc(C->ncapacity_*sizeof(Index))); C->h_csrVal_ = reinterpret_cast(malloc(C->ncapacity_*sizeof(c))); - C->need_update_ = true; // Set flag that we need to copy data from GPU C->csr_initialized_ = true; C->csc_initialized_ = false; return GrB_SUCCESS; } - - template Info cusparse_spgemm(SparseMatrix* C, From dac2a22f40b49b702ba34462007201f0ae50f690 Mon Sep 17 00:00:00 2001 From: Richard Lettich Date: Fri, 2 Jul 2021 21:26:16 +0000 Subject: [PATCH 12/12] formatting/spelling fixes --- .gitmodules | 4 ++-- CMakeLists.txt | 4 ++-- graphblas/backend/cuda/operations.hpp | 2 +- graphblas/backend/cuda/spgemm.hpp | 2 +- 4 files changed, 6 insertions(+), 6 deletions(-) diff --git a/.gitmodules b/.gitmodules index 455917c..b6ccde3 100644 --- a/.gitmodules +++ b/.gitmodules @@ -1,6 +1,6 @@ [submodule "ext/moderngpu"] path = ext/moderngpu - url = git@github.com:ctcyang/moderngpu.git + url = https://ctcyang@github.com/ctcyang/moderngpu.git [submodule "ext/GALATIC"] path = ext/GALATIC - url = git@github.com:richardlett/GALATIC.git + url = https://ctcyang@github.com/richardlett/GALATIC.git diff --git a/CMakeLists.txt b/CMakeLists.txt index 4653712..e3ce001 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -13,13 +13,13 @@ set( PROJ_PATH ${CMAKE_SOURCE_DIR}) set( PROJ_OUT_PATH ${CMAKE_BINARY_DIR}) set( PROJ_HEADERS "" ) set( PROJ_LIBRARIES "" ) -set( PROJ_INCLUDES "./" "ext/moderngpu/include" "ext") +set( PROJ_INCLUDES "./" "ext/moderngpu/include" "ext/GALATIC") set( mgpu_SRC_FILES "ext/moderngpu/src/mgpucontext.cu" "ext/moderngpu/src/mgpuutil.cpp") set( CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/../bin ) #set( CUDA_CURAND_LIBRARY "$ENV{CUDA_HOME}/lib64/libcurand.so" ) #set( CUDA_CUBLAS_LIBRARY "$ENV{CUDA_HOME}/lib64/libcublas.so" ) set( CUDA_CUSPARSE_LIBRARY "$ENV{CUDA_HOME}/lib64/libcusparse.so" ) -#FILE( G LOB_RECURSE PROJ_SOURCES graphblas/*.cu ../graphblas/*.cpp ) +#FILE( GLOB_RECURSE PROJ_SOURCES graphblas/*.cu ../graphblas/*.cpp ) #FILE( GLOB_RECURSE PROJ_LIBRARIES ext/cublas1.1/*.cu ) FILE( GLOB_RECURSE PROJ_HEADERS graphblas/*.hpp) # nvcc flags diff --git a/graphblas/backend/cuda/operations.hpp b/graphblas/backend/cuda/operations.hpp index 951f190..33af949 100644 --- a/graphblas/backend/cuda/operations.hpp +++ b/graphblas/backend/cuda/operations.hpp @@ -52,7 +52,7 @@ Info mxm(Matrix* C, &B->sparse_, desc)); else { if (s_mode != GrB_GALATIC) { - std::cout << R"(Unknown mode (Options are: "cuspare2" and "galatic"; defaulting to galatic)" << std::endl; + std::cout << R"(Unknown mode (Options are: "cusparse2" and "galatic"; defaulting to galatic)" << std::endl; } CHECK(GALATIC_spgemm(&C->sparse_, op, diff --git a/graphblas/backend/cuda/spgemm.hpp b/graphblas/backend/cuda/spgemm.hpp index c811356..66e78e6 100644 --- a/graphblas/backend/cuda/spgemm.hpp +++ b/graphblas/backend/cuda/spgemm.hpp @@ -1,7 +1,7 @@ #ifndef GRAPHBLAS_BACKEND_CUDA_SPGEMM_HPP_ #define GRAPHBLAS_BACKEND_CUDA_SPGEMM_HPP_ -#include +#include "GALATICMinimumIncludes.cuh" #include "graphblas/backend/cuda/sparse_matrix.hpp"