From 0d11df1d35e902fbaab0dfb70735d59bf7a5d378 Mon Sep 17 00:00:00 2001
From: Johannes M Dieterich <Johannes.Dieterich@amd.com>
Date: Fri, 24 Aug 2018 01:26:25 -0500
Subject: [PATCH 1/2] Correctly type cast everything.

---
 aten/src/THC/generic/THCTensorSort.cu | 16 ++++++++--------
 1 file changed, 8 insertions(+), 8 deletions(-)

diff --git a/aten/src/THC/generic/THCTensorSort.cu b/aten/src/THC/generic/THCTensorSort.cu
index b6bcf6aecb4f78..764a508073f5df 100644
--- a/aten/src/THC/generic/THCTensorSort.cu
+++ b/aten/src/THC/generic/THCTensorSort.cu
@@ -56,21 +56,21 @@ THC_API void THCTensor_(sortKeyValueInplace)(THCState* state,
       bitonicSortKVInPlace<real, int64_t, A, -1, GTComp<real>, TYPE, SIZE> \
         <<<grid, block, 0, THCState_getCurrentStream(state)>>>(         \
           keyInfo,                                                      \
-          keySlices,                                                    \
-          (TYPE) keySliceSize,                                          \
-          (TYPE) keyInfo.strides[collapseKeyDim],                       \
+          static_cast<TYPE>(keySlices),                                 \
+          static_cast<TYPE>(keySliceSize),                              \
+          static_cast<TYPE>(keyInfo.strides[collapseKeyDim]),           \
           valueInfo,                                                    \
-          (TYPE) valueInfo.strides[collapseValueDim],                   \
+          static_cast<TYPE>(valueInfo.strides[collapseValueDim]),       \
           GTComp<real>());                                              \
     } else {                                                            \
       bitonicSortKVInPlace<real, int64_t, A, -1, LTComp<real>, TYPE, SIZE> \
         <<<grid, block, 0, THCState_getCurrentStream(state)>>>(         \
           keyInfo,                                                      \
-          keySlices,                                                    \
-          (TYPE) keySliceSize,                                          \
-          (TYPE) keyInfo.strides[collapseKeyDim],                       \
+          static_cast<TYPE>(keySlices),                                 \
+          static_cast<TYPE>(keySliceSize),                              \
+          static_cast<TYPE>(keyInfo.strides[collapseKeyDim]),           \
           valueInfo,                                                    \
-          (TYPE) valueInfo.strides[collapseValueDim],                   \
+          static_cast<TYPE>(valueInfo.strides[collapseValueDim]),       \
           LTComp<real>());                                              \
     }                                                                   \
   } while (0)

From 3ddabadb37a0ee23ebe53fa845a7fc89e268126b Mon Sep 17 00:00:00 2001
From: Johannes M Dieterich <Johannes.Dieterich@amd.com>
Date: Sun, 2 Sep 2018 18:34:28 -0500
Subject: [PATCH 2/2] CUDA does not really allow passing objects by reference
 to __global__ functions.

Submitted by: AlexVIx
---
 aten/src/THC/THCSortUtils.cuh | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/aten/src/THC/THCSortUtils.cuh b/aten/src/THC/THCSortUtils.cuh
index 518063a229c723..891e77067686e7 100644
--- a/aten/src/THC/THCSortUtils.cuh
+++ b/aten/src/THC/THCSortUtils.cuh
@@ -142,7 +142,7 @@ bitonicSortKVInPlace(TensorInfo<K, IndexType> keys,
                      IndexType keySliceStride,
                      TensorInfo<V, IndexType> values,
                      IndexType valueSliceStride,
-                     const Comparator& comp) {
+                     const Comparator comp) {
   // Find the slice of the tensor that we are sorting
   const IndexType linearIndex = getLinearBlockId<IndexType>();
   // Tiling the slices could have us be out of bounds, if there are a