2

如果元素遵守许多约束，我想使用 Thrust 的流压缩功能 (copy_if) 从向量中提取元素的索引。这些约束之一取决于相邻元素的值（2D 中的 8 和 3D 中的 26）。我的问题是：如何在 Thrust 中获取元素的邻居？

'copy_if' 的仿函数的函数调用运算符基本上如下所示：

__host__ __device__ bool operator()(float x) {
    bool mark = x < 0.0f;
    if (mark) {
        if (left neighbor of x > 1.0f) return false;
        if (right neighbor of x > 1.0f) return false;
        if (top neighbor of x > 1.0f) return false;
        //etc.
    }
    return mark;
}

目前我通过首先启动一个 CUDA 内核（在其中很容易访问邻居）来适当地标记元素来使用一种解决方法。之后，我将标记的元素传递给 Thrust 的 copy_if 以提取标记元素的索引。

我遇到了counting_iterator 作为一种直接使用threadIdx 和blockIdx 来获取已处理元素的索引的替代品。我尝试了下面的解决方案，但是在编译它时，它给了我一个“/usr/include/cuda/thrust/detail/device/cuda/copy_if.inl(151): Error: Unaligned memory accesses not supported”。据我所知，我并没有试图以不对齐的方式访问内存。任何人都知道发生了什么和/或如何解决这个问题？

struct IsEmpty2 {
    float* xi;

    IsEmpty2(float* pXi) { xi = pXi; }

    __host__ __device__ bool operator()(thrust::tuple<float, int> t) {
        bool mark = thrust::get<0>(t) < -0.01f;
        if (mark) {
            int countindex = thrust::get<1>(t);
            if (xi[countindex] > 1.01f) return false;
            //etc.
        }
        return mark;
    }
};


thrust::copy_if(indices.begin(),
                indices.end(),
                thrust::make_zip_iterator(thrust::make_tuple(xi, thrust::counting_iterator<int>())),
                indicesEmptied.begin(),
                IsEmpty2(rawXi));
4

1 回答 1

1

@phoad：您对共享内存是正确的，在我发布了我的回复之后，我感到很震惊，随后认为缓存可能会对我有所帮助。但是你的快速反应打败了我。然而，if 语句在不到 5% 的情况下执行，因此使用共享内存或依赖缓存可能对性能的影响可以忽略不计。

元组仅支持 10 个值，因此这意味着在 3D 情况下我需要元组的元组来处理 26 个值。使用元组和 zip_iterator 已经很麻烦了，所以我将通过这个选项（也是从代码可读性的角度来看）。我通过在设备函数中直接使用 threadIdx.x 等来尝试您的建议，但 Thrust 不喜欢这样。我似乎得到了一些无法解释的结果，有时我最终会出现推力错误。例如，以下程序会生成带有“未指定启动失败”的“thrust::system::system_error”，尽管它首先正确地将“Processing 10”打印到“Processing 41”：

struct printf_functor {
    __host__ __device__ void operator()(int e) {
        printf("Processing %d\n", threadIdx.x);
    }
};

int main() {
    thrust::device_vector<int> dVec(32);
    for (int i = 0; i < 32; ++i)
        dVec[i] = i + 10;

    thrust::for_each(dVec.begin(), dVec.end(), printf_functor());

    return 0;
}

同样适用于打印 blockIdx.x 打印 blockDim.x 但是不会产生错误。我希望有一个干净的解决方案，但我想我坚持使用我目前的解决方案。

于 2012-10-08T10:19:00.970 回答