diff --git a/CHANGELOG.md b/CHANGELOG.md index 96a30057540..0ef62d942da 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -5,6 +5,7 @@ ## Improvements - PR #1227 Pin cmake policies to cmake 3.17 version +- PR #1269 Removed old db code that was not being used ## Bug Fixes - PR #1242 Calling gunrock cmake using explicit -D options, re-enabling C++ tests diff --git a/build.sh b/build.sh index ae3ad575227..ffe3f72418c 100755 --- a/build.sh +++ b/build.sh @@ -39,6 +39,7 @@ CUGRAPH_BUILD_DIR=${REPODIR}/python/build BUILD_DIRS="${LIBCUGRAPH_BUILD_DIR} ${CUGRAPH_BUILD_DIR}" # Set defaults for vars modified by flags to this script +ARG_COUNT=${NUMARGS} VERBOSE="" BUILD_TYPE=Release INSTALL_TARGET=install @@ -48,7 +49,7 @@ BUILD_DISABLE_DEPRECATION_WARNING=ON # FIXME: if PREFIX is not set, check CONDA_PREFIX, but there is no fallback # from there! INSTALL_PREFIX=${PREFIX:=${CONDA_PREFIX}} -PARALLEL_LEVEL=${PARALLEL_LEVEL:=""} +PARALLEL_LEVEL=${PARALLEL_LEVEL:=`nproc`} BUILD_ABI=${BUILD_ABI:=ON} function hasArg { @@ -73,15 +74,19 @@ fi # Process flags if hasArg -v; then VERBOSE=1 + ARG_COUNT=$((ARG_COUNT -1)) fi if hasArg -g; then BUILD_TYPE=Debug + ARG_COUNT=$((ARG_COUNT -1)) fi if hasArg -n; then INSTALL_TARGET="" + ARG_COUNT=$((ARG_COUNT -1)) fi if hasArg --show_depr_warn; then BUILD_DISABLE_DEPRECATION_WARNING=OFF + ARG_COUNT=$((ARG_COUNT -1)) fi # If clean given, run it prior to any other steps @@ -100,7 +105,7 @@ fi ################################################################################ # Configure, build, and install libcugraph -if (( ${NUMARGS} == 0 )) || hasArg libcugraph; then +if (( ${ARG_COUNT} == 0 )) || hasArg libcugraph; then mkdir -p ${LIBCUGRAPH_BUILD_DIR} cd ${LIBCUGRAPH_BUILD_DIR} @@ -111,7 +116,7 @@ if (( ${NUMARGS} == 0 )) || hasArg libcugraph; then fi # Build and install the cugraph Python package -if (( ${NUMARGS} == 0 )) || hasArg cugraph; then +if (( ${ARG_COUNT} == 0 )) || hasArg cugraph; then cd ${REPODIR}/python if [[ ${INSTALL_TARGET} != "" ]]; then @@ -124,7 +129,7 @@ fi ################################################################################ # Build the docs -if (( ${NUMARGS} == 0 )) || hasArg docs; then +if (( ${ARG_COUNT} == 0 )) || hasArg docs; then if [ ! -d ${LIBCUGRAPH_BUILD_DIR} ]; then mkdir -p ${LIBCUGRAPH_BUILD_DIR} diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 2d6b9facd8b..d4ea1cd5a89 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -158,19 +158,6 @@ if(OpenMP_FOUND) set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Xcompiler=${OpenMP_CXX_FLAGS}") endif(OpenMP_FOUND) -################################################################################################### -# - find libcypher-parser ------------------------------------------------------------------------- - -find_path(LIBCYPHERPARSER_INCLUDE "cypher-parser.h" - HINTS "$ENV{CONDA_PREFIX}/include") - -find_library(LIBCYPHERPARSER_LIBRARY STATIC "libcypher-parser.a" - HINTS "$ENV{CONDA_PREFIX}/lib") - -add_library(libcypher-parser STATIC IMPORTED ${LIBCYPHERPARSER_LIBRARY}) -if (LIBCYPHERPARSER_INCLUDE AND LIBCYPHERPARSER_LIBRARY) - set_target_properties(libcypher-parser PROPERTIES IMPORTED_LOCATION ${LIBCYPHERPARSER_LIBRARY}) -endif (LIBCYPHERPARSER_INCLUDE AND LIBCYPHERPARSER_LIBRARY) ################################################################################################### # - find gtest ------------------------------------------------------------------------------------ @@ -359,9 +346,6 @@ link_directories( "${CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES}") add_library(cugraph SHARED - src/db/db_object.cu - src/db/db_parser_integration_test.cu - src/db/db_operators.cu src/utilities/spmv_1D.cu src/utilities/cython.cu src/structure/graph.cu @@ -410,7 +394,6 @@ target_include_directories(cugraph "${CUCO_INCLUDE_DIR}" "${LIBCUDACXX_INCLUDE_DIR}" "${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES}" - "${LIBCYPHERPARSER_INCLUDE}" "${Boost_INCLUDE_DIRS}" "${RMM_INCLUDE}" "${CMAKE_CURRENT_SOURCE_DIR}/../thirdparty" @@ -430,7 +413,7 @@ target_include_directories(cugraph # - link libraries -------------------------------------------------------------------------------- target_link_libraries(cugraph PRIVATE - gunrock cublas cusparse curand cusolver cudart cuda ${LIBCYPHERPARSER_LIBRARY} ${MPI_CXX_LIBRARIES} ${NCCL_LIBRARIES}) + gunrock cublas cusparse curand cusolver cudart cuda ${NCCL_LIBRARIES}) if(OpenMP_CXX_FOUND) target_link_libraries(cugraph PRIVATE diff --git a/cpp/src/db/db_object.cu b/cpp/src/db/db_object.cu deleted file mode 100644 index 31c149f3503..00000000000 --- a/cpp/src/db/db_object.cu +++ /dev/null @@ -1,444 +0,0 @@ -/* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include - -#include - -#include -#include - -#include -#include - -#include - -namespace cugraph { -namespace db { -// Define kernel for copying run length encoded values into offset slots. -template -__global__ void offsetsKernel(T runCounts, T* unique, T* counts, T* offsets) -{ - uint64_t tid = threadIdx.x + blockIdx.x * blockDim.x; - if (tid < runCounts) offsets[unique[tid]] = counts[tid]; -} - -template -db_pattern_entry::db_pattern_entry(std::string variable) -{ - is_var = true; - variableName = variable; -} - -template -db_pattern_entry::db_pattern_entry(idx_t constant) -{ - is_var = false; - constantValue = constant; -} - -template -db_pattern_entry::db_pattern_entry(const db_pattern_entry& other) -{ - is_var = other.is_var; - constantValue = other.constantValue; - variableName = other.variableName; -} - -template -db_pattern_entry& db_pattern_entry::operator=(const db_pattern_entry& other) -{ - is_var = other.is_var; - constantValue = other.constantValue; - variableName = other.variableName; - return *this; -} - -template -bool db_pattern_entry::isVariable() const -{ - return is_var; -} - -template -idx_t db_pattern_entry::getConstant() const -{ - return constantValue; -} - -template -std::string db_pattern_entry::getVariable() const -{ - return variableName; -} - -template class db_pattern_entry; -template class db_pattern_entry; - -template -db_pattern::db_pattern() -{ -} - -template -db_pattern::db_pattern(const db_pattern& other) -{ - for (size_t i = 0; i < other.entries.size(); i++) { entries.push_back(other.getEntry(i)); } -} - -template -db_pattern& db_pattern::operator=(const db_pattern& other) -{ - entries = other.entries; - return *this; -} - -template -int db_pattern::getSize() const -{ - return entries.size(); -} - -template -const db_pattern_entry& db_pattern::getEntry(int position) const -{ - return entries[position]; -} - -template -void db_pattern::addEntry(db_pattern_entry& entry) -{ - entries.push_back(entry); -} - -template -bool db_pattern::isAllConstants() -{ - for (size_t i = 0; i < entries.size(); i++) - if (entries[i].isVariable()) return false; - return true; -} - -template class db_pattern; -template class db_pattern; - -template -db_column_index::db_column_index(rmm::device_buffer&& off, rmm::device_buffer&& ind) -{ - offsets = std::move(off); - indirection = std::move(ind); -} - -template -void db_column_index::resetData(rmm::device_buffer&& off, rmm::device_buffer&& ind) -{ - offsets = std::move(off); - indirection = std::move(ind); -} - -template -idx_t* db_column_index::getOffsets() -{ - return reinterpret_cast(offsets.data()); -} - -template -idx_t db_column_index::getOffsetsSize() -{ - return offsets.size() / sizeof(idx_t); -} - -template -idx_t* db_column_index::getIndirection() -{ - return reinterpret_cast(indirection.data()); -} - -template -idx_t db_column_index::getIndirectionSize() -{ - return indirection.size() / sizeof(idx_t); -} - -template -std::string db_column_index::toString() -{ - std::stringstream ss; - ss << "db_column_index:\n"; - ss << "Offsets: "; - std::vector hostOff(getOffsetsSize()); - idx_t* hostOffsets = reinterpret_cast(hostOff.data()); - CUDA_TRY( - cudaMemcpy(hostOffsets, offsets.data(), sizeof(idx_t) * getOffsetsSize(), cudaMemcpyDefault)); - for (idx_t i = 0; i < getOffsetsSize(); i++) { ss << hostOff[i] << " "; } - ss << "\nIndirection: "; - std::vector hostInd(getIndirectionSize()); - idx_t* hostIndirection = reinterpret_cast(hostInd.data()); - CUDA_TRY(cudaMemcpy( - hostIndirection, indirection.data(), sizeof(idx_t) * getIndirectionSize(), cudaMemcpyDefault)); - for (idx_t i = 0; i < getIndirectionSize(); i++) { ss << hostInd[i] << " "; } - ss << "\n"; - return ss.str(); -} - -template class db_column_index; -template class db_column_index; - -template -db_result::db_result() -{ - dataValid = false; - columnSize = 0; -} - -template -db_result::db_result(db_result&& other) -{ - dataValid = other.dataValid; - columns = std::move(other.columns); - names = std::move(other.names); - other.dataValid = false; -} - -template -db_result& db_result::operator=(db_result&& other) -{ - dataValid = other.dataValid; - columns = std::move(other.columns); - names = std::move(other.names); - other.dataValid = false; - return *this; -} - -template -idx_t db_result::getSize() -{ - return columnSize; -} - -template -idx_t* db_result::getData(std::string idx) -{ - CUGRAPH_EXPECTS(dataValid, "Data not valid"); - - idx_t* returnPtr = nullptr; - for (size_t i = 0; i < names.size(); i++) - if (names[i] == idx) returnPtr = reinterpret_cast(columns[i].data()); - return returnPtr; -} - -template -void db_result::addColumn(std::string columnName) -{ - CUGRAPH_EXPECTS(!dataValid, "Cannot add a column to an allocated result."); - names.push_back(columnName); -} - -template -void db_result::allocateColumns(idx_t size) -{ - CUGRAPH_EXPECTS(!dataValid, "Already allocated columns"); - - for (size_t i = 0; i < names.size(); i++) { - rmm::device_buffer col(sizeof(idx_t) * size); - columns.push_back(std::move(col)); - } - dataValid = true; - columnSize = size; -} - -template -std::string db_result::toString() -{ - std::stringstream ss; - ss << "db_result with " << columns.size() << " columns of length " << columnSize << "\n"; - for (size_t i = 0; i < columns.size(); i++) ss << names[i] << " "; - ss << "\n"; - std::vector> hostColumns; - hostColumns.resize(columns.size()); - for (size_t i = 0; i < columns.size(); i++) { - hostColumns[i].resize(columnSize); - CUDA_TRY(cudaMemcpy( - hostColumns[i].data(), columns[i].data(), sizeof(idx_t) * columnSize, cudaMemcpyDefault)); - } - for (idx_t i = 0; i < columnSize; i++) { - for (size_t j = 0; j < hostColumns.size(); j++) ss << hostColumns[j][i] << " "; - ss << "\n"; - } - return ss.str(); -} - -template class db_result; -template class db_result; - -template -db_table::db_table() -{ - column_size = 0; -} - -template -void db_table::addColumn(std::string name) -{ - CUGRAPH_EXPECTS(column_size == 0, "Can't add a column to a non-empty table"); - - rmm::device_buffer _col; - columns.push_back(std::move(_col)); - names.push_back(name); - indices.resize(indices.size() + 1); -} - -template -void db_table::addEntry(db_pattern& pattern) -{ - CUGRAPH_EXPECTS(pattern.isAllConstants(), "Can't add an entry that isn't all constants"); - CUGRAPH_EXPECTS(static_cast(pattern.getSize()) == columns.size(), - "Can't add an entry that isn't the right size"); - inputBuffer.push_back(pattern); -} - -template -void db_table::rebuildIndices() -{ - for (size_t i = 0; i < columns.size(); i++) { - // Copy the column's data to a new array - idx_t size = column_size; - rmm::device_buffer tempColumn(sizeof(idx_t) * size); - cudaMemcpy(tempColumn.data(), columns[i].data(), sizeof(idx_t) * size, cudaMemcpyDefault); - - // Construct an array of ascending integers - rmm::device_buffer indirection(sizeof(idx_t) * size); - thrust::sequence(rmm::exec_policy(nullptr)->on(nullptr), - reinterpret_cast(indirection.data()), - reinterpret_cast(indirection.data()) + size); - - // Sort the arrays together - thrust::sort_by_key(rmm::exec_policy(nullptr)->on(nullptr), - reinterpret_cast(tempColumn.data()), - reinterpret_cast(tempColumn.data()) + size, - reinterpret_cast(indirection.data())); - - // Compute offsets array based on sorted column - idx_t maxId; - CUDA_TRY(cudaMemcpy(&maxId, - reinterpret_cast(tempColumn.data()) + size - 1, - sizeof(idx_t), - cudaMemcpyDefault)); - rmm::device_buffer offsets(sizeof(idx_t) * (maxId + 2)); - thrust::lower_bound(rmm::exec_policy(nullptr)->on(nullptr), - reinterpret_cast(tempColumn.data()), - reinterpret_cast(tempColumn.data()) + size, - thrust::counting_iterator(0), - thrust::counting_iterator(maxId + 2), - reinterpret_cast(offsets.data())); - - // Assign new offsets array and indirection vector to index - indices[i].resetData(std::move(offsets), std::move(indirection)); - } -} - -template -void db_table::flush_input() -{ - if (inputBuffer.size() == size_t{0}) return; - idx_t tempSize = inputBuffer.size(); - std::vector> tempColumns(columns.size()); - for (size_t i = 0; i < columns.size(); i++) { - tempColumns[i].resize(tempSize); - for (idx_t j = 0; j < tempSize; j++) { - tempColumns[i][j] = inputBuffer[j].getEntry(i).getConstant(); - } - } - inputBuffer.clear(); - idx_t currentSize = column_size; - idx_t newSize = currentSize + tempSize; - std::vector newColumns; - for (size_t i = 0; i < columns.size(); i++) { newColumns.emplace_back(sizeof(idx_t) * newSize); } - for (size_t i = 0; i < columns.size(); i++) { - if (currentSize > 0) - CUDA_TRY(cudaMemcpy( - newColumns[i].data(), columns[i].data(), sizeof(idx_t) * currentSize, cudaMemcpyDefault)); - CUDA_TRY(cudaMemcpy(reinterpret_cast(newColumns[i].data()) + currentSize, - tempColumns[i].data(), - sizeof(idx_t) * tempSize, - cudaMemcpyDefault)); - columns[i] = std::move(newColumns[i]); - column_size = newSize; - } - - rebuildIndices(); -} - -template -std::string db_table::toString() -{ - idx_t columnSize = 0; - if (columns.size() > 0) columnSize = column_size; - std::stringstream ss; - ss << "Table with " << columns.size() << " columns of length " << columnSize << "\n"; - for (size_t i = 0; i < names.size(); i++) ss << names[i] << " "; - ss << "\n"; - std::vector> hostColumns; - hostColumns.resize(columns.size()); - for (size_t i = 0; i < columns.size(); i++) { - hostColumns[i].resize(columnSize); - CUDA_TRY(cudaMemcpy( - hostColumns[i].data(), columns[i].data(), sizeof(idx_t) * columnSize, cudaMemcpyDefault)); - } - for (idx_t i = 0; i < columnSize; i++) { - for (size_t j = 0; j < hostColumns.size(); j++) ss << hostColumns[j][i] << " "; - ss << "\n"; - } - return ss.str(); -} - -template -db_column_index& db_table::getIndex(int idx) -{ - return indices[idx]; -} - -template -idx_t* db_table::getColumn(int idx) -{ - return reinterpret_cast(columns[idx].data()); -} - -template class db_table; -template class db_table; - -template -db_object::db_object() -{ - next_id = 0; - relationshipsTable.addColumn("begin"); - relationshipsTable.addColumn("end"); - relationshipsTable.addColumn("type"); - relationshipPropertiesTable.addColumn("id"); - relationshipPropertiesTable.addColumn("name"); - relationshipPropertiesTable.addColumn("value"); -} - -template -std::string db_object::query(std::string query) -{ - return ""; -} - -template class db_object; -template class db_object; -} // namespace db -} // namespace cugraph diff --git a/cpp/src/db/db_object.cuh b/cpp/src/db/db_object.cuh deleted file mode 100644 index a9b1f461f85..00000000000 --- a/cpp/src/db/db_object.cuh +++ /dev/null @@ -1,192 +0,0 @@ -/* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#pragma once - -#include -#include -#include "rmm/device_buffer.hpp" -#include "utilities/graph_utils.cuh" - -namespace cugraph { -namespace db { -/** - * Class for representing an entry in a pattern, which may either be a variable or constant value - * See description of db_pattern for more info on how this is used. - */ -template -class db_pattern_entry { - bool is_var; - idx_t constantValue; - std::string variableName; - - public: - db_pattern_entry(std::string variable); - db_pattern_entry(idx_t constant); - db_pattern_entry(const db_pattern_entry& other); - db_pattern_entry& operator=(const db_pattern_entry& other); - bool isVariable() const; - idx_t getConstant() const; - std::string getVariable() const; -}; - -/** - * Class for representing a pattern (usually a triple pattern, but it's extensible) - * A pattern in this sense consists of a sequence of entries each element is either a constant - * value (an integer, since we dictionary encode everything) or a variable. Variables stand - * in for unknown values that are being searched for. For example: if we have a pattern like - * {'a', :haslabel, Person} (Where :haslabel and Person are dictionary encoded constants and - * 'a' is a variable) We are looking for all nodes that have the label Person. - */ -template -class db_pattern { - std::vector> entries; - - public: - db_pattern(); - db_pattern(const db_pattern& other); - db_pattern& operator=(const db_pattern& other); - int getSize() const; - const db_pattern_entry& getEntry(int position) const; - void addEntry(db_pattern_entry& entry); - bool isAllConstants(); -}; - -/** - * Class which encapsulates a CSR-style index on a column - */ -template -class db_column_index { - rmm::device_buffer offsets; - rmm::device_buffer indirection; - - public: - db_column_index() = default; - db_column_index(rmm::device_buffer&& off, rmm::device_buffer&& ind); - db_column_index(const db_column_index& other) = delete; - db_column_index(db_column_index&& other) = default; - ~db_column_index() = default; - db_column_index& operator=(const db_column_index& other) = delete; - db_column_index& operator=(db_column_index&& other) = default; - void resetData(rmm::device_buffer&& offsets, rmm::device_buffer&& indirection); - idx_t* getOffsets(); - idx_t getOffsetsSize(); - idx_t* getIndirection(); - idx_t getIndirectionSize(); - - /** - * For debugging purposes only. - * @return Human readable representation - */ - std::string toString(); -}; - -/** - * Class which encapsulates a result set binding - */ -template -class db_result { - std::vector columns; - std::vector names; - bool dataValid; - idx_t columnSize; - - public: - db_result(); - db_result(db_result&& other); - db_result(db_result& other) = delete; - db_result(const db_result& other) = delete; - ~db_result() = default; - db_result& operator =(db_result&& other); - db_result& operator=(db_result& other) = delete; - db_result& operator=(const db_result& other) = delete; - idx_t getSize(); - idx_t* getData(std::string idx); - void addColumn(std::string columnName); - void allocateColumns(idx_t size); - /** - * For debugging purposes - * @return Human readable representation - */ - std::string toString(); -}; - -/** - * Class which glues an arbitrary number of columns together to form a table - */ -template -class db_table { - std::vector columns; - idx_t column_size; - std::vector names; - std::vector> inputBuffer; - std::vector> indices; - - public: - db_table(); - ~db_table() = default; - void addColumn(std::string name); - void addEntry(db_pattern& pattern); - - /** - * This method will rebuild the indices for each column in the table. This is done by - * sorting a copy of the column along with an array which is a 0..n sequence, where - * n is the number of entries in the column. The sorted column is used to produce an - * offsets array and the sequence array becomes a permutation which maps the offset - * position into the original table. - */ - void rebuildIndices(); - - /** - * This method takes all the temporary input in the input buffer and appends it onto - * the existing table. - */ - void flush_input(); - - /** - * This method is for debugging purposes. It returns a human readable string representation - * of the table. - * @return Human readable string representation - */ - std::string toString(); - db_column_index& getIndex(int idx); - idx_t* getColumn(int idx); - idx_t getColumnSize(); -}; - -/** - * The main database object. It stores the needed tables and provides a method hook to run - * a query on the data. - */ -template -class db_object { - // The dictionary and reverse dictionary encoding strings to ids and vice versa - std::map valueToId; - std::map idToValue; - idx_t next_id; - - // The relationship table - db_table relationshipsTable; - - // The relationship property table - db_table relationshipPropertiesTable; - - public: - db_object(); - std::string query(std::string query); -}; -} // namespace db -} // namespace cugraph diff --git a/cpp/src/db/db_operators.cu b/cpp/src/db/db_operators.cu deleted file mode 100644 index d67f7ef9140..00000000000 --- a/cpp/src/db/db_operators.cu +++ /dev/null @@ -1,394 +0,0 @@ -/* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include - -#include - -#include - -#include - -namespace cugraph { -namespace db { -template -struct degree_iterator { - IndexType* offsets; - degree_iterator(IndexType* _offsets) : offsets(_offsets) {} - - __host__ __device__ IndexType operator[](IndexType place) - { - return offsets[place + 1] - offsets[place]; - } -}; - -template -struct deref_functor { - It iterator; - deref_functor(It it) : iterator(it) {} - - __host__ __device__ IndexType operator()(IndexType in) { return iterator[in]; } -}; - -template -struct notNegativeOne { - __host__ __device__ flag_t operator()(idx_t in) { return in != -1; } -}; - -template -__device__ IndexType -binsearch_maxle(const IndexType* vec, const IndexType val, IndexType low, IndexType high) -{ - while (true) { - if (low == high) return low; // we know it exists - if ((low + 1) == high) return (vec[high] <= val) ? high : low; - - IndexType mid = low + (high - low) / 2; - - if (vec[mid] > val) - high = mid - 1; - else - low = mid; - } -} - -template -__global__ void compute_bucket_offsets_kernel(const IndexType* frontier_degrees_exclusive_sum, - IndexType* bucket_offsets, - const IndexType frontier_size, - IndexType total_degree) -{ - IndexType end = ((total_degree - 1 + FIND_MATCHES_BLOCK_SIZE) / FIND_MATCHES_BLOCK_SIZE); - - for (IndexType bid = blockIdx.x * blockDim.x + threadIdx.x; bid <= end; - bid += gridDim.x * blockDim.x) { - IndexType eid = min(bid * FIND_MATCHES_BLOCK_SIZE, total_degree - 1); - - bucket_offsets[bid] = - binsearch_maxle(frontier_degrees_exclusive_sum, eid, (IndexType)0, frontier_size - 1); - } -} - -template -__global__ void findMatchesKernel(idx_t inputSize, - idx_t outputSize, - idx_t maxBlock, - idx_t* offsets, - idx_t* indirection, - idx_t* blockStarts, - idx_t* expandCounts, - idx_t* frontier, - idx_t* columnA, - idx_t* columnB, - idx_t* columnC, - idx_t* outputA, - idx_t* outputB, - idx_t* outputC, - idx_t* outputD, - idx_t patternA, - idx_t patternB, - idx_t patternC) -{ - __shared__ idx_t blockRange[2]; - __shared__ idx_t localExSum[FIND_MATCHES_BLOCK_SIZE * 2]; - __shared__ idx_t localFrontier[FIND_MATCHES_BLOCK_SIZE * 2]; - - for (idx_t bid = blockIdx.x; bid < maxBlock; bid += gridDim.x) { - // Copy in the block's section of the expand counts - if (threadIdx.x == 0) { - blockRange[0] = blockStarts[bid]; - blockRange[1] = blockStarts[bid + 1]; - if (blockRange[0] > 0) { blockRange[0] -= 1; } - } - __syncthreads(); - - idx_t sectionSize = blockRange[1] - blockRange[0]; - for (int tid = threadIdx.x; tid <= sectionSize; tid += blockDim.x) { - localExSum[tid] = expandCounts[blockRange[0] + tid]; - localFrontier[tid] = frontier[blockRange[0] + tid]; - } - __syncthreads(); - - // Do the work item for each thread of this virtual block: - idx_t tid = bid * blockDim.x + threadIdx.x; - if (tid < outputSize) { - // Figure out which row this thread/iteration is working on - idx_t sourceIdx = binsearch_maxle(localExSum, tid, (idx_t)0, (idx_t)sectionSize); - idx_t source = localFrontier[sourceIdx]; - idx_t rank = tid - localExSum[sourceIdx]; - idx_t row_id = indirection[offsets[source] + rank]; - - // Load in values from the row for A, B, and C columns - idx_t valA = columnA[row_id]; - idx_t valB = columnB[row_id]; - idx_t valC = columnC[row_id]; - - // Compare the row values with constants in the pattern - bool matchA = outputA != nullptr ? true : patternA == valA; - bool matchB = outputB != nullptr ? true : patternB == valB; - bool matchC = outputC != nullptr ? true : patternC == valC; - - // If row doesn't match, set row values to -1 before writing out - if (!(matchA && matchB && matchC)) { - valA = -1; - valB = -1; - valC = -1; - row_id = -1; - } - - // Write out values to non-null outputs - if (outputA != nullptr) outputA[tid] = valA; - if (outputB != nullptr) outputB[tid] = valB; - if (outputC != nullptr) outputC[tid] = valC; - if (outputD != nullptr) outputD[tid] = row_id; - } - } -} - -template -db_result findMatches(db_pattern& pattern, - db_table& table, - idx_t* frontier, - idx_t frontier_size, - int indexPosition) -{ - // Find out if the indexPosition is a variable or constant - bool indexConstant = !pattern.getEntry(indexPosition).isVariable(); - - db_column_index& theIndex = table.getIndex(indexPosition); - - // Check to see whether we are going to be saving out the row ids from matches - bool saveRowIds = false; - if (pattern.getSize() == 4) saveRowIds = true; - - // Check if we have a frontier to use, if we don't make one up - bool givenInputFrontier = frontier != nullptr; - idx_t frontierSize; - idx_t* frontier_ptr = nullptr; - rmm::device_buffer frontierBuffer; - if (givenInputFrontier) { - frontier_ptr = frontier; - frontierSize = frontier_size; - } else { - if (indexConstant) { - // Use a single value equal to the constant in the pattern - idx_t constantValue = pattern.getEntry(indexPosition).getConstant(); - frontierBuffer.resize(sizeof(idx_t)); - thrust::fill(rmm::exec_policy(nullptr)->on(nullptr), - reinterpret_cast(frontierBuffer.data()), - reinterpret_cast(frontierBuffer.data()) + 1, - constantValue); - frontier_ptr = reinterpret_cast(frontierBuffer.data()); - frontierSize = 1; - } else { - // Making a sequence of values from zero to n where n is the highest ID present in the index. - idx_t highestId = theIndex.getOffsetsSize() - 2; - frontierBuffer.resize(sizeof(idx_t) * (highestId + 1)); - thrust::sequence(rmm::exec_policy(nullptr)->on(nullptr), - reinterpret_cast(frontierBuffer.data()), - reinterpret_cast(frontierBuffer.data()) + highestId + 1); - frontier_ptr = reinterpret_cast(frontierBuffer.data()); - frontierSize = highestId + 1; - } - } - - // Collect all the pointers needed to run the main kernel - idx_t* columnA = table.getColumn(0); - idx_t* columnB = table.getColumn(1); - idx_t* columnC = table.getColumn(2); - idx_t* offsets = theIndex.getOffsets(); - idx_t* indirection = theIndex.getIndirection(); - - // Load balance the input - rmm::device_buffer exsum_degree(sizeof(idx_t) * (frontierSize + 1)); - degree_iterator deg_it(offsets); - deref_functor, idx_t> deref(deg_it); - thrust::fill(rmm::exec_policy(nullptr)->on(nullptr), - reinterpret_cast(exsum_degree.data()), - reinterpret_cast(exsum_degree.data()) + 1, - 0); - thrust::transform(rmm::exec_policy(nullptr)->on(nullptr), - frontier_ptr, - frontier_ptr + frontierSize, - reinterpret_cast(exsum_degree.data()) + 1, - deref); - thrust::inclusive_scan(rmm::exec_policy(nullptr)->on(nullptr), - reinterpret_cast(exsum_degree.data()) + 1, - reinterpret_cast(exsum_degree.data()) + frontierSize + 1, - reinterpret_cast(exsum_degree.data()) + 1); - idx_t output_size; - CUDA_TRY(cudaMemcpy(&output_size, - reinterpret_cast(exsum_degree.data()) + frontierSize, - sizeof(idx_t), - cudaMemcpyDefault)); - - idx_t num_blocks = (output_size + FIND_MATCHES_BLOCK_SIZE - 1) / FIND_MATCHES_BLOCK_SIZE; - rmm::device_buffer block_bucket_offsets(sizeof(idx_t) * (num_blocks + 1)); - - dim3 grid, block; - block.x = 512; - grid.x = min((idx_t)MAXBLOCKS, (num_blocks / 512) + 1); - compute_bucket_offsets_kernel<<>>( - reinterpret_cast(exsum_degree.data()), - reinterpret_cast(block_bucket_offsets.data()), - frontierSize, - output_size); - - // Allocate space for the result - idx_t* outputA = nullptr; - idx_t* outputB = nullptr; - idx_t* outputC = nullptr; - idx_t* outputD = nullptr; - rmm::device_buffer outputABuffer; - rmm::device_buffer outputBBuffer; - rmm::device_buffer outputCBuffer; - rmm::device_buffer outputDBuffer; - if (pattern.getEntry(0).isVariable()) { - outputABuffer.resize(sizeof(idx_t) * output_size); - outputA = reinterpret_cast(outputABuffer.data()); - } - if (pattern.getEntry(1).isVariable()) { - outputBBuffer.resize(sizeof(idx_t) * output_size); - outputB = reinterpret_cast(outputBBuffer.data()); - } - if (pattern.getEntry(2).isVariable()) { - outputCBuffer.resize(sizeof(idx_t) * output_size); - outputC = reinterpret_cast(outputCBuffer.data()); - } - if (saveRowIds) { - outputDBuffer.resize(sizeof(idx_t) * output_size); - outputD = reinterpret_cast(outputDBuffer.data()); - } - - // Get the constant pattern entries from the pattern to pass into the main kernel - idx_t patternA = -1; - idx_t patternB = -1; - idx_t patternC = -1; - if (!pattern.getEntry(0).isVariable()) { patternA = pattern.getEntry(0).getConstant(); } - if (!pattern.getEntry(1).isVariable()) { patternB = pattern.getEntry(1).getConstant(); } - if (!pattern.getEntry(2).isVariable()) { patternC = pattern.getEntry(2).getConstant(); } - - // Call the main kernel - block.x = FIND_MATCHES_BLOCK_SIZE; - grid.x = min((idx_t)MAXBLOCKS, - (output_size + (idx_t)FIND_MATCHES_BLOCK_SIZE - 1) / (idx_t)FIND_MATCHES_BLOCK_SIZE); - findMatchesKernel<<>>( - frontierSize, - output_size, - num_blocks, - offsets, - indirection, - reinterpret_cast(block_bucket_offsets.data()), - reinterpret_cast(exsum_degree.data()), - frontier_ptr, - columnA, - columnB, - columnC, - outputA, - outputB, - outputC, - outputD, - patternA, - patternB, - patternC); - - // Get the non-null output columns - std::vector columns; - std::vector names; - if (outputA != nullptr) { - columns.push_back(outputA); - names.push_back(pattern.getEntry(0).getVariable()); - } - if (outputB != nullptr) { - columns.push_back(outputB); - names.push_back(pattern.getEntry(1).getVariable()); - } - if (outputC != nullptr) { - columns.push_back(outputC); - names.push_back(pattern.getEntry(2).getVariable()); - } - if (outputD != nullptr) { - columns.push_back(outputD); - names.push_back(pattern.getEntry(3).getVariable()); - } - - // Remove non-matches from result - rmm::device_buffer flags(sizeof(int8_t) * output_size); - - idx_t* col_ptr = columns[0]; - thrust::transform(rmm::exec_policy(nullptr)->on(nullptr), - col_ptr, - col_ptr + output_size, - reinterpret_cast(flags.data()), - notNegativeOne()); - - size_t tempSpaceSize = 0; - rmm::device_buffer compactSize_d(sizeof(idx_t)); - cub::DeviceSelect::Flagged(nullptr, - tempSpaceSize, - col_ptr, - reinterpret_cast(flags.data()), - col_ptr, - reinterpret_cast(compactSize_d.data()), - output_size); - rmm::device_buffer tempSpace(tempSpaceSize); - cub::DeviceSelect::Flagged(tempSpace.data(), - tempSpaceSize, - col_ptr, - reinterpret_cast(flags.data()), - col_ptr, - reinterpret_cast(compactSize_d.data()), - output_size); - idx_t compactSize_h; - cudaMemcpy(&compactSize_h, compactSize_d.data(), sizeof(idx_t), cudaMemcpyDefault); - - for (size_t i = 1; i < columns.size(); i++) { - col_ptr = columns[i]; - cub::DeviceSelect::Flagged(tempSpace.data(), - tempSpaceSize, - col_ptr, - reinterpret_cast(flags.data()), - col_ptr, - reinterpret_cast(compactSize_d.data()), - output_size); - } - - // Put together the result to return - db_result result; - for (size_t i = 0; i < names.size(); i++) { result.addColumn(names[i]); } - result.allocateColumns(compactSize_h); - for (size_t i = 0; i < columns.size(); i++) { - idx_t* outputPtr = result.getData(names[i]); - idx_t* inputPtr = columns[i]; - CUDA_TRY(cudaMemcpy(outputPtr, inputPtr, sizeof(idx_t) * compactSize_h, cudaMemcpyDefault)); - } - - // Return the result - return result; -} - -template db_result findMatches(db_pattern& pattern, - db_table& table, - int32_t* frontier, - int32_t frontier_size, - int indexPosition); -template db_result findMatches(db_pattern& pattern, - db_table& table, - int64_t* frontier, - int64_t frontier_size, - int indexPosition); -} // namespace db -} // namespace cugraph diff --git a/cpp/src/db/db_operators.cuh b/cpp/src/db/db_operators.cuh deleted file mode 100644 index 6a2e8322069..00000000000 --- a/cpp/src/db/db_operators.cuh +++ /dev/null @@ -1,46 +0,0 @@ -/* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#pragma once - -#include - -#define MAXBLOCKS 65535 -#define FIND_MATCHES_BLOCK_SIZE 512 - -namespace cugraph { -namespace db { -/** - * Method to find matches to a pattern against an indexed table. - * @param pattern The pattern to match against. It is assumed that the order of the entries - * matches the order of the columns in the table being searched. - * @param table The table to find matching entries within. - * @param frontier The frontier of already bound values. The search is restricted to entries in the - * table which match at least the frontier entry. If the frontier is null, then the entire table - * will be scanned. - * @param indexColumn The name of the variable in the pattern which is bound to the frontier - * and which indicates which index should be used on the table. - * @return A result table with columns for each variable in the given pattern containing the bound - * values to those variables. - */ -template -db_result findMatches(db_pattern& pattern, - db_table& table, - idx_t* frontier, - idx_t frontier_size, - int indexPosition); -} // namespace db -} // namespace cugraph diff --git a/cpp/src/db/db_parser_integration_test.cu b/cpp/src/db/db_parser_integration_test.cu deleted file mode 100644 index aa395bf8a4c..00000000000 --- a/cpp/src/db/db_parser_integration_test.cu +++ /dev/null @@ -1,27 +0,0 @@ -/* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include - -namespace cugraph { -namespace db { -std::string getParserVersion() -{ - std::string version = libcypher_parser_version(); - return version; -} -} // namespace db -} // namespace cugraph \ No newline at end of file diff --git a/cpp/src/db/db_parser_integration_test.cuh b/cpp/src/db/db_parser_integration_test.cuh deleted file mode 100644 index 63da8805164..00000000000 --- a/cpp/src/db/db_parser_integration_test.cuh +++ /dev/null @@ -1,24 +0,0 @@ -/* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include -#include - -namespace cugraph { -namespace db { -std::string getParserVersion(); -} -} // namespace cugraph diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index a8c789210e0..f26fe9c022f 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -245,13 +245,6 @@ set(SCC_TEST_SRC ConfigureTest(SCC_TEST "${SCC_TEST_SRC}" "") -################################################################################################### -# - FIND_MATCHES tests ---------------------------------------------------------------------------- - -set(FIND_MATCHES_TEST_SRC - "${CMAKE_CURRENT_SOURCE_DIR}/db/find_matches_test.cu") - -ConfigureTest(FIND_MATCHES_TEST "${FIND_MATCHES_TEST_SRC}" "") ################################################################################################### # - Experimental Graph tests ---------------------------------------------------------------------- diff --git a/cpp/tests/db/find_matches_test.cu b/cpp/tests/db/find_matches_test.cu deleted file mode 100644 index c1373bb8bf2..00000000000 --- a/cpp/tests/db/find_matches_test.cu +++ /dev/null @@ -1,232 +0,0 @@ -/* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#include -#include -#include -#include - -#include - -#include - -class Test_FindMatches : public ::testing::Test { - public: - Test_FindMatches() {} - virtual void SetUp() - { - cugraph::db::db_pattern p; - cugraph::db::db_pattern_entry p1(0); - cugraph::db::db_pattern_entry p2(1); - cugraph::db::db_pattern_entry p3(2); - p.addEntry(p1); - p.addEntry(p2); - p.addEntry(p3); - table.addColumn("Start"); - table.addColumn("EdgeType"); - table.addColumn("End"); - table.addEntry(p); - table.flush_input(); - } - virtual void TearDown() {} - void insertConstantEntry(int32_t a, int32_t b, int32_t c) - { - cugraph::db::db_pattern p; - cugraph::db::db_pattern_entry p1(a); - cugraph::db::db_pattern_entry p2(b); - cugraph::db::db_pattern_entry p3(c); - p.addEntry(p1); - p.addEntry(p2); - p.addEntry(p3); - table.addEntry(p); - } - cugraph::db::db_table table; -}; - -TEST_F(Test_FindMatches, verifyIndices) -{ - insertConstantEntry(0, 1, 1); - insertConstantEntry(2, 0, 1); - table.flush_input(); - - std::cout << table.toString(); - std::cout << "Index[0]: " << table.getIndex(0).toString(); - std::cout << "Index[1]: " << table.getIndex(1).toString(); - std::cout << "Index[2]: " << table.getIndex(2).toString(); -} - -TEST_F(Test_FindMatches, firstTest) -{ - cugraph::db::db_pattern p; - cugraph::db::db_pattern_entry p1(0); - cugraph::db::db_pattern_entry p2("a"); - cugraph::db::db_pattern_entry p3("b"); - p.addEntry(p1); - p.addEntry(p2); - p.addEntry(p3); - cugraph::db::db_result result = - cugraph::db::findMatches(p, table, nullptr, 0, 1); - ASSERT_EQ(result.getSize(), 1); - std::vector resultA(result.getSize()); - std::vector resultB(result.getSize()); - CUDA_TRY(cudaMemcpy( - resultA.data(), result.getData("a"), sizeof(int32_t) * result.getSize(), cudaMemcpyDefault)); - CUDA_TRY(cudaMemcpy( - resultB.data(), result.getData("b"), sizeof(int32_t) * result.getSize(), cudaMemcpyDefault)); - ASSERT_EQ(resultA[0], 1); - ASSERT_EQ(resultB[0], 2); -} - -TEST_F(Test_FindMatches, secondTest) -{ - insertConstantEntry(0, 1, 1); - insertConstantEntry(2, 0, 1); - table.flush_input(); - - std::cout << table.toString() << "\n\n"; - - std::cout << table.getIndex(2).toString() << "\n"; - - cugraph::db::db_pattern q; - cugraph::db::db_pattern_entry q1(0); - cugraph::db::db_pattern_entry q2("a"); - cugraph::db::db_pattern_entry q3("b"); - q.addEntry(q1); - q.addEntry(q2); - q.addEntry(q3); - - cugraph::db::db_result result = - cugraph::db::findMatches(q, table, nullptr, 0, 2); - - std::cout << result.toString(); - - ASSERT_EQ(result.getSize(), 2); - std::vector resultA(result.getSize()); - std::vector resultB(result.getSize()); - CUDA_TRY(cudaMemcpy( - resultA.data(), result.getData("a"), sizeof(int32_t) * result.getSize(), cudaMemcpyDefault)); - CUDA_TRY(cudaMemcpy( - resultB.data(), result.getData("b"), sizeof(int32_t) * result.getSize(), cudaMemcpyDefault)); - - ASSERT_EQ(resultA[0], 1); - ASSERT_EQ(resultB[0], 1); - ASSERT_EQ(resultA[1], 1); - ASSERT_EQ(resultB[1], 2); -} - -TEST_F(Test_FindMatches, thirdTest) -{ - insertConstantEntry(1, 1, 2); - insertConstantEntry(2, 1, 2); - table.flush_input(); - - cugraph::db::db_pattern q; - cugraph::db::db_pattern_entry q1("a"); - cugraph::db::db_pattern_entry q2(1); - cugraph::db::db_pattern_entry q3(2); - q.addEntry(q1); - q.addEntry(q2); - q.addEntry(q3); - - rmm::device_buffer frontier(sizeof(int32_t)); - int32_t* frontier_ptr = reinterpret_cast(frontier.data()); - thrust::fill(thrust::device, frontier_ptr, frontier_ptr + 1, 0); - - cugraph::db::db_result result = - cugraph::db::findMatches(q, table, frontier_ptr, 1, 0); - - ASSERT_EQ(result.getSize(), 1); - std::vector resultA(result.getSize()); - CUDA_TRY(cudaMemcpy( - resultA.data(), result.getData("a"), sizeof(int32_t) * result.getSize(), cudaMemcpyDefault)); - - std::cout << result.toString(); - - ASSERT_EQ(resultA[0], 0); -} - -TEST_F(Test_FindMatches, fourthTest) -{ - insertConstantEntry(1, 1, 2); - insertConstantEntry(2, 1, 2); - table.flush_input(); - - cugraph::db::db_pattern q; - cugraph::db::db_pattern_entry q1("a"); - cugraph::db::db_pattern_entry q2(1); - cugraph::db::db_pattern_entry q3(2); - cugraph::db::db_pattern_entry q4("r"); - q.addEntry(q1); - q.addEntry(q2); - q.addEntry(q3); - q.addEntry(q4); - - cugraph::db::db_result result = - cugraph::db::findMatches(q, table, nullptr, 0, 0); - std::cout << result.toString(); - ASSERT_EQ(result.getSize(), 3); - - std::vector resultA(result.getSize()); - std::vector resultR(result.getSize()); - - CUDA_TRY(cudaMemcpy( - resultA.data(), result.getData("a"), sizeof(int32_t) * result.getSize(), cudaMemcpyDefault)); - CUDA_TRY(cudaMemcpy( - resultR.data(), result.getData("r"), sizeof(int32_t) * result.getSize(), cudaMemcpyDefault)); - - ASSERT_EQ(resultA[0], 0); - ASSERT_EQ(resultA[1], 1); - ASSERT_EQ(resultA[2], 2); - ASSERT_EQ(resultR[0], 0); - ASSERT_EQ(resultR[1], 1); - ASSERT_EQ(resultR[2], 2); -} - -TEST_F(Test_FindMatches, fifthTest) -{ - insertConstantEntry(0, 1, 3); - insertConstantEntry(0, 2, 1); - insertConstantEntry(0, 2, 2); - table.flush_input(); - - cugraph::db::db_pattern q; - cugraph::db::db_pattern_entry q1("a"); - cugraph::db::db_pattern_entry q2(1); - cugraph::db::db_pattern_entry q3("b"); - q.addEntry(q1); - q.addEntry(q2); - q.addEntry(q3); - - cugraph::db::db_result result = - cugraph::db::findMatches(q, table, nullptr, 0, 1); - std::cout << result.toString(); - - ASSERT_EQ(result.getSize(), 2); - std::vector resultA(result.getSize()); - std::vector resultB(result.getSize()); - - CUDA_TRY(cudaMemcpy( - resultA.data(), result.getData("a"), sizeof(int32_t) * result.getSize(), cudaMemcpyDefault)); - CUDA_TRY(cudaMemcpy( - resultB.data(), result.getData("b"), sizeof(int32_t) * result.getSize(), cudaMemcpyDefault)); - - ASSERT_EQ(resultA[0], 0); - ASSERT_EQ(resultA[1], 0); - ASSERT_EQ(resultB[0], 2); - ASSERT_EQ(resultB[1], 3); -} - -CUGRAPH_TEST_PROGRAM_MAIN()