diff --git a/Jenkinsfile b/Jenkinsfile
index 66dea2c7d..a038d2726 100644
--- a/Jenkinsfile
+++ b/Jenkinsfile
@@ -64,7 +64,7 @@ pipeline {
// The build-gpu-* builds below use Ubuntu image
'build-gpu-cuda11.0': { BuildCUDA(cuda_version: '11.0', build_rmm: true) },
'build-gpu-rpkg': { BuildRPackageWithCUDA(cuda_version: '10.1') },
- 'build-jvm-packages-gpu-cuda10.1': { BuildJVMPackagesWithCUDA(spark_version: '3.0.0', cuda_version: '10.1') },
+ 'build-jvm-packages-gpu-cuda10.1': { BuildJVMPackagesWithCUDA(spark_version: '3.0.0', cuda_version: '11.0') },
'build-jvm-packages': { BuildJVMPackages(spark_version: '3.0.0') },
'build-jvm-doc': { BuildJVMDoc() }
])
diff --git a/jvm-packages/CMakeLists.txt b/jvm-packages/CMakeLists.txt
index c7aa32f97..247c44378 100644
--- a/jvm-packages/CMakeLists.txt
+++ b/jvm-packages/CMakeLists.txt
@@ -1,10 +1,20 @@
find_package(JNI REQUIRED)
-add_library(xgboost4j SHARED
- ${PROJECT_SOURCE_DIR}/jvm-packages/xgboost4j/src/native/xgboost4j.cpp)
+list(APPEND JVM_SOURCES
+ ${PROJECT_SOURCE_DIR}/jvm-packages/xgboost4j/src/native/xgboost4j.cpp
+ ${PROJECT_SOURCE_DIR}/jvm-packages/xgboost4j-gpu/src/native/xgboost4j-gpu.cpp)
+
+if (USE_CUDA)
+ list(APPEND JVM_SOURCES
+ ${PROJECT_SOURCE_DIR}/jvm-packages/xgboost4j-gpu/src/native/xgboost4j-gpu.cu)
+endif (USE_CUDA)
+
+add_library(xgboost4j SHARED ${JVM_SOURCES} ${XGBOOST_OBJ_SOURCES})
+
if (ENABLE_ALL_WARNINGS)
target_compile_options(xgboost4j PUBLIC -Wall -Wextra)
endif (ENABLE_ALL_WARNINGS)
+
target_link_libraries(xgboost4j PRIVATE objxgboost)
target_include_directories(xgboost4j
PRIVATE
@@ -15,8 +25,4 @@ target_include_directories(xgboost4j
${PROJECT_SOURCE_DIR}/rabit/include)
set_output_directory(xgboost4j ${PROJECT_SOURCE_DIR}/lib)
-set_target_properties(
- xgboost4j PROPERTIES
- CXX_STANDARD 14
- CXX_STANDARD_REQUIRED ON)
target_link_libraries(xgboost4j PRIVATE ${JAVA_JVM_LIBRARY})
diff --git a/jvm-packages/xgboost4j-gpu/pom.xml b/jvm-packages/xgboost4j-gpu/pom.xml
index 81b536e8e..6c2e22b7d 100644
--- a/jvm-packages/xgboost4j-gpu/pom.xml
+++ b/jvm-packages/xgboost4j-gpu/pom.xml
@@ -12,7 +12,24 @@
1.5.0-SNAPSHOT
jar
+
+ 21.08.2
+ cuda11
+
+
+
+ ai.rapids
+ cudf
+ ${cudf.version}
+ ${cudf.classifier}
+ provided
+
+
+ com.fasterxml.jackson.core
+ jackson-databind
+ 2.10.5.1
+
org.apache.hadoop
hadoop-hdfs
diff --git a/jvm-packages/xgboost4j-gpu/src/main/java b/jvm-packages/xgboost4j-gpu/src/main/java
deleted file mode 120000
index 2e2be8ff1..000000000
--- a/jvm-packages/xgboost4j-gpu/src/main/java
+++ /dev/null
@@ -1 +0,0 @@
-../../../xgboost4j/src/main/java/
\ No newline at end of file
diff --git a/jvm-packages/xgboost4j-gpu/src/main/java/ml/dmlc/xgboost4j/gpu/java/CudfColumn.java b/jvm-packages/xgboost4j-gpu/src/main/java/ml/dmlc/xgboost4j/gpu/java/CudfColumn.java
new file mode 100644
index 000000000..06501cbfa
--- /dev/null
+++ b/jvm-packages/xgboost4j-gpu/src/main/java/ml/dmlc/xgboost4j/gpu/java/CudfColumn.java
@@ -0,0 +1,110 @@
+/*
+ Copyright (c) 2021 by Contributors
+
+ 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.
+ */
+
+package ml.dmlc.xgboost4j.gpu.java;
+
+import ai.rapids.cudf.BaseDeviceMemoryBuffer;
+import ai.rapids.cudf.BufferType;
+import ai.rapids.cudf.ColumnVector;
+import ai.rapids.cudf.DType;
+
+import ml.dmlc.xgboost4j.java.Column;
+
+/**
+ * This class is composing of base data with Apache Arrow format from Cudf ColumnVector.
+ * It will be used to generate the cuda array interface.
+ */
+class CudfColumn extends Column {
+
+ private final long dataPtr; // gpu data buffer address
+ private final long shape; // row count
+ private final long validPtr; // gpu valid buffer address
+ private final int typeSize; // type size in bytes
+ private final String typeStr; // follow array interface spec
+ private final long nullCount; // null count
+
+ private String arrayInterface = null; // the cuda array interface
+
+ public static CudfColumn from(ColumnVector cv) {
+ BaseDeviceMemoryBuffer dataBuffer = cv.getDeviceBufferFor(BufferType.DATA);
+ BaseDeviceMemoryBuffer validBuffer = cv.getDeviceBufferFor(BufferType.VALIDITY);
+ long validPtr = 0;
+ if (validBuffer != null) {
+ validPtr = validBuffer.getAddress();
+ }
+ DType dType = cv.getType();
+ String typeStr = "";
+ if (dType == DType.FLOAT32 || dType == DType.FLOAT64 ||
+ dType == DType.TIMESTAMP_DAYS || dType == DType.TIMESTAMP_MICROSECONDS ||
+ dType == DType.TIMESTAMP_MILLISECONDS || dType == DType.TIMESTAMP_NANOSECONDS ||
+ dType == DType.TIMESTAMP_SECONDS) {
+ typeStr = " table.getColumn(i))
+ .map(CudfColumn::from)
+ .toArray(CudfColumn[]::new);
+ }
+
+}
diff --git a/jvm-packages/xgboost4j-gpu/src/main/java/ml/dmlc/xgboost4j/gpu/java/CudfUtils.java b/jvm-packages/xgboost4j-gpu/src/main/java/ml/dmlc/xgboost4j/gpu/java/CudfUtils.java
new file mode 100644
index 000000000..748024bd9
--- /dev/null
+++ b/jvm-packages/xgboost4j-gpu/src/main/java/ml/dmlc/xgboost4j/gpu/java/CudfUtils.java
@@ -0,0 +1,100 @@
+/*
+ Copyright (c) 2021 by Contributors
+
+ 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.
+ */
+
+package ml.dmlc.xgboost4j.gpu.java;
+
+import java.io.ByteArrayOutputStream;
+import java.io.IOException;
+
+import com.fasterxml.jackson.core.JsonFactory;
+import com.fasterxml.jackson.core.JsonGenerator;
+import com.fasterxml.jackson.databind.ObjectMapper;
+import com.fasterxml.jackson.databind.node.ArrayNode;
+import com.fasterxml.jackson.databind.node.JsonNodeFactory;
+import com.fasterxml.jackson.databind.node.ObjectNode;
+
+/**
+ * Cudf utilities to build cuda array interface against {@link CudfColumn}
+ */
+class CudfUtils {
+
+ /**
+ * Build the cuda array interface based on CudfColumn(s)
+ * @param cudfColumns the CudfColumn(s) to be built
+ * @return the json format of cuda array interface
+ */
+ public static String buildArrayInterface(CudfColumn... cudfColumns) {
+ return new Builder().add(cudfColumns).build();
+ }
+
+ // Helper class to build array interface string
+ private static class Builder {
+ private JsonNodeFactory nodeFactory = new JsonNodeFactory(false);
+ private ArrayNode rootArrayNode = nodeFactory.arrayNode();
+
+ private Builder add(CudfColumn... columns) {
+ if (columns == null || columns.length <= 0) {
+ throw new IllegalArgumentException("At least one ColumnData is required.");
+ }
+ for (CudfColumn cd : columns) {
+ rootArrayNode.add(buildColumnObject(cd));
+ }
+ return this;
+ }
+
+ private String build() {
+ try {
+ ByteArrayOutputStream bos = new ByteArrayOutputStream();
+ JsonGenerator jsonGen = new JsonFactory().createGenerator(bos);
+ new ObjectMapper().writeTree(jsonGen, rootArrayNode);
+ return bos.toString();
+ } catch (IOException ie) {
+ ie.printStackTrace();
+ throw new RuntimeException("Failed to build array interface. Error: " + ie);
+ }
+ }
+
+ private ObjectNode buildColumnObject(CudfColumn column) {
+ if (column.getDataPtr() == 0) {
+ throw new IllegalArgumentException("Empty column data is NOT accepted!");
+ }
+ if (column.getTypeStr() == null || column.getTypeStr().isEmpty()) {
+ throw new IllegalArgumentException("Empty type string is NOT accepted!");
+ }
+ ObjectNode colDataObj = buildMetaObject(column.getDataPtr(), column.getShape(),
+ column.getTypeStr());
+
+ if (column.getValidPtr() != 0 && column.getNullCount() != 0) {
+ ObjectNode validObj = buildMetaObject(column.getValidPtr(), column.getShape(), "
+
+#include "../../../../src/common/common.h"
+#include "../../../../src/c_api/c_api_error.h"
+
+namespace xgboost {
+namespace jni {
+XGB_DLL int XGDeviceQuantileDMatrixCreateFromCallbackImpl(JNIEnv *jenv, jclass jcls,
+ jobject jiter,
+ jfloat jmissing,
+ jint jmax_bin, jint jnthread,
+ jlongArray jout) {
+ API_BEGIN();
+ common::AssertGPUSupport();
+ API_END();
+}
+} // namespace jni
+} // namespace xgboost
+#endif // XGBOOST_USE_CUDA
diff --git a/jvm-packages/xgboost4j-gpu/src/native/xgboost4j-gpu.cu b/jvm-packages/xgboost4j-gpu/src/native/xgboost4j-gpu.cu
new file mode 100644
index 000000000..de7a1fc41
--- /dev/null
+++ b/jvm-packages/xgboost4j-gpu/src/native/xgboost4j-gpu.cu
@@ -0,0 +1,398 @@
+#include
+#include
+
+#include "../../../../src/common/device_helpers.cuh"
+#include "../../../../src/data/array_interface.h"
+#include "jvm_utils.h"
+#include
+
+namespace xgboost {
+namespace jni {
+
+template
+T const *RawPtr(std::vector const &data) {
+ return data.data();
+}
+
+template T *RawPtr(std::vector &data) {
+ return data.data();
+}
+
+template T const *RawPtr(dh::device_vector const &data) {
+ return data.data().get();
+}
+
+template T *RawPtr(dh::device_vector &data) {
+ return data.data().get();
+}
+
+template T CheckJvmCall(T const &v, JNIEnv *jenv) {
+ if (!v) {
+ CHECK(jenv->ExceptionOccurred());
+ jenv->ExceptionDescribe();
+ }
+ return v;
+}
+
+template
+void CopyColumnMask(xgboost::ArrayInterface const &interface,
+ std::vector const &columns, cudaMemcpyKind kind,
+ size_t c, VCont *p_mask, Json *p_out, cudaStream_t stream) {
+ auto &mask = *p_mask;
+ auto &out = *p_out;
+ auto size = sizeof(typename VCont::value_type) * interface.num_rows *
+ interface.num_cols;
+ mask.resize(size);
+ CHECK(RawPtr(mask));
+ CHECK(size);
+ CHECK(interface.valid.Data());
+ dh::safe_cuda(
+ cudaMemcpyAsync(RawPtr(mask), interface.valid.Data(), size, kind, stream));
+ auto const &mask_column = columns[c]["mask"];
+ out["mask"] = Object();
+ std::vector mask_data{
+ Json{reinterpret_cast(RawPtr(mask))},
+ Json{get(mask_column["data"][1])}};
+ out["mask"]["data"] = Array(std::move(mask_data));
+ if (get(mask_column["shape"]).size() == 2) {
+ std::vector mask_shape{
+ Json{get(mask_column["shape"][0])},
+ Json{get(mask_column["shape"][1])}};
+ out["mask"]["shape"] = Array(std::move(mask_shape));
+ } else if (get(mask_column["shape"]).size() == 1) {
+ std::vector mask_shape{
+ Json{get(mask_column["shape"][0])}};
+ out["mask"]["shape"] = Array(std::move(mask_shape));
+ } else {
+ LOG(FATAL) << "Invalid shape of mask";
+ }
+ out["mask"]["typestr"] = String("
+void CopyInterface(std::vector &interface_arr,
+ std::vector const &columns, cudaMemcpyKind kind,
+ std::vector *p_data, std::vector *p_mask,
+ std::vector *p_out, cudaStream_t stream) {
+ p_data->resize(interface_arr.size());
+ p_mask->resize(interface_arr.size());
+ p_out->resize(interface_arr.size());
+ for (size_t c = 0; c < interface_arr.size(); ++c) {
+ auto &interface = interface_arr.at(c);
+ size_t element_size = interface.ElementSize();
+ size_t size = element_size * interface.num_rows * interface.num_cols;
+
+ auto &data = (*p_data)[c];
+ auto &mask = (*p_mask)[c];
+ data.resize(size);
+ dh::safe_cuda(cudaMemcpyAsync(RawPtr(data), interface.data, size, kind, stream));
+
+ auto &out = (*p_out)[c];
+ out = Object();
+ std::vector j_data{
+ Json{Integer(reinterpret_cast(RawPtr(data)))},
+ Json{Boolean{false}}};
+
+ out["data"] = Array(std::move(j_data));
+ out["shape"] = Array(std::vector{Json(Integer(interface.num_rows)),
+ Json(Integer(interface.num_cols))});
+
+ if (interface.valid.Data()) {
+ CopyColumnMask(interface, columns, kind, c, &mask, &out, stream);
+ }
+ out["typestr"] = String(" *out, cudaStream_t stream) {
+ auto &j_interface = *p_interface;
+ CHECK_EQ(get(j_interface).size(), 1);
+ auto object = get