grid.cu 1.17 KB
Newer Older
rusty1s's avatar
rusty1s committed
1
2
#include <THC.h>

rusty1s's avatar
rusty1s committed
3
#include "grid.h"
rusty1s's avatar
rusty1s committed
4

rusty1s's avatar
rusty1s committed
5
6
7
#include "common.cuh"
#include "THCIndex.cuh"

rusty1s's avatar
bugfix  
rusty1s committed
8
#define cluster_(NAME) TH_CONCAT_4(cluster_, NAME, _kernel_, Real)
rusty1s's avatar
rusty1s committed
9
#define thc_(NAME) TH_CONCAT_4(thc_, NAME, _, Real)
rusty1s's avatar
rusty1s committed
10

rusty1s's avatar
rusty1s committed
11
#include "generic/common.cu"
rusty1s's avatar
rusty1s committed
12
#include "THCGenerateAllTypes.h"
rusty1s's avatar
rusty1s committed
13
14

template<typename Real, int Dims>
15
__global__ void gridKernel(int64_t *output, TensorInfo<Real> position, Real *size, int64_t *count, const int C, const int N) {
rusty1s's avatar
rusty1s committed
16
  KERNEL_LOOP(i, N) {
17
    int positionOffset = 0; int tmp = C; int64_t c = 0;
rusty1s's avatar
rusty1s committed
18
    IndexToOffset<Real, Dims>::compute(i, position, &positionOffset);
19
20
21
    for (int d = 0; d < position.size[position.dims - 1]; d++) {
      tmp = tmp / count[d];
      c += tmp * (int64_t) (position.data[positionOffset + d] / size[d]);
rusty1s's avatar
rusty1s committed
22
    }
23
    output[i] = c;
rusty1s's avatar
rusty1s committed
24
25
26
  }
}

rusty1s's avatar
rusty1s committed
27
#include "generic/grid.cu"
rusty1s's avatar
rusty1s committed
28
#include "THCGenerateFloatType.h"
rusty1s's avatar
rusty1s committed
29
#include "generic/grid.cu"
rusty1s's avatar
rusty1s committed
30
#include "THCGenerateDoubleType.h"
rusty1s's avatar
rusty1s committed
31
#include "generic/grid.cu"
rusty1s's avatar
rusty1s committed
32
#include "THCGenerateByteType.h"
rusty1s's avatar
rusty1s committed
33
#include "generic/grid.cu"
rusty1s's avatar
rusty1s committed
34
#include "THCGenerateCharType.h"
rusty1s's avatar
rusty1s committed
35
#include "generic/grid.cu"
rusty1s's avatar
rusty1s committed
36
#include "THCGenerateShortType.h"
rusty1s's avatar
rusty1s committed
37
#include "generic/grid.cu"
rusty1s's avatar
rusty1s committed
38
#include "THCGenerateIntType.h"
rusty1s's avatar
rusty1s committed
39
#include "generic/grid.cu"
rusty1s's avatar
rusty1s committed
40
#include "THCGenerateLongType.h"