wykvictor / cs344-cuda-udacity Goto Github PK

__global__ void scanHillisSteele(int *d_out, int *d_in, int n) {
  int idx = threadIdx.x;
  extern __shared__ int temp[];
  int pout = 0, pin = 1;

  temp[idx] = (idx > 0) ? d_in[idx - 1] : 0;
  __syncthreads();

  for (int offset = 1; offset < n; offset *= 2) {
    // swap double buffer indices
    pout = 1 - pout;
    pin = 1 - pout;
    if (idx >= offset) {
      temp[pout*n+idx] = temp[pin*n+idx - offset] + temp[pin*n+idx];  // changed line
    } else {
      temp[pout*n+idx] = temp[pin*n+idx];
    }
    __syncthreads();
  }
  d_out[idx] = temp[pout*n+idx];
}