이웃에 대한 추력 접근

Question

요소가 여러 제약 조건을 준수하면 벡터에서 요소의 색인을 추출하는 데 Thrust의 스트림 압축 기능 (copy_if)을 사용하고 싶습니다. 이러한 제한 중 하나는 인접 요소의 값에 따라 다릅니다 (2D는 8, 3D는 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 커널을 실행하여 주위에 작업-A를 사용 요소를 적절하게 표시 할 수 있습니다. 그 후, 표시된 요소를 Thrust의 copy_if에 전달하여 표시된 요소의 색인을 추출합니다.

---

I는 직접 처리 요소의 인덱스를 취득하고 threadIdx blockIdx를 사용하는 대신에 일종의 counting_iterator 우연히. 아래의 솔루션을 시도했지만 컴파일 할 때 "/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)); 

Answer 1

@phoad : 당신은 공유 된 mem에 대해 옳았습니다. 이미 회신을 게시 한 후 캐시가 아마도 도움이 될 것이라고 생각한 것 같습니다. 그러나 당신은 당신의 빠른 반응으로 나를 때렸습니다. 그러나 if 문은 모든 경우의 5 % 미만에서 실행되므로 공유 mem을 사용하거나 캐시를 사용하면 성능에 거의 영향을 미치지 않습니다.

튜플은 10 개의 값만 지원하므로 3D 경우에는 26 개의 값에 대해 튜플의 튜플이 필요합니다. 튜플과 zip_iterator를 사용한 작업은 이미 꽤 번거롭기 때문에이 옵션을 전달할 것입니다 (코드 가독성 관점에서도). 나는 당신의 제안을 장치 기능에서 threadIdx.x 등을 사용하여 직접 시도했지만 추력은 그것을 좋아하지 않습니다. 나는 설명 할 수없는 결과를 얻는 것 같고 때로는 추력 오류로 끝납니다. 먼저 올바르게 "처리 10"를 "가공 (41)"를 출력하지만, 예를 들면 다음과 같은 프로그램은,에 '불특정 진입 실패'와 '추력 :: 시스템 :: SYSTEM_ERROR'를 생성

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를 인쇄하면 오류가 발생하지 않습니다. 나는 깨끗한 솔루션을 원했지만, 현재의 해결 방법을 고수하고있는 것 같아요.