1

我正在尝试执行以下操作（简化）：请阅读编辑部分！

__shared__ int currentPos = 0;
__global__ myThreadedFunction(float *int, float *out)
{
    // do calculations with in values
    ...

    // now first thread reach this:
    //suspend other threads up here

    out += currentPos;
    for (int i = 0; i < size; ++i)
    {
        *(currentPos++) =  calculation[i];
    }
    currentPos +=  size;

    // now thread is finish, other threads can
    // go on with writing
}

那么如何在写入相同的内存之前挂起线程呢？我不能同时写，因为我不知道每个计算数组的大小（calculation[i] - size）。

我知道有同步线程和线程围栏，但我不知道如何正确使用它们来解决这个问题。

编辑： 我想做的是：

我有 2 个线程（仅作为示例）。每个线程都在一个新数组中使用浮点数*进行计算。

线程 1 计算：{ 1, 3, 2, 4 }

线程 2 计算：{ 3, 2, 5, 6, 3, 4 }

这些数组的大小在计算后是已知的。现在我想在 float *out 中写入这些数组。

如果第一个线程 1 或线程 2 正在写入，对我来说没有必要。输出可能是： * { 1, 3, 2, 4, 3, 2, 5, 6, 3, 4 } 或 { 3, 2, 5, 6, 3, 4, 1, 3, 2, 4} *

那么如何计算输出数组的位置呢？

我不想使用固定的“数组大小”，以便输出为： * { 1, 3, 2, 4, ?, ?, 3, 2, 5, 6, 3, 4 } *

我想我可以为下一个写作位置使用一个共享变量 POSITION。

线程 1 到达写入点（计算新数组后）。线程 1 在共享变量 POSITION 中写入他的数组大小 (4)。

当线程 1 现在将他的临时数组写入输出数组时，线程 2 读取变量 POSITION 并添加他的 tmp。数组大小 (6) 到这个变量，并从线程 1 结束的位置开始写入

如果有线程 3，他还将读取 POSITION，添加他的数组大小并写入输出，线程 2 结束

所以有人有想法吗？

4

2 回答 2

2

从概念上讲，您将如何使用共享数组来存储每个线程的索引来执行并发输出。

__global__ myThreadedFunction(float *int, float *out)
{

    __shared__ index[blockDim.x];//replace the size with an constant
    // do calculations with in values
    ...



    index[tid] = size;// assuming size is the size of the array you output
    //you could do a reduction on this for loop for better performance.
    for(int i = 1; i < blockDim.x; ++i) {
        __syncthreads();
        if(tid == i) {
            index[tid] += index[tid-1];
        }
    }
    int startposition = index[tid] - size; // you want to start at the start, not where the index ends

    //do your output for all threads concurrently where startposition is the first index you output to

}

所以你要做的是分配index[tid]你想要输出的大小，tid线程索引在哪里threadIdx.x，然后向上求和数组（增加索引），最后index[tid]是你的输出数组中从线程0开始的偏移量起始索引。求和可以很容易地使用归约来完成。

于 2013-01-13T21:37:25.887 回答
0

此代码按您的预期工作。它同时读取input[]. 对于每个输入元素 ，它按照存储在 中的顺序size写入次数。sizesizeresultinput[]

请注意，写入过程可能比在 CPU 上执行此操作需要更长的时间。由于您已经知道每个线程要写入的数据大小，您可能希望使用并行前缀和先计算每个线程的写入位置，然后并发写入数据。

有关在代码中使用的更多信息，请参阅Memory Fence Functions 。__threadfence()

#include <thrust/device_vector.h>
#include <thrust/device_ptr.h>

volatile __device__ int count = 0;
volatile __device__ int pos = 0;
__global__ void serial(const float* input, const int N, float* result)
{
    int id = threadIdx.x + blockIdx.x * blockDim.x;

    //parallel part
    int size = (int) input[id];

    //serial output
    for (int i = 0; i < N; i++)
    {
        int localcount = count;
        if (localcount == id)
        {
            int localpos = pos;
            for (int j = 0; j < size; j++)
            {
                result[localpos + j] = (float) j + 1;
            }
            pos = localpos + size;
            count = localcount + 1;
            __threadfence();
        }
        while (count == localcount)
        {
            __syncthreads();
        };

    }
}

int main()
{
    int N = 6;
    thrust::device_vector<float> input(
            thrust::counting_iterator<float>(1),
            thrust::counting_iterator<float>(1) + N);

    thrust::device_vector<float> result(N * (N + 1) / 2);
    serial<<<2, 3>>>(
            thrust::raw_pointer_cast(&input[0]),
            N,
            thrust::raw_pointer_cast(&result[0]));

    thrust::copy(
            result.begin(), result.end(),
            std::ostream_iterator<float>(std::cout, " "));

    return 0;

}

按预期输出：

1 1 2 1 2 3 1 2 3 4 1 2 3 4 5 1 2 3 4 5 6
于 2013-01-13T19:45:27.827 回答