Fix empty subspan. (#4151)

* Silent the death tests.

src/common/span.h

@@ -69,6 +69,7 @@ namespace common {
 
 // Usual logging facility is not available inside device code.
 // TODO(trivialfis): Make dmlc check more generic.
+// assert is not supported in mac as of CUDA 10.0
 #define KERNEL_CHECK(cond)                                      \
   do {                                                          \
     if (!(cond)) {                                              \
@@ -543,7 +544,7 @@ class Span {
   XGBOOST_DEVICE auto subspan() const ->                   // NOLINT
       Span<element_type,
            detail::ExtentValue<Extent, Offset, Count>::value> {
-    SPAN_CHECK(Offset >= 0 && Offset < size());
+    SPAN_CHECK(Offset >= 0 && (Offset < size() || size() == 0));
     SPAN_CHECK(Count == dynamic_extent ||
                Count >= 0 && Offset + Count <= size());
 
@@ -553,7 +554,7 @@ class Span {
   XGBOOST_DEVICE Span<element_type, dynamic_extent> subspan(  // NOLINT
       detail::ptrdiff_t _offset,
       detail::ptrdiff_t _count = dynamic_extent) const {
-    SPAN_CHECK(_offset >= 0 && _offset < size());
+    SPAN_CHECK(_offset >= 0 && (_offset < size() || size() == 0));
     SPAN_CHECK((_count == dynamic_extent) ||
                (_count >= 0 && _offset + _count <= size()));
 

tests/cpp/common/test_span.cc

@@ -419,5 +419,29 @@ TEST(Span, AsWritableBytes) {
   ASSERT_EQ(status, 1);
 }
 
+TEST(Span, Empty) {
+  {
+    Span<float> s {nullptr, static_cast<Span<float>::index_type>(0)};
+    auto res = s.subspan(0);
+    ASSERT_EQ(res.data(), nullptr);
+    ASSERT_EQ(res.size(), 0);
+
+    res = s.subspan(0, 0);
+    ASSERT_EQ(res.data(), nullptr);
+    ASSERT_EQ(res.size(), 0);
+  }
+
+  {
+    Span<float, 0> s {nullptr, static_cast<Span<float>::index_type>(0)};
+    auto res = s.subspan(0);
+    ASSERT_EQ(res.data(), nullptr);
+    ASSERT_EQ(res.size(), 0);
+
+    res = s.subspan(0, 0);
+    ASSERT_EQ(res.data(), nullptr);
+    ASSERT_EQ(res.size(), 0);
+  }
+}
+
 }  // namespace common
 }  // namespace xgboost

tests/cpp/common/test_span.cu

@@ -216,15 +216,20 @@ struct TestElementAccess {
 
 TEST(GPUSpan, ElementAccess) {
   dh::safe_cuda(cudaSetDevice(0));
-  EXPECT_DEATH({
+  auto test_element_access = []() {
-      thrust::host_vector<float> h_vec (16);
+    thrust::host_vector<float> h_vec (16);
-      InitializeRange(h_vec.begin(), h_vec.end());
+    InitializeRange(h_vec.begin(), h_vec.end());
 
-      thrust::device_vector<float> d_vec (h_vec.size());
+    thrust::device_vector<float> d_vec (h_vec.size());
-      thrust::copy(h_vec.begin(), h_vec.end(), d_vec.begin());
+    thrust::copy(h_vec.begin(), h_vec.end(), d_vec.begin());
 
-      Span<float> span (d_vec.data().get(), d_vec.size());
+    Span<float> span (d_vec.data().get(), d_vec.size());
-      dh::LaunchN(0, 17, TestElementAccess{span});}, "");
+    dh::LaunchN(0, 17, TestElementAccess{span});
+  };
+
+  testing::internal::CaptureStdout();
+  EXPECT_DEATH(test_element_access(), "");
+  std::string output = testing::internal::GetCapturedStdout();
 }
 
 __global__ void test_first_dynamic_kernel(Span<float> _span) {
@@ -253,7 +258,9 @@ TEST(GPUSpan, FirstLast) {
     Span<float> span (d_vec.data().get(), d_vec.size());
     test_first_dynamic_kernel<<<1, 1>>>(span);
   };
+  testing::internal::CaptureStdout();
   EXPECT_DEATH(lambda_first_dy(), "");
+  std::string output = testing::internal::GetCapturedStdout();
 
   auto lambda_first_static = []() {
     thrust::host_vector<float> h_vec (4);
@@ -265,7 +272,9 @@ TEST(GPUSpan, FirstLast) {
     Span<float> span (d_vec.data().get(), d_vec.size());
     test_first_static_kernel<<<1, 1>>>(span);
   };
+  testing::internal::CaptureStdout();
   EXPECT_DEATH(lambda_first_static(), "");
+  output = testing::internal::GetCapturedStdout();
 
   auto lambda_last_dy = []() {
     thrust::host_vector<float> h_vec (4);
@@ -277,7 +286,9 @@ TEST(GPUSpan, FirstLast) {
     Span<float> span (d_vec.data().get(), d_vec.size());
     test_last_dynamic_kernel<<<1, 1>>>(span);
   };
+  testing::internal::CaptureStdout();
   EXPECT_DEATH(lambda_last_dy(), "");
+  output = testing::internal::GetCapturedStdout();
 
   auto lambda_last_static = []() {
     thrust::host_vector<float> h_vec (4);
@@ -289,7 +300,9 @@ TEST(GPUSpan, FirstLast) {
     Span<float> span (d_vec.data().get(), d_vec.size());
     test_last_static_kernel<<<1, 1>>>(span);
   };
+  testing::internal::CaptureStdout();
   EXPECT_DEATH(lambda_last_static(), "");
+  output = testing::internal::GetCapturedStdout();
 }
 
 
@@ -310,7 +323,9 @@ TEST(GPUSpan, Subspan) {
     Span<float> span (d_vec.data().get(), d_vec.size());
     test_subspan_dynamic_kernel<<<1, 1>>>(span);
   };
+  testing::internal::CaptureStdout();
   EXPECT_DEATH(lambda_subspan_dynamic(), "");
+  std::string output = testing::internal::GetCapturedStdout();
 
   auto lambda_subspan_static = []() {
     thrust::host_vector<float> h_vec (4);
@@ -322,7 +337,9 @@ TEST(GPUSpan, Subspan) {
     Span<float> span (d_vec.data().get(), d_vec.size());
     test_subspan_static_kernel<<<1, 1>>>(span);
   };
+  testing::internal::CaptureStdout();
   EXPECT_DEATH(lambda_subspan_static(), "");
+  output = testing::internal::GetCapturedStdout();
 }
 
 TEST(GPUSpanIter, Construct) {