THCPropose.cuh 2.31 KB
Newer Older
rusty1s's avatar
rusty1s committed
1
2
3
#ifndef THC_PROPOSE_INC
#define THC_PROPOSE_INC

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

__global__ void proposeKernel(int64_t *color, int64_t *prop, int64_t *row, int64_t *col,
rusty1s's avatar
rusty1s committed
8
                              int64_t *degree, int64_t *cumDegree, ptrdiff_t nNodes) {
rusty1s's avatar
rusty1s committed
9
  KERNEL_LOOP(i, nNodes) {
rusty1s's avatar
rusty1s committed
10
11
    if (color[i] != -1) { continue; }  // Only visit blue nodes.
    ptrdiff_t c; bool isDead = true;
rusty1s's avatar
rusty1s committed
12
    for (ptrdiff_t e = cumDegree[i] - degree[i]; e < cumDegree[i]; e++) {
rusty1s's avatar
rusty1s committed
13
      c = col[e];
rusty1s's avatar
rusty1s committed
14
15
      if (isDead && color[c] < 0) { isDead = false; }  // Unmatched neighbor found.
      if (color[c] == -2) { prop[i] = c; break; }  // Propose to first red neighbor.
rusty1s's avatar
rusty1s committed
16
    }
rusty1s's avatar
rusty1s committed
17
    if (isDead) { color[i] = i; }  // Mark node as dead.
rusty1s's avatar
rusty1s committed
18
19
20
  }
}

rusty1s's avatar
rusty1s committed
21
22
23
24
25
template<typename T>
__global__ void weightedProposeKernel(int64_t *color, int64_t *prop, int64_t *row, int64_t *col,
                                      T *weight, int64_t *degree, int64_t *cumDegree,
                                      ptrdiff_t nNodes) {
  KERNEL_LOOP(i, nNodes) {
rusty1s's avatar
rusty1s committed
26
27
    if (color[i] != -1) { continue; }  // Only visit blue nodes.
    ptrdiff_t c; bool isDead = true;
rusty1s's avatar
bugfix  
rusty1s committed
28
29
    T maxWeight = ScalarConvert<int, T>::to(0), tmp;
    int64_t matchedValue = -1;
rusty1s's avatar
rusty1s committed
30
31
32
33
    for (ptrdiff_t e = cumDegree[i] - degree[i]; e < cumDegree[i]; e++) {
      c = col[e];
      tmp = weight[e];
      if (isDead && color[c] < 0) { isDead = false; }  // Unmatched neighbor found.
rusty1s's avatar
rusty1s committed
34
      // Find maximum weighted red neighbor.
rusty1s's avatar
rusty1s committed
35
36
37
38
39
      if (color[c] == -2 && THCNumerics<T>::gt(tmp, maxWeight)) {
        matchedValue = c;
        maxWeight = tmp;
      }
    }
rusty1s's avatar
rusty1s committed
40
    prop[i] = matchedValue;  // Propose.
rusty1s's avatar
rusty1s committed
41
    if (isDead) { color[i] = i; }  // Mark node as dead.
rusty1s's avatar
rusty1s committed
42
43
44
  }
}

rusty1s's avatar
rusty1s committed
45
void THCTensor_propose(THCState *state, THCudaLongTensor *color, THCudaLongTensor *prop,
rusty1s's avatar
rusty1s committed
46
47
48
49
50
                       THCudaLongTensor *row, THCudaLongTensor *col, THCudaLongTensor *degree,
                       THCudaLongTensor *cumDegree) {
  KERNEL_RUN(proposeKernel, THCudaLongTensor_nElement(state, color),
             THCudaLongTensor_data(state, color), THCudaLongTensor_data(state, prop),
             THCudaLongTensor_data(state, row), THCudaLongTensor_data(state, col),
rusty1s's avatar
rusty1s committed
51
             THCudaLongTensor_data(state, degree), THCudaLongTensor_data(state, cumDegree));
rusty1s's avatar
rusty1s committed
52
}
rusty1s's avatar
rusty1s committed
53

rusty1s's avatar
rusty1s committed
54
55
56
#include "generic/THCPropose.cuh"
#include "THC/THCGenerateAllTypes.h"

rusty1s's avatar
rusty1s committed
57
#endif  // THC_PROPOSE_INC