THCBasis.cu 2.95 KB
Newer Older
rusty1s's avatar
rusty1s committed
1
2
#include "THCBasis.h"

rusty1s's avatar
rusty1s committed
3
#include "THCBasisForward.cuh"
rusty1s's avatar
rusty1s committed
4
#include "THCBasisBackward.cuh"
rusty1s's avatar
rusty1s committed
5
6
7
8
9
10

template<typename T>
__global__ void linearBasisForwardKernel(TensorInfo<T> basis, TensorInfo<int64_t>weightIndex,
                                         TensorInfo<T> pseudo, int64_t *kernelSize,
                                         uint8_t *isOpenSpline, ptrdiff_t n) {
  THC_TENSOR_BASIS_FORWARD_KERNEL(1, basis, weightIndex, pseudo, kernelSize, isOpenSpline, n,
rusty1s's avatar
rusty1s committed
11
                                  BasisForward<T>::linear(v, kMod))
rusty1s's avatar
rusty1s committed
12
13
14
15
16
17
18
}

template<typename T>
__global__ void quadraticBasisForwardKernel(TensorInfo<T> basis, TensorInfo<int64_t>weightIndex,
                                            TensorInfo<T> pseudo, int64_t *kernelSize,
                                            uint8_t *isOpenSpline, ptrdiff_t n) {
  THC_TENSOR_BASIS_FORWARD_KERNEL(2, basis, weightIndex, pseudo, kernelSize, isOpenSpline, n,
rusty1s's avatar
rusty1s committed
19
                                  BasisForward<T>::quadratic(v, kMod))
rusty1s's avatar
rusty1s committed
20
21
22
23
24
25
26
}

template<typename T>
__global__ void cubicBasisForwardKernel(TensorInfo<T> basis, TensorInfo<int64_t>weightIndex,
                                        TensorInfo<T> pseudo, int64_t *kernelSize,
                                        uint8_t *isOpenSpline, ptrdiff_t n) {
  THC_TENSOR_BASIS_FORWARD_KERNEL(3, basis, weightIndex, pseudo, kernelSize, isOpenSpline, n,
rusty1s's avatar
rusty1s committed
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
                                  BasisForward<T>::cubic(v, kMod))
}

template<typename T>
__global__ void linearBasisBackwardKernel(TensorInfo<T> self, TensorInfo<T>gradBasis,
                                          TensorInfo<T> pseudo, int64_t *kernelSize,
                                          uint8_t *isOpenSpline, ptrdiff_t n) {
  THC_TENSOR_BASIS_BACKWARD_KERNEL(1, self, gradBasis, pseudo, kernelSize, isOpenSpline, n,
                                   BasisForward<T>::linear(v, kMod),
                                   BasisBackward<T>::linear(v, kMod))
}

template<typename T>
__global__ void quadraticBasisBackwardKernel(TensorInfo<T> self, TensorInfo<T>gradBasis,
                                             TensorInfo<T> pseudo, int64_t *kernelSize,
                                             uint8_t *isOpenSpline, ptrdiff_t n) {
  THC_TENSOR_BASIS_BACKWARD_KERNEL(2, self, gradBasis, pseudo, kernelSize, isOpenSpline, n,
                                   BasisForward<T>::quadratic(v, kMod),
                                   BasisBackward<T>::quadratic(v, kMod))
}

template<typename T>
__global__ void cubicBasisBackwardKernel(TensorInfo<T> self, TensorInfo<T>gradBasis,
                                         TensorInfo<T> pseudo, int64_t *kernelSize,
                                         uint8_t *isOpenSpline, ptrdiff_t n) {
  THC_TENSOR_BASIS_BACKWARD_KERNEL(3, self, gradBasis, pseudo, kernelSize, isOpenSpline, n,
                                   BasisForward<T>::cubic(v, kMod),
                                   BasisBackward<T>::cubic(v, kMod))
rusty1s's avatar
rusty1s committed
55
56
57
58
}

#include "generic/THCBasis.cu"
#include "THC/THCGenerateFloatTypes.h"