   CHECK_EQ(values.CheckContiguous(), true);
 #if CUDA_VERSION >= 7000
   cudaStream_t stream = mshadow::Stream<gpu>::GetStream(keys.stream_);
+#ifndef SORT_WITH_THRUST
   if (workspace != NULL) {
     // Workspace given, sort using CUB
     CHECK_EQ(workspace->CheckContiguous(), true);

    mshadow::Copy(keys, keys_out, keys.stream_);
    mshadow::Copy(values, values_out, values.stream_);
  } else {
    // No workspace, sort using thrust
    thrust::device_ptr<KDType> key_iter = thrust::device_pointer_cast(keys.dptr_);
    thrust::device_ptr<VDType> value_iter = thrust::device_pointer_cast(values.dptr_);

        thrust::cuda::par.on(stream),
        key_iter, key_iter + keys.size(0), value_iter, thrust::greater<KDType>());
    }
#ifndef SORT_WITH_THRUST
  }
#endif // SORT_WITH_THRUST
  MSHADOW_CUDA_POST_KERNEL_CHECK(SortByKey);
#else
  LOG(FATAL) << "SortByKey is only supported for CUDA version >=7.0!";
