 #define MXNET_OPERATOR_TENSOR_SORT_OP_INL_CUH_
 #include <thrust/device_ptr.h>
 #include <thrust/sort.h>
+#if defined(_MSC_VER) && __CUDACC_VER__ != 80044
+// Many CUDA compilers other than V8.0.44 crash on Windows 
+#pragma warning("Potential crash on CUDA compiler detected. Switching sorting from CUB to Thrust")
+#define SORT_WITH_THRUST
+#else
 #include <cub/device/device_radix_sort.cuh>
+#undef SORT_WITH_THRUST
+#endif
 #if CUDA_VERSION >= 7000
 #include <thrust/system/cuda/execution_policy.h>
 #endif

template <typename KDType, typename VDType, typename xpu>
inline typename std::enable_if<std::is_same<xpu, gpu>::value, size_t>::type
SortByKeyWorkspaceSize(const size_t num_keys) {
  size_t sortpairs_bytes = 0;
  cub::DeviceRadixSort::SortPairs<KDType, VDType>(NULL, sortpairs_bytes,
      NULL, NULL, NULL, NULL, num_keys);
  size_t keys_bytes = num_keys*sizeof(KDType);
  size_t values_bytes = num_keys*sizeof(VDType);
  return (keys_bytes + values_bytes + sortpairs_bytes);
}

template<typename KDType, typename VDType>

        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!";
