CUDA

Question

를 사용하여 배열의 크기를 알 수없는 지역의 최대 값을 찾기

내가 [45,21,764,234,7,0,12,55,...]

그럼 난으로 배열 A에 지역의 위치를 ​​나타냅니다 다른 배열 B[4000]이 모두 다른 숫자가 포함 된 값의 배열 A[4000]을 말해봐 지역의 일부인 경우 1, 그렇지 않은 경우 0입니다. 1's이 서로 인접 해 동일한 지역에 속해 있다는 것을 의미합니다. 서로 인접하지 않으면 (사이에 0이 있음) 서로 다른 지역의 일부입니다.

ex. B = [1,1,1,0,1,1,0,0...] 내가 first three numbers in array A의 영역에서 최대 값을 찾고, 5th and 6th numbers in array A, etc. 의 최대 값을 찾으려면 B으로 표시된 각 영역에서 최대 값이 A 인 배열 C[4000]을 생성 할 수 있다는 것을 의미합니다. 지역의 일부가 아닌 영역에 0.

그래서이 경우 C = [764,764,764,0,7,7,0,0...]

는 0 to 2,000 regions 어디서나있을 수 있고, 지역의 길이는 2 to 4,000 numbers long에 이르기까지 다양 할 수 있습니다. 얼마나 많은 지역이 있는지 또는 지역의 크기가 다른지 미리 알 수 없습니다.

나는이 결과를 얻을 수있는 CUDA 커널을 만들려고 노력해 왔습니다. 실제로는 이미지에 사용되므로 가능한 한 빨리 처리해야합니다. 이는 단순한 예입니다. 축소를 사용하는 것과 같은 나의 모든 아이디어는 하나의 영역 만있는 경우에만 4000 배열 번호가 모두 A 인 경우에만 작동합니다. 그러나 배열에서 여러 지역이 1에서 3996 공백 (0's)으로 분리되어 있기 때문에 여기에서 감산을 사용할 수 있다고 생각하지 않습니다. 축소하면 분리 된 지역의 느슨한 궤도가 나을 것입니다. 또는, 커널은 너무 많은 루프를 가지고 있으며, 거기에 문이 분명히 코드도 공유 메모리와 느린, 정말 CUDA의 병렬 특성을 활용하지 않는 빠른 같은

int intR = 0; 
while(B[blockIdx.x * blockDim.x + threadIdx.x + intR] > 0){ 
    intMaxR = intMaxR < A[blockIdx.x * blockDim.x + threadIdx.x + intR] ? A[blockIdx.x * blockDim.x + threadIdx.x + intR] : intMaxR; 
    intR++; 
} 

int intL = 0; 
while(B[blockIdx.x * blockDim.x + threadIdx.x - intL] > 0){ 
    intMaxL = intMaxL < A[blockIdx.x * blockDim.x + threadIdx.x - intL] ? A[blockIdx.x * blockDim.x + threadIdx.x + intL] : intMaxL; 
    intL++; 
} 

intMax = intMaxR > intMaxL ? intMaxR : intMaxL; 

for(int i = 0; i < intR; i++){ 
    C[blockIdx.x * blockDim.x + threadIdx.x + i] = intMax; 
} 
for(int i = 0; i < intL; i++){ 
    C[blockIdx.x * blockDim.x + threadIdx.x - i] = intMax; 
} 

로 할 수있는 경우. CUDA에서 이것이 어떻게 효율적으로 수행 될 수 있는지에 대한 아이디어가 있습니까?

미리 감사드립니다.

Answer 1

한 가지 가능한 접근 방법은 thrust을 사용하는 것입니다.

가능한 시퀀스는 다음과 같이 될 것이다 :

1. 사용 thrust::reduce_by_key 각 범위의 최대 값을 생성한다.
2. 각 범위의 시작을 윤곽으로 나타 내기 위해 thrust :: adjacent_difference를 사용하십시오.
3. 단계 2의 결과에 대한 포괄적 인 스캔을 사용하여 수집 색인, 즉 감소 된 값을 선택하는 데 사용될 색인을 생성하십시오 1)는 출력 벡터의 각 위치로 갈 것입니다.
4. thrust::gather_if을 사용하면 3 단계에서 생성 한 gather 인덱스를 사용하여 출력 벡터에 적절한 값 (B 벡터에 1이있는 위치)에 축소 값을 선택적으로 배치 할 수 있습니다.대한
   1. reduce_by_key이 감소 생성하는 값 (최대 값) : 예에 대한

      #include <iostream> 
      #include <thrust/device_vector.h> 
      #include <thrust/adjacent_difference.h> 
      #include <thrust/reduce.h> 
      #include <thrust/copy.h> 
      #include <thrust/transform_scan.h> 
      #include <thrust/iterator/discard_iterator.h> 
      #include <thrust/iterator/transform_iterator.h> 
      #include <thrust/functional.h> 
      
      #define DSIZE 8 
      
      template <typename T> 
      struct abs_val : public thrust::unary_function<T, T> 
      { 
          __host__ __device__ 
          T operator()(const T& x) const 
          { 
          if (x<0) return -x; 
          else return x; 
          } 
      }; 
      
      template <typename T> 
      struct subtr : public thrust::unary_function<T, T> 
      { 
          const T val; 
          subtr(T _val): val(_val) {} 
          __host__ __device__ 
          T operator()(const T& x) const 
          { 
          return x-val; 
          } 
      }; 
      
      int main(){ 
      
          int A[DSIZE] = {45,21,764,234,7,0,12,55}; 
          int B[DSIZE] = {1,1,1,0,1,1,0,0}; 
          thrust::device_vector<int> dA(A, A+DSIZE); 
          thrust::device_vector<int> dB(B, B+DSIZE); 
          thrust::device_vector<int> dRed(DSIZE); 
          thrust::device_vector<int> diffB(DSIZE); 
          thrust::device_vector<int> dRes(DSIZE); 
      
          thrust::reduce_by_key(dB.begin(), dB.end(), dA.begin(), thrust::make_discard_iterator(), dRed.begin(), thrust::equal_to<int>(), thrust::maximum<int>()); 
          thrust::adjacent_difference(dB.begin(), dB.end(), diffB.begin()); 
          thrust::transform_inclusive_scan(diffB.begin(), diffB.end(), diffB.begin(), abs_val<int>(), thrust::plus<int>()); 
          thrust::gather_if(thrust::make_transform_iterator(diffB.begin(), subtr<int>(B[0])), thrust::make_transform_iterator(diffB.end(), subtr<int>(B[0])), dB.begin(), dRed.begin(), dRes.begin()); 
          thrust::copy(dRes.begin(), dRes.end(), std::ostream_iterator<int>(std::cout, " ")); 
          std::cout << std::endl; 
          return 0; 
      } 
      

      주 :

   여기에 완벽하게 작동 코드가 예처럼 A와 B 벡터를 사용하여,이 시연입니다 각 연속 0 시퀀스 또는 1 시퀀스가 ​​B에 있습니다. 실제로는 개의 시퀀스에 대해 최대 값만 필요합니다. gather_if 함수를 통해 0 시퀀스 을 무시합니다.

5. 2 단계의 벡터 결과를 transform_iterator 처리하여 B 벡터의 값을 각각 빼서 1 시퀀스 또는 0 시퀀스로 시작하는 가능성을 허용합니다. 색인을 수집하십시오.
6. adjacent_difference 연산은 새로운 시퀀스의 시작 부분을 나타내는 을 1 또는 -1로 만듭니다. abs_val 펑 터와 함께 transform_inclusive_scan 변형을 사용하여 스캔 목적 (즉, 수집 색인 생성)을 위해 균등하게 처리합니다.

7. 위의 코드는 다음과 같이 원하는 C 출력 벡터와 일치하는 결과를 생성해야합니다

   $ nvcc -arch=sm_20 -o t53 t53.cu 
   $ ./t53 
   764 764 764 0 7 7 0 0 
   $ 
   

우리는 더 위의 코드를 단순화하기 위해 thrust::placeholders을 사용할 수 있습니다, 여분의 펑터 정의에 대한 필요성을 제거 :

#include <iostream> 
#include <thrust/device_vector.h> 
#include <thrust/adjacent_difference.h> 
#include <thrust/reduce.h> 
#include <thrust/copy.h> 
#include <thrust/transform_scan.h> 
#include <thrust/iterator/discard_iterator.h> 
#include <thrust/iterator/transform_iterator.h> 
#include <thrust/functional.h> 

#define DSIZE 2000000 
using namespace thrust::placeholders; 

typedef int mytype; 

int main(){ 

    mytype *A = (mytype *)malloc(DSIZE*sizeof(mytype)); 
    int *B = (int *)malloc(DSIZE*sizeof(int)); 
    for (int i = 0; i < DSIZE; i++){ 
    A[i] = (rand()/(float)RAND_MAX)*10.0f; 
    B[i] = rand()%2;} 
    thrust::device_vector<mytype> dA(A, A+DSIZE); 
    thrust::device_vector<int> dB(B, B+DSIZE); 
    thrust::device_vector<mytype> dRed(DSIZE); 
    thrust::device_vector<int> diffB(DSIZE); 
    thrust::device_vector<mytype> dRes(DSIZE); 

    cudaEvent_t start, stop; 
    cudaEventCreate(&start); 
    cudaEventCreate(&stop); 
    cudaEventRecord(start); 
    thrust::reduce_by_key(dB.begin(), dB.end(), dA.begin(), thrust::make_discard_iterator(), dRed.begin(), thrust::equal_to<mytype>(), thrust::maximum<mytype>()); 
    thrust::adjacent_difference(dB.begin(), dB.end(), diffB.begin()); 
    thrust::transform_inclusive_scan(diffB.begin(), diffB.end(), diffB.begin(), _1*_1, thrust::plus<int>()); 
    thrust::gather_if(thrust::make_transform_iterator(diffB.begin(), _1 - B[0]), thrust::make_transform_iterator(diffB.end(), _1 - B[0]), dB.begin(), dRed.begin(), dRes.begin()); 
    cudaEventRecord(stop); 
    cudaEventSynchronize(stop); 
    float et; 
    cudaEventElapsedTime(&et, start, stop); 
    std::cout<< "elapsed time: " << et << "ms " << std::endl; 
    thrust::copy(dRes.begin(), dRes.begin()+10, std::ostream_iterator<mytype>(std::cout, " ")); 
    std::cout << std::endl; 
    return 0; 
} 

은 (I은 또한 더 큰 크기의 데이터 세트의 생성뿐만 아니라, 기본적인 타이밍 APPA를 포함하는 상기 자리 코드를 수정 한)