Merge pull request #36 from sirius-db/devesh/optimized_result_formatter

Optimized result formatter

Author: sarda-devesh

CMakeLists.txt

@@ -59,6 +59,7 @@ set(EXTENSION_SOURCES
   src/gpu_buffer_manager.cpp
   src/gpu_columns.cpp
   src/gpu_expression_executor.cpp
+  src/gpu_query_result.cpp
 )
 add_subdirectory(src/operator)
 add_subdirectory(src/plan)

docs/README.md

@@ -43,7 +43,7 @@ conda install -c rapidsai -c conda-forge -c nvidia rapidsai::libcudf
 Set the environment variables `USE_CUDF` to 1 and `LIBCUDF_ENV_PREFIX` to the conda environment's path. For example, if we installed miniconda in `~/miniconda3` and installed libcudf in the conda environment `libcudf-env`, then we would set the `LIBCUDF_ENV_PREFIX` to `~/miniconda3/envs/libcudf-env`.
 ```
 export USE_CUDF=1
-export LIBCUDF_ENV_PREFIX = {PATH to libcudf-env}
+export LIBCUDF_ENV_PREFIX={PATH to libcudf-env}
 ```
 
 ### Clone the Sirius repository

queries/devesh_test_queries/q10_group_by_test.sql

@@ -153,7 +153,7 @@ __global__ void create_cpu_strings(duckdb_string_type* gpu_strings, char* cpu_ch
     }
 }
 
-void materializeStringColumnToDuckdbFormat(GPUColumn* column, char* column_char_write_buffer, string_t* column_string_write_buffer) {
+void materializeStringColumnToDuckdbFormat(shared_ptr<GPUColumn> column, char* column_char_write_buffer, string_t* column_string_write_buffer) {
     // First copy the characters from the GPU to the CPU
     SIRIUS_LOG_DEBUG("Materialize String Column to Duckdb format");
     SETUP_TIMING();

queries/devesh_test_queries/string_group_and_order_by.txt

@@ -0,0 +1,73 @@
+#include "gpu_query_result.hpp"
+#include "duckdb/common/to_string.hpp"
+#include "duckdb/main/client_context.hpp"
+#include "duckdb/common/box_renderer.hpp"
+
+namespace duckdb {
+
+void GPUResultCollection::SetCapacity(size_t capacity) { 
+    data_chunks = new DataChunk[capacity];
+}
+
+void GPUResultCollection::AddChunk(DataChunk& chunk) {
+    num_rows += chunk.size();
+    data_chunks[write_idx].Move(chunk);
+    write_idx += 1;
+}
+
+unique_ptr<DataChunk> GPUResultCollection::GetNext() {
+    // We have returned all of the values then return the empty result
+    if(read_idx >= write_idx) {
+        return nullptr;
+    }
+
+    // Create a result that references the value in the buffer
+    DataChunk& return_chunk = data_chunks[read_idx];
+    unique_ptr<DataChunk> result_value = make_uniq<DataChunk>();
+
+    result_value->InitializeEmpty(return_chunk.GetTypes());
+    result_value->SetCardinality(return_chunk.size());
+    for(int col = 0; col < return_chunk.data.size(); col++) {
+        result_value->data[col].Reference(return_chunk.data[col]);
+    }
+    read_idx += 1;
+    num_rows -= result_value->size();
+
+    return result_value;
+}
+
+GPUQueryResult::GPUQueryResult(StatementType statement_type, StatementProperties properties,
+    vector<string> names, vector<LogicalType> types, ClientProperties client_properties, unique_ptr<GPUResultCollection> result_collection)
+    : QueryResult(QueryResultType::MATERIALIZED_RESULT, statement_type, std::move(properties), types,
+std::move(names), std::move(client_properties)), result_collection(std::move(result_collection)) {
+
+}
+
+string GPUQueryResult::ToString() {
+    std::string result("GPUQueryResult");
+    return result;
+}
+
+string GPUQueryResult::ToBox(ClientContext &context, const BoxRendererConfig &config) {
+    std::string result("BoxedGPUQueryResult");
+    return result;
+}
+
+idx_t GPUQueryResult::RowCount() const {
+    idx_t row_count = (idx_t) result_collection->size();
+    return row_count;
+}
+
+unique_ptr<DataChunk> GPUQueryResult::Fetch() {
+	return FetchRaw();
+}
+
+Value GPUQueryResult::GetValue(idx_t column, idx_t index) {
+    throw InternalException("GetValue not implemented for GPUQueryResult");
+}
+
+unique_ptr<DataChunk> GPUQueryResult::FetchRaw() {
+    return result_collection->GetNext();
+}
+
+}

src/cuda/operator/materialize.cu

@@ -0,0 +1,64 @@
+#pragma once
+
+#include "helper/common.h"
+#include "utils.hpp"
+#include "duckdb/common/winapi.hpp"
+#include "duckdb/main/query_result.hpp"
+
+namespace duckdb {
+
+class GPUResultCollection { 
+public: 
+	GPUResultCollection() : read_idx(0), write_idx(0), num_rows(0), data_chunks(nullptr) {
+
+	}
+
+	void SetCapacity(size_t capacity);
+	void AddChunk(DataChunk& chunk);
+	unique_ptr<DataChunk> GetNext();
+
+	size_t size() {
+		return num_rows;
+	}
+
+	~GPUResultCollection() {
+		if(data_chunks != nullptr) {
+			delete[] data_chunks;
+		}
+	}
+
+	DataChunk* data_chunks;
+	size_t num_rows;
+	size_t write_idx;
+	size_t read_idx;
+};
+
+// The reason we need to implement our own QueryResult is that duckdb's MaterializedQueryResult
+// in constructed by passing in a ColumnDataCollection which makes a copy of the data when we try
+// to Append DataChunks to it. By implement our own basic QueryResult we can bypass this unncessary
+// copy
+class GPUQueryResult : public QueryResult {
+public:
+    DUCKDB_API GPUQueryResult(StatementType statement_type, StatementProperties properties, vector<string> names, 
+		vector<LogicalType> types, ClientProperties client_properties, unique_ptr<GPUResultCollection> result_collection);
+    
+        //! Fetches a DataChunk from the query result.
+	//! This will consume the result (i.e. the result can only be scanned once with this function)
+	DUCKDB_API unique_ptr<DataChunk> Fetch() override;
+	DUCKDB_API unique_ptr<DataChunk> FetchRaw() override;
+	//! Converts the QueryResult to a string
+	DUCKDB_API string ToString() override;
+	DUCKDB_API string ToBox(ClientContext &context, const BoxRendererConfig &config) override;
+
+	//! Gets the (index) value of the (column index) column.
+	//! Note: this is very slow. Scanning over the underlying collection is much faster.
+	DUCKDB_API Value GetValue(idx_t column, idx_t index);
+
+	DUCKDB_API idx_t RowCount() const;
+
+private:
+    // The actual chunks we want to return
+    unique_ptr<GPUResultCollection> result_collection;
+};
+
+}

src/gpu_query_result.cpp

@@ -3,6 +3,7 @@
 #include "gpu_physical_operator.hpp"
 #include "duckdb/common/enums/statement_type.hpp"
 #include "gpu_buffer_manager.hpp"
+#include "gpu_query_result.hpp"
 #include "duckdb/common/types/column/column_data_collection.hpp"
 
 namespace duckdb {
@@ -47,7 +48,7 @@ class GPUPhysicalResultCollector : public GPUPhysicalOperator {
 class GPUPhysicalMaterializedCollector : public GPUPhysicalResultCollector {
 public:
 	GPUPhysicalMaterializedCollector(GPUPreparedStatementData &data);
-	unique_ptr<ColumnDataCollection> collection;
+	unique_ptr<GPUResultCollection> result_collection;
 	// ColumnDataAppendState append_state;
 	// bool parallel;