推力 - 访问邻居(Thrust - accessing neighbors)

我想用推力的流压缩功能（copy_if），用于蒸馏从向量元素的索引，如果要素坚持一些限制。一个的这些限制取决于相邻元件（8在二维和三维中的26）的值。我的问题是：我怎么能获得推力元素的邻居？

函子的“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核心（其中很容易访问邻居）以正确标记的元素。在那之后，我通过标记元素推力的copy_if提炼标记元素的索引。

我跨越counting_iterator来作为一种替代的直接使用threadIdx和blockIdx获取处理元素的索引。我尝试了以下解决方案，但在编译的时候，它给了我一个“/usr/include/cuda/thrust/detail/device/cuda/copy_if.inl(151）：错误：未对齐的内存访问，不支持”。据我所知，我并不想在不对齐的方式来访问内存。任何人都知道发生了什么和/或如何解决这一问题？

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));

@phoad：你说得对共享的纪念品，它打了我，我已经张贴了我的答复，后来以为缓存可能会帮助我。但是，你有你的快速反应打我。 if语句不过是在所有情况下都小于5％执行，因此，无论是使用共享MEM或依赖缓存可能对性能的影响微乎其微。

元组只支持10个值，因此这将意味着我需要的元组在3D情况下，26个值的元组。与元组和zip_iterator工作已经很麻烦，所以我会通过此选项（也可从代码的可读性立场来看）。我想你的建议在设备功能直接使用threadIdx.x等，但推力不喜欢这样。我似乎得到一些无法解释的结果，有时我结束了一个推力错误。例如下面的程序产生一个“推力::系统:: SYSTEM_ERROR”与“未指定的发射失败”，虽然它首先正确打印“处理10”到“处理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产生任何错误。我希望有一个干净的解决方案，但我想我坚持我目前的工作，它的解决方案。