[WIP] Initial commit for cudf integration.

include/xgboost/data.h

@@ -16,10 +16,11 @@
 #include <algorithm>
 #include <string>
 #include <vector>
+
 #include "./base.h"
+
 #include "../../src/common/span.h"
 #include "../../src/common/group_data.h"
-
 #include "../../src/common/host_device_vector.h"
 
 namespace xgboost {
@@ -34,6 +35,17 @@ enum DataType {
   kUInt64 = 4
 };
 
+typedef unsigned char foreign_valid_type;
+typedef int foreign_size_type;
+
+struct ForeignColumn {
+  void * data;
+  foreign_valid_type * valid;
+  foreign_size_type size;
+  foreign_size_type num_nonzero;
+  foreign_size_type null_count;
+};
+
 /*!
  * \brief Meta information about dataset, always sit in memory.
  */
@@ -122,6 +134,13 @@ class MetaInfo {
    * \param num Number of elements in the source array.
    */
   void SetInfo(const char* key, const void* dptr, DataType dtype, size_t num);
+  /*!
+   * \brief Set information in the meta info for foreign columns buffer.
+   * \param key The key of the information.
+   * \param cols The foreign columns buffer used to set the info.
+   * \param n_cols The number of foreign columns.
+   */
+  void SetInfo(const char * key, ForeignColumn ** cols, foreign_size_type n_cols);
 
  private:
   /*! \brief argsort of labels */
@@ -151,6 +170,14 @@ struct Entry {
   }
 };
 
+struct ForeignCSR {
+  Entry * data;
+  size_t * offsets;
+  size_t n_nonzero;
+  size_t n_cols;
+  size_t n_rows;
+};
+
 /*!
  * \brief In-memory storage unit of sparse batch, stored in CSR format.
  */

python-package/xgboost/core.py

@@ -526,6 +526,33 @@ def _init_from_dt(self, data, nthread):
             nthread))
         self.handle = handle
 
+    def _init_from_columnar(df):
+        '''Initialize DMatrix from columnar memory format.  For now assuming
+        it's cudf.DataFrame.
+
+        '''
+        print('_init_from_columnar')
+        col_ptrs = [ctypes.c_void_p(df[col]._column._pointer) for col in
+                    df.columns]
+        validity_masks = []
+        for col in df.columns:
+            if df[col].has_null_mask:
+                validity_masks.append(df[col].nullmask)
+            else:
+                validity_masks.append(False)
+        col_pairs = list(zip(col_ptrs, validity_masks))
+        interfaces = []
+        for pointers in col_pairs:
+            col = {'data': (pointers[0], False)}
+            if pointers[1] is not False:
+                col['mask'] = pointers[1].mem.__cuda_array_interface__
+            else:
+                col['mask'] = ''
+            interfaces.append(str(col))
+            print(col)
+        interfaces = from_pystr_to_cstr(interfaces)
+        return interfaces
+
     def __del__(self):
         if hasattr(self, "handle") and self.handle is not None:
             _check_call(_LIB.XGDMatrixFree(self.handle))

src/c_api/c_api.cc

@@ -4,6 +4,7 @@
 #include <xgboost/learner.h>
 #include <xgboost/c_api.h>
 #include <xgboost/logging.h>
+#include <xgboost/json.h>
 
 #include <dmlc/thread_local.h>
 #include <rabit/rabit.h>
@@ -189,6 +190,29 @@ int XGDMatrixCreateFromDataIter(
   API_END();
 }
 
+int XGDMatrixCreateFromForeignColumns(ForeignColumn ** cols,
+                                      foreign_size_type n_cols,
+                                      DMatrixHandle * out) {
+  API_BEGIN();
+  std::unique_ptr<data::SimpleCSRSource> source(new data::SimpleCSRSource());
+  source->CopyFrom(cols, n_cols);
+  *out = new std::shared_ptr<DMatrix>(DMatrix::Create(std::move(source)));
+  API_END();
+}
+
+XGB_DLL int XGDMatrixCreateFromArrayInterface(char** c_json_strs, size_t n_columns, DMatrixHandle* out) {
+  API_BEGIN();
+  std::vector<std::string> json_strs;
+  for (size_t i = 0; i < n_columns; ++i) {
+    json_strs.emplace_back(c_json_strs[i]);
+  }
+  std::vector<Json> interfaces(n_columns);
+  for (size_t i = 0; i < n_columns; ++i) {
+    interfaces[i] = Json::Load({json_strs[i].c_str(), json_strs[i].size()});
+  }
+  API_END();
+}
+
 XGB_DLL int XGDMatrixCreateFromCSREx(const size_t* indptr,
                                      const unsigned* indices,
                                      const bst_float* data,
@@ -689,6 +713,17 @@ XGB_DLL int XGDMatrixSetFloatInfo(DMatrixHandle handle,
   API_END();
 }
 
+XGB_DLL int XGDMatrixSetForeignInfo(DMatrixHandle handle,
+                          const char * field,
+                          ForeignColumn ** cols,
+                          foreign_size_type n_cols) {
+  API_BEGIN();
+  CHECK_HANDLE();
+  static_cast<std::shared_ptr<DMatrix>*>(handle)
+      ->get()->Info().SetInfo(field, cols, n_cols);
+  API_END();
+}
+
 XGB_DLL int XGDMatrixSetUIntInfo(DMatrixHandle handle,
                          const char* field,
                          const unsigned* info,

src/common/hist_util.cu

@@ -470,4 +470,4 @@ size_t DeviceSketch
 }
 
 }  // namespace common
-}  // namespace xgboost
+}  // namespace xgboost

src/common/host_device_vector.cu

@@ -623,4 +623,4 @@ template class HostDeviceVector<int>;
 template class HostDeviceVector<Entry>;
 template class HostDeviceVector<size_t>;
 
-}  // namespace xgboost
+}  // namespace xgboost

src/common/host_device_vector.h

@@ -1,7 +1,6 @@
 /*!
  * Copyright 2017-2019 XGBoost contributors
  */
-
 /**
  * @file host_device_vector.h
  * @brief A device-and-host vector abstraction layer.

src/data/data.cc

@@ -146,7 +146,6 @@ void MetaInfo::SetInfo(const char* key, const void* dptr, DataType dtype, size_t
   }
 }
 
-
 DMatrix* DMatrix::Load(const std::string& uri,
                        bool silent,
                        bool load_row_split,

src/data/data.cu

@@ -0,0 +1,67 @@
+/*!
+ * Copyright 2019 by XGBoost Contributors
+ * \file data.cuh
+ * \brief An extension for the data interface to support foreign columnar data buffers
+    This file adds the necessary functions to fill the meta information for the columnar buffers
+ * \author Andrey Adinets
+ * \author Matthew Jones
+ */
+#include <xgboost/data.h>
+#include <xgboost/logging.h>
+
+#include "../common/device_helpers.cuh"
+#include "../common/host_device_vector.h"
+
+namespace xgboost {
+
+__global__ void ReadColumn(ForeignColumn * col,
+                           foreign_size_type n_cols,
+                           float * data) {
+  int tid = threadIdx.x + blockDim.x * blockIdx.x;
+  foreign_size_type n_rows = col->size;
+  if (n_rows <= tid) {
+    return;
+  } else {
+    float * d = (float *) (col->data);
+    data[n_cols * tid] = float(d[tid]);
+ }
+}
+
+void SetInfoFromForeignColumns(MetaInfo * info,
+                               const char * key,
+                               ForeignColumn ** cols,
+                               foreign_size_type n_cols) {
+  CHECK_GT(n_cols, 0);
+  foreign_size_type n_rows = cols[0]->size;
+  for (foreign_size_type i = 0; i < n_cols; ++i) {
+    CHECK_EQ(n_rows, cols[i]->size) << "all foreign columns must be the same size";
+    CHECK_EQ(cols[i]->null_count, 0) << "all labels and weights must be valid";
+  }
+  HostDeviceVector<bst_float> * field = nullptr;
+  if(!strcmp(key, "label")) {
+    field = &info->labels_;
+  } else if (!strcmp(key, "weight")) {
+    CHECK_EQ(n_cols, 1) << "only one foreign column permitted for weights";
+    field = &info->weights_;
+  } else {
+    LOG(WARNING) << key << ": invalid key value for MetaInfo field";
+  }
+
+   GPUSet devices = GPUSet::Range(0, 1);
+  field->Reshard(GPUDistribution::Granular(devices, n_cols));
+  field->Resize(n_cols * n_rows);
+  bst_float * data = field->DevicePointer(0);
+
+  int threads = 1024;
+  int blocks = (n_rows + threads - 1) / threads;
+
+  for (foreign_size_type i = 0; i < n_cols; ++i) {
+    ReadColumn <<<threads, blocks>>> (cols[i], n_cols, data + i);
+    dh::safe_cuda(cudaGetLastError());
+  }
+}
+
+void MetaInfo::SetInfo(const char * key, ForeignColumn ** cols, foreign_size_type n_cols) {
+  SetInfoFromForeignColumns(this, key, cols, n_cols);
+}
+}  // namespace xgboost

src/data/simple_csr_source.cu

@@ -0,0 +1,127 @@
+/*!
+ * Copyright 2019 by XGBoost Contributors
+ * \file simple_csr_source.cuh
+ * \brief An extension for the simple CSR source in-memory data structure to accept
+    foreign columnar data buffers, and convert them to XGBoost's internal DMatrix
+ * \author Andrey Adinets
+ * \author Matthew Jones
+ */
+#include <thrust/device_ptr.h>
+#include <thrust/device_vector.h>
+#include <thrust/execution_policy.h>
+#include <thrust/scan.h>
+
+#include <xgboost/base.h>
+#include <xgboost/data.h>
+#include <vector>
+#include <algorithm>
+
+#include "simple_csr_source.h"
+
+namespace xgboost {
+namespace data {
+
+__device__ int which_bit (int bit) {
+  return bit % 8;
+}
+__device__ int which_bitmap (int record) {
+  return record / 8;
+}
+
+__device__ int check_bit (foreign_valid_type bitmap, int bid) {
+  foreign_valid_type bitmask[8] = {1, 2, 4, 8, 16, 32, 64, 128};
+  return bitmap & bitmask[bid];
+}
+
+__device__ bool is_valid(foreign_valid_type * valid, int tid) {
+  if (valid == nullptr) {
+    return true;
+  }
+  int bmid = which_bitmap(tid);
+  int bid = which_bit(tid);
+  foreign_valid_type bitmap = valid[bmid];
+  return check_bit(bitmap, bid);
+}
+
+__global__ void CountValid(foreign_valid_type * valid,
+                           foreign_size_type n_rows,
+                           foreign_size_type n_cols,
+                           size_t * offsets) {
+  int tid = threadIdx.x + blockDim.x * blockIdx.x;
+  if (n_rows <= tid) {
+    return;
+  } else if (is_valid(valid, tid)) {
+    ++offsets[tid];
+  }
+}
+
+__global__ void CreateCSR(ForeignColumn * col, int col_idx, ForeignCSR * csr) {
+  int tid = threadIdx.x + blockDim.x * blockIdx.x;
+  if (col->size <= tid) {
+    return;
+  } else if (is_valid(col->valid, tid)) {
+    foreign_size_type oid = csr->offsets[tid];
+    float * d = (float *) (col->data);
+    csr->data[oid].fvalue = float(d[tid]);
+    csr->data[oid].index = col_idx;
+    ++csr->offsets[tid];
+  }
+}
+
+void ForeignColsToCSR(ForeignColumn ** cols, foreign_size_type n_cols, ForeignCSR * csr) {
+  foreign_size_type n_rows = cols[0]->size;
+  int threads = 1024;
+  int blocks = (n_rows + threads - 1) / threads;
+
+  dh::safe_cuda(cudaMemset(csr->offsets, 0, sizeof(foreign_size_type) * (n_rows + 1)));
+  if (0 < blocks) {
+    for (foreign_size_type i = 0 ; i < n_cols; ++i) {
+      CountValid <<<blocks, threads>>> (cols[i]->valid, n_rows, n_cols, csr->offsets);
+      dh::safe_cuda(cudaGetLastError());
+      dh::safe_cuda(cudaDeviceSynchronize());
+    }
+
+    thrust::device_ptr<size_t> offsets(csr->offsets);
+    int64_t n_valid = thrust::reduce(offsets, offsets + n_rows, 0ull, thrust::plus<size_t>());
+    thrust::exclusive_scan(offsets, offsets + n_rows + 1, offsets);
+
+    csr->n_nonzero = n_valid;
+    csr->n_rows = n_rows;
+    csr->n_cols = n_cols;
+
+    for (foreign_size_type i = 0; i < n_cols; ++i) {
+      CreateCSR <<<blocks, threads>>> (cols[i], i, csr);
+    }
+  }
+}
+
+void SimpleCSRSource::CopyFrom(ForeignColumn ** cols, foreign_size_type n_cols) {
+  CHECK_GT(n_cols, 0);
+  foreign_size_type n_valid = 0;
+  for (foreign_size_type i = 0; i < n_cols; ++i) {
+    CHECK_EQ(cols[0]->size, cols[i]->size);
+    n_valid += cols[i]->size - cols[i]->null_count;
+  }
+
+  info.num_col_ = n_cols;
+  info.num_row_ = cols[0]->size;
+  info.num_nonzero_ = n_valid;
+
+  GPUSet devices = GPUSet::Range(0, 1);
+  page_.offset.Reshard(GPUDistribution::Overlap(devices, 1));
+  page_.offset.Resize(cols[0]->size + 1);
+
+  std::vector<size_t> device_offsets{0, (size_t) n_valid};
+  page_.data.Reshard(GPUDistribution::Explicit(devices, device_offsets));
+  page_.data.Reshard(GPUDistribution::Overlap(devices, 1));
+  page_.data.Resize(n_valid);
+
+  ForeignCSR csr;
+  csr.data = page_.data.DevicePointer(0);
+  csr.offsets = page_.offset.DevicePointer(0);
+
+  ForeignColsToCSR(cols, n_cols, &csr);
+}
+
+}  // namespace data
+}  // namespace xgboost

src/data/simple_csr_source.h

@@ -35,7 +35,7 @@ class SimpleCSRSource : public DataSource {
   /*! \brief destructor */
   ~SimpleCSRSource() override = default;
   /*! \brief clear the data structure */
-  void Clear();
+  void Clear();  
   /*!
    * \brief copy content of data from src
    * \param src source data iter.
@@ -47,6 +47,12 @@ class SimpleCSRSource : public DataSource {
    * \param info The additional information reflected in the parser.
    */
   void CopyFrom(dmlc::Parser<uint32_t>* src);
+    /*!
+   * \brief copy content of data from foreign columns buffer.
+   * \param cols foreign columns data buffer.
+   * \param n_cols the number of foreign columns.
+   */
+  void CopyFrom(ForeignColumn ** cols, foreign_size_type n_cols);
   /*!
    * \brief Load data from binary stream.
    * \param fi the pointer to load data from.