1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
|
#ifndef THC_TENSORSORT_CUH
#define THC_TENSORSORT_CUH
#include "THCReduceApplyUtils.cuh"
#include "THCSortUtils.cuh"
#include "THCTensorCopy.h"
#include "THCTensorTypeUtils.cuh"
#include <thrust/device_ptr.h>
#include <thrust/sort.h>
#if CUDA_VERSION >= 7000
#include <thrust/system/cuda/execution_policy.h>
#endif
template <typename T>
struct ThrustGTOp {
__device__ bool operator()(const T& lhs, const T& rhs) const {
return THCNumerics<T>::gt(lhs, rhs);
}
};
template <typename T>
struct ThrustLTOp {
__device__ bool operator()(const T& lhs, const T& rhs) const {
return THCNumerics<T>::lt(lhs, rhs);
}
};
// `base` is the base address of a tensor
// For each slice (defined as a linear point of `out`, from 0 ->
// (sliceSize - 1) * sliceStride, we fill that slice from `0` to
// `sliceSize - 1`.
template <typename IndexType, int Dim>
__global__ void
fillSliceWithIndex(TensorInfo<long, IndexType> out,
IndexType totalSlices,
IndexType sliceSize,
IndexType sliceStride) {
IndexType slice = getLinearBlockId<IndexType>();
if (slice >= totalSlices) {
return;
}
const unsigned long offset =
IndexToOffset<long, IndexType, Dim>::get(slice, out);
long* base = &out.data[offset];
for (long i = threadIdx.x; i < sliceSize; i += blockDim.x) {
// Torch indices are 1-based (hence the +1)
base[i * sliceStride] = i + TH_INDEX_BASE;
}
}
// For slice sorting in Thrust; extracts a slice index from a linear
// index and uses that for comparison
struct SliceComp {
SliceComp(long size) : sliceSize(size) {}
__device__ bool operator()(const long& a, const long& b) const {
// Since the slices are guaranteed to be innermost, the segment is
// just via long division
long segA = a / sliceSize;
long segB = b / sliceSize;
return segA < segB;
}
const long sliceSize;
};
// For sorting in Thurst; extracts a within-slice index from a linear index
struct GlobalIndexToPerSliceIndex {
GlobalIndexToPerSliceIndex(long size) : sliceSize(size) {}
__device__ inline void operator()(long& v) const {
v = v % sliceSize + TH_INDEX_BASE;
}
const long sliceSize;
};
unsigned long nextHighestPowerOf2(unsigned long n);
void THCudaLongTensor_fillSliceWithIndex(THCState* state,
THCudaLongTensor* t,
int dim);
#endif // THC_TENSORSORT_CUH
|
