1

有人可以帮助了解 Hillis & Steele：内核函数如何为每个线程执行工作吗？

__global__ void scan(float *g_odata, float *g_idata, int n)
 {
    extern __shared__ float temp[]; // allocated on invocation
    int thid = threadIdx.x;
    int pout = 0, pin = 1;
    // load input into shared memory.
    // This is exclusive scan, so shift right by one and set first elt to 0
    temp[pout*n + thid] = (thid > 0) ? g_idata[thid-1] : 0;
    __syncthreads();
    for (int offset = 1; offset < n; offset *= 2)
    {
      pout = 1 - pout; // swap double buffer indices
      pin = 1 - pout;
      if (thid >= offset)
        temp[pout*n+thid] += temp[pin*n+thid - offset];
      else
        temp[pout*n+thid] = temp[pin*n+thid];
     __syncthreads();
    }
    g_odata[thid] = temp[pout*n+thid1]; // write output
}

从现在开始，我已经了解以下内容：首先，我们有pout=0, pin=1 and thid = [1,bockDim.x]. 所以在第一次同步之前，我们有一个简单的向右移动，例如，如果我们有数组[1 | 2 | 5 | 7 ]，新数组是[0 |1 | 2 | 5 | 7 ].

我将执行for loop视为多个实例，每个实例对应每个thId. 例如，如果thId=0我们要执行以下操作：

  1. thid=0

    • offset=1
    • pout = 1-0=1 （在函数开头使用 pout 初始化）
    • 引脚 = 1 - 1 =0；（使用刚刚计算的噘嘴，ei 1）
    • temp[4] = temp[0]（else 语句）

    • [0 | 1 | 2 | 5 | 0]

    • offset=2

    • pout = 1-1=0 （使用循环中上一步的 pout）
    • 引脚 = 1 - 0 =1；（刚刚计算的值）
    • temp[0] = temp[4] ( else 语句)
    • [0 | 1 | 2 | 5 | 0]

    pout 和 pin 变量是根据 for 循环内部的信息更改的，而不是
    在开始时考虑这些变量的初始化。以我想象的相同方式
    执行thid=1.

  2. thid=1

    • offset=1
    • pout = 1 - 0 = 1（在函数开头使用 pout 初始化）
    • 引脚 = 1 - 1 = 0
    • temp[4+1] = temp[0+1-1] ( if 语句) ???? 内存超出温度范围？？？？

谁能给出一个直观的例子来说明它是如何执行的？另外，当最后一个代码语句将被执行时，将使用哪个 pout 值？

4

2 回答 2

2

如果您的意思是并行扫描算法，您可以在此处查看直观的解释。

http://en.wikipedia.org/wiki/Prefix_sum

在此处输入图像描述

我相信这些链接也有帮助。

http://nvlabs.github.io/cub/structcub_1_1_device_scan.html

http://nvlabs.github.io/cub/classcub_1_1_block_scan.html

http://nvlabs.github.io/cub/namespacecub.html#abec44bba36037c547e7e84906d0d23aba0fa6cac57b7df2f475a67af053b9371c

于 2013-11-02T17:18:19.090 回答
-1

好吧，这已经5年了。不过，我想回答这个问题，因为我自己花了一些时间来解决这个问题。也许这仍然可以帮助某人......

  1. 关于“临时内存超出范围”：您需要分配两倍于输入数组的共享内存g_idata。

  2. 在我看来，代码需要稍微改变才能工作。我附上了一个工作示例。用 编译nvcc -std=c++11。

https://github.com/fg91/learning_cuda/blob/master/notes/lesson_3/scanHillisSteele.cu

#include <cuda.h>
#include <cuda_runtime.h>
#include <stdio.h>
#include <iostream>

__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];
}

int main() {
  const int ARRAY_SIZE = 10;
  const int ARRAY_BYTES = ARRAY_SIZE * sizeof(int);

  // generate the input array on the host
  int h_in[ARRAY_SIZE]{1, 2, 5, 7, 8, 10, 11, 12, 15, 19};
  int h_out[ARRAY_SIZE];

  // declare GPU memory pointers
  int * d_in;
  int * d_out;

  // allocate GPU memory
  cudaMalloc((void **) &d_in, ARRAY_BYTES);
  cudaMalloc((void **) &d_out, ARRAY_BYTES);

  // transfer the array to the GPU
  cudaMemcpy(d_in, h_in, ARRAY_BYTES, cudaMemcpyHostToDevice);

  // launch the kernel
  scanHillisSteele<<<1, ARRAY_SIZE, 2 * ARRAY_BYTES>>>(d_out, d_in, ARRAY_SIZE);
  cudaDeviceSynchronize();

  // transfer the resulting array to the cpu
  cudaMemcpy(h_out, d_out, ARRAY_BYTES, cudaMemcpyDeviceToHost);

  // print out the input and resulting array
  std::cout << "Input:" << std::endl;
  for (int i = 0; i < ARRAY_SIZE; ++i) {
    std::cout << h_in[i] << " " << std::flush;
  }
  std::cout << std::endl << "Exclusive scan with operation +:" << std::endl;
  for (int i = 0; i < ARRAY_SIZE; ++i) {
    std::cout << h_out[i] << " " << std::flush;
  }
  std::cout << std::endl;

  // free GPU memory allocation
  cudaFree(d_in);
  cudaFree(d_out);

  return 0;
}
于 2018-11-12T12:39:53.437 回答