Thrust - access to neighbors

Question

I would like to use the Thrust stream compression function (copy_if) to extract element indices from a vector if the elements adhere to a number of restrictions. One of these restrictions depends on the values ​​of neighboring elements (8 in 2D and 26 in 3D). My question is: how can I get the neighbors of an element in Thrust?

The functor function call operator for 'copy_if' basically looks like this:

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

I am currently using a workaround, starting up the CUDA kernel first (which makes it easy to access neighbors) to mark the elements accordingly. After that, I pass the marked items to Thrust copy_if to overtake the indices of the marked items.

---

counting_iterator threadIdx blockIdx . , "/usr/include/cuda/thrust/detail/device/cuda/copy_if.inl(151): : Unaligned ". , . - , / ?

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

Answer 1

@phoad: mem, , , , , , . . if-statement 5% , shared mem, , , .

10 , , 26 . zip_iterator , ( ). , threadIdx.x .. , Thrust . , , Thrust. , , "thrust:: system:: 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 . , , , .