AMD OpenCL 예 읽 기 노트 시리즈 의 Binary Search

우선 더 잘 이해 하기 위해 서 첨부 된 문 서 를 대충 넘 겨 보 세 요.이 사례 의 주요 사상 은 빅 데 이 터 를 트 리 축소 와 유사 한 형식 으로 검색 범 위 를 최종 생 성 결과 로 축소 하 는 것 이다.
소개 하 다.
         이 예 는 정렬 된 배열 에서 특정한 값 을 찾 습 니 다.이 값 이 배열 에 없 으 면 보고 하 세 요.이분법 검색 과 달리 매번 검색 범 위 를 반 으로 줄 입 니 다. 검색 범 위 를 N 부 로 나 누 는 것 을 설계 하 였 습 니 다. 우 리 는 그 를 N 로 검색 알고리즘 이 라 고 부 릅 니 다.이분법
log 는 2 를 밑 으로 하 는 복잡 도 이 고 N 로 검색 은 log 를 밑 으로 한다.
디 테 일 구현
 
 
이 특별한 실현 방식 은 배열 이 256 의 배수 가 되 어야 한다.100000 개의 질서 있 는 데이터 에서 원 하 는 값 을 검색 하 는 것 을 고려 하여 우선 이 배열 을 10000 크기 의 데이터 그룹 으로 나 눕 니 다.우 리 는 먼저 검색 할 데이터 가 어느 세그먼트 에 있 는 지 판단 한 다음 에 이 단락 을 10 개의 데이터 (데 이 터 량 은 1000 개의 구성원 / 각 그룹) 로 나 누 었 다.이렇게 하면 우 리 는 우리 의 검색 범 위 를 점차 줄 일 수 있다.
      예 를 들 어, 우 리 는 우리 의 입력 배열 이 2, 4 라 고 가정 합 니 다................................................................
       우선 모든 스 레 드 의 데이터 세그먼트 는:
       Thread 0:     2....2*10^4,     [0,10^4]
       Thread 1:    2*10^4.+2  ..... 3*10^4     [10^4,2*10^4]
       Thread 2:  3*10^4+2……4*10^4     [2*10^4,3*10^4]
        등등
       이 위의 단락 은 사실 매우 보기 좋 습 니 다. 10 ^ 4 를 1 로 보면 됩 니 다. 즉, 끊 은 후에 그룹 크기 를 곱 하 는 것 입 니 다. [] 는 상하 계 입 니 다.
        분명히 42 는 첫 번 째 스 레 드 의 상하 계 에 만 있 기 때문에 첫 번 째 스 레 드 만 상하 계 를 출력 버퍼 에 기록 합 니 다. 42 와 상하 계 가 모두 다 르 기 때문에 세 번 째 요소 에 0 을 쓰 십시오. 즉, 이 단계 가 완 료 된 출력 결 과 는 0, 10 ^ 4, 0 입 니 다.
        이와 유사 하 게 다음 출력 결 과 는 0, 10 ^ 3, 0 입 니 다. 다음 출력 결 과 는 0, 10 ^ 2, 0 입 니 다. 현재 검색 범 위 는 2...............................................
        Thread0:2...20
        Thread1:22....40
Thread2:42....60
         이 때 스 레 드 2 출력 만 있 고 세 번 째 요소 의 값 은 1 입 니 다. 검색 이 끝 났 음 을 의미 합 니 다. 43 을 검색 하면 계속 규정 하고 모든 스 레 드 는 데 이 터 를 처리 합 니 다.
         예 를 들 어 42, 44, 46...............................................................................
Binary Search kernels. cl 중 세 가지 실현:
1) binary Search: 본 고 에서 기술 한 알고리즘
2) binary Search mulkys: 여러 값 찾기
3) binarysearch mulkys Concurrent: 여러 값 을 동시에 찾 습 니 다.
우선 binary Search 커 널:

__kernel void
binarySearch(        __global uint4 * outputArray,
             __const __global uint2  * sortedArray, 
             const   unsigned int findMe)
{
    unsigned int tid = get_global_id(0);

    /* Then we find the elements  for this thread */
    uint2 element = sortedArray[tid];

    /* If the element to be found does not lie between them, then nothing left to do in this thread */
    if( (element.x > findMe) || (element.y < findMe)) 
    {
        return;
    }
    else
    {
        /* However, if the element does lie between the lower and upper bounds of this thread's searchspace
         * we need to narrow down the search further in this search space 
         */
 
        /* The search space for this thread is marked in the output as being the total search space for the next pass */
        outputArray[0].x = tid;
        outputArray[0].w = 1;

    }
}

      입력 량 sortedArray 는 세그먼트 뒤의 상하 계 정보 입 니 다. findMe 는 당연히 찾 아야 할 값 입 니 다. 출력 량 outputArray 는 스 레 드 번 호 를 이용 하여 표 시 됩 니 다. 이 코드 는 비교적 간단 합 니 다. 찾 을 때 출력 에 필요 한 정 보 를 기록 하 는 것 입 니 다. 그렇지 않 으 면 바로 되 돌아 갑 니 다.
      호스트 엔 드 코드 를 보 겠 습 니 다:
       우선 setupBinary Search 에서 테스트 데 이 터 를 초기 화 합 니 다. 핵심 코드 는:
      

    input[0] = 0;
    for(cl_uint i = 1; i < length; i++)
    {
        input[i] = input[i - 1] + (cl_uint) ((max * rand()) / (float)RAND_MAX);
    }

       임 의 수 를 통 해 초기 화 됩 니 다. 커 널 의 실행 은 주로 runCLKernels 함수 에 있 습 니 다. 함 수 는 선택 한 장치 의 GroupSize 와 비교 한 다음 최종 분할 수 subdivsize 를 확인 합 니 다. 입력 량 상하 계 의 코드 는 다음 과 같 습 니 다.

        cl_uint *in=NULL;

        // Set input data
        status = mapBuffer( inputBuffer, in, inlength, CL_MAP_WRITE );
        CHECK_ERROR(status, SDK_SUCCESS, "Failed to map device buffer.(inputBuffer)");

        for(cl_uint i=0 ; i<numSubdivisions; i++)
        {
            int indexa = i*subdivSize;
            int indexb = (i+1)*subdivSize-1;
            in[2*i] = input[indexa];
            in[2*i+1] = input[indexb];
        }

        status = unmapBuffer( inputBuffer, in);
        CHECK_ERROR(status, SDK_SUCCESS, "Failed to unmap device buffer.(inputBuffer)");

      OpenCL 의 메모리 맵 기술 을 사용 합 니 다. 이 방식 은 호스트 에 장치 메모 리 를 매 핑 하여 호스트 의 지침 을 받 고 포인터 로 이 메모리 영역 을 조작 할 수 있 습 니 다. 그리고 커 널 매개 변수 전달 을 보십시오.

    status = clSetKernelArg(
                 kernel,
                 0,
                 sizeof(cl_mem),
                 (void *)&outputBuffer);
    CHECK_OPENCL_ERROR(status, "clSetKernelArg 0(OutputBuffer) failed.");

        /*
        * Second argument is input buffer
        */
        status = clSetKernelArg(
                     kernel,
                     1,
                     sizeof(cl_mem),
                     (void *)&inputBuffer);
        CHECK_OPENCL_ERROR(status, "clSetKernelArg 1(inputBuffer) failed.");

        /*
        * Third is the element we are looking for
        */
        status = clSetKernelArg(
                     kernel,
                     2,
                     sizeof(cl_uint),
                     (void *)&findMe);
        CHECK_OPENCL_ERROR(status, "clSetKernelArg 2(findMe) failed.");

       이 두 부분 을 특별히 나 눈 이 유 는 첫 번 째 매개 변 수 는 실제 적 으로 공용 출력 매개 변수 이기 때문에 한 번 만 들 어 와 야 하고 두 번 째 와 매번 변화 가 필요 하 다 고 생각 합 니 다. 그리고 세 번 째 매개 변 수 는 순환 밖 에 두 는 것 이 좋 을 것 이 라 고 생각 합 니 다. 설정 이 완료 되면 이 커 널 의 입 대 를 통 해 처리 해 야 합 니 다.

        cl_event ndrEvt;
        status = clEnqueueNDRangeKernel(commandQueue,
                                        kernel,
                                        1,
                                        NULL,
                                        globalThreads,
                                        localThreads,
                                        0,
                                        NULL,
                                        &ndrEvt);

         매번 처리 한 후에 가장 중요 한 것 은 검색 범위 의 축소 이다. 축소 과정 은 다음 과 같다.

        globalLowerBound = output[0]*subdivSize;
        globalUpperBound = globalLowerBound+subdivSize-1;
        subdivSize = (globalUpperBound - globalLowerBound + 1)/numSubdivisions;

         output [0] 커 널 에서 볼 수 있 듯 이 범위 내 스 레 드 를 검색 할 때마다 표 시 를 하기 때문에 다음 스 레 드 의 상하 계 와 분 단 량 으로 쉽게 전환 할 수 있 습 니 다. 전체적인 운행 과정 은 아래 의 의사 코드 로 설명 할 수 있 습 니 다.
     while(subdivSize>1 && output[3]!=0)
    {
  새 상하 경계 값 설정
          커 널 매개 변수 업데이트
         출력 재 설정 [3] = 0
          실행 커 널
         읽 기 출력 값
         검색 상하 계 및 분할 크기 업데이트
      }
      이러한 최종 목적 은 당연히 우리 가 필요 로 하 는 값 을 찾 는 것 이다. 실제로 위 에서 도 최소 비교 범 위 를 최종 적 으로 확 정 했 을 뿐 이지 만 이렇게 하면 찾기 에 있어 서도 많이 간소화 되 었 다. 구체 적 으로 는 다음 과 같다.

    for(cl_uint i=globalLowerBound; i<= globalUpperBound; i++)
    {
        if(input[i] == findMe)
        {
            elementIndex = i;
            globalLowerBound = i;
            globalUpperBound = i+1;
            isElementFound = 1;
            break;
        }
    }

     하지만 유일 하 게 이해 하지 못 할 때 output 의 Map 조작 이 이상 하 다. while 순환 에서 unmap 를 해 야 할 것 같다.
     마지막 으로 우 리 는 다른 두 개의 커 널 을 다시 보 자.
     

__kernel void
binarySearch_mulkeys(__global int *keys,
                    __global uint *input,
                    const unsigned int numKeys,
                    __global int *output)
{
    int gid = get_global_id(0);
    int lBound = gid * 256;
    int uBound = lBound + 255;
    //int uBound
    

    for(int i = 0; i < numKeys; i++)
    {
        
        if(keys[i] >= input[lBound] && keys[i] <=input[uBound])
            output[i]=lBound;       
        
    }
    
}

     위의 설명 이 있 으 면 이 예 는 많이 이해 할 수 있 습 니 다. 찾 은 값 을 여러 개 로 바 꾸 었 을 뿐 입 니 다. 출력 대응 도 여러 개 로 바 뀌 었 을 뿐 입 니 다. 병렬 버 전 을 살 펴 보 겠 습 니 다.

__kernel void
binarySearch_mulkeysConcurrent(__global uint *keys,
                    __global uint *input,
                    const unsigned int inputSize, // number of input elements
                    const unsigned int numSubdivisions,
                    __global int *output)
{
    int lBound = (get_global_id(0) % numSubdivisions) * (inputSize / numSubdivisions);
    int uBound = lBound + inputSize / numSubdivisions;
    int myKey = keys[get_global_id(0) / numSubdivisions];
    int mid;
    while(uBound >= lBound)
    {
        mid = (lBound + uBound) / 2;
        if(input[mid] == myKey)
        {
            output[get_global_id(0) / numSubdivisions] = mid;
            return;
        }
        else if(input[mid] > myKey)
            uBound = mid - 1;
        else
            lBound = mid + 1;
    }
}

       많이 복잡 해진 것 같 습 니 다. 전체적으로 보면 그 while 는 2 분 에 찾 는 것 입 니 다. 그 위 에 있 는 매개 변 수 는 무엇 에 쓰 입 니까? 우 리 는 하나씩 봅 시다.
      lBound 는 말 그대로 하계 가 분명 합 니 다. input Size / numSubdivisions. 이것 은 분 단 된 세그먼트 크기 입 니 다. get global id (0)% numSubdivisions 는 0 ~ numSubdivisions 의 수 를 얻 었 습 니 다. 우리 가 합 쳐 보면 0 ~ numSubdivision 의 수 에 분 단 된 세그먼트 크기 를 곱 하면 다음 단락 의 하계 가 분명 합 니 다.
      uBound 라 는 간단 한 것 은 lBound 를 바탕 으로 분 단 된 세그먼트 길 이 를 추가 하 는 것 입 니 다. my Key 는 찾 으 려 는 keys 에서 이 세그먼트 데 이 터 를 얻 으 려 면 어떤 키 워드 를 찾 아야 합 니까?
      이렇게 많은 병렬 법 을 말 했 는데, 실제로는 0 ~ 100, 100 ~ 200, 200 ~ 300 세 범위 의 자연 수 는 각각 10, 120, 256 이다. 그러면 numSubdivisions 를 3 으로 설정 하고, 첫 번 째 스 레 드 는 계산 범 위 를 0 ~ 100 으로 계산 할 수 있 으 며, 찾 아야 할 값 은 10 이 고, 두 번 째 스 레 드 가 찾 아야 할 값 은 120 이다.이 단락 에 대응 하 는 key 값 입 니 다. 이해 하 는 것 이 맞 는 지 모 르 겠 습 니 다. 그렇지 않 으 면 다음 에 보 세 요.