OpenCL 커널 코드를 사용하여 장치 레지스터 개수 추측

OpenCL 표준에서 계산 장치가 모두 몇 개의 레지스터를 가지고 있는지, 최소한work-item에 몇 개의 레지스터에 사용할 수 있는지를 볼 수 있는 특징 조회를 제공하지 않았다.한 단락의 내장 코드가 레지스터 부족으로 인해 성능이 심각하게 떨어졌는지도 비교적 중요한 요소이기 때문에 저는 현재의 계산 장치가 최소한 모든work-item에 사용할 수 있는 레지스터를 얼마나 분배할 수 있는지 추측하는 비교적 기본적인 방법을 제공합니다.
이 방법의 사고방식은 먼저 네 개의 임시 변수를 정의한 다음에 하나의 대규모 순환 안에서 일정한 규모의 계산을 하는 것이다.그리고 시간을 통계해 낸다.그 다음에 여덟 개의 임시 변수를 정의하고 전자와 같은 횟수의 순환에서 일정한 규모의 계산을 한 다음에 시간을 통계한다.일반적으로 레지스터가 터지지 않거나 Cache 때문에 성능에 큰 영향을 미치지 않는다면 두 가지 소모 시간은 보통 2배 정도이다.만약 후자가 전자보다 2.2배 이상 많다면 우리는 레지스터가 폭발했다고 생각할 수 있다~
이 방법은 일반적인 GPU에 더욱 유용하다.CPU는 L1 Data Cache를 가지고 있기 때문에 레지스터가 부족할 때 컴파일러는 자주 사용하지 않는 데이터를 창고에 넣고 창고는 이때 높은 명중률의 Cache 접근을 얻을 수 있기 때문에 성능에 영향을 주지 않는다.GPU 측에서 레지스터가 충분하지 않을 때 컴파일러는 자주 사용하지 않는 데이터를 VRAM에 직접 저장하고 외부 VRAM에 대한 접근은 비교적 느리기 때문에 임시 변수가 너무 많아서 외부 메모리에 빈번하게 접근하면 전체적인 계산 성능이 크게 떨어진다.물론 현재 많은 GPU에도 L1 Cache가 있지만 공간도 매우 제한되어 있다.그래서 여기에는'맞히기'라는 단어를 쓰는데, 허허~
먼저 4개의 임시 변수에 대한 kernel 코드가 제공됩니다.

__kernel void QueryRegisterCount(__global int *pInOut)
{
    int index = get_global_id(0);

    int i0 = pInOut[(index * 4 + 0) * 4];
    int i1 = pInOut[(index * 4 + 1) * 4];
    int i2 = pInOut[(index * 4 + 2) * 4];
    int i3 = pInOut[(index * 4 + 3) * 4];
    
    for(int i = 0; i < 100000; i++)
    {
        i1 += i0 << 1;
        i2 += i1 << 1;
        i3 += i2 << 1;
        i0 += i3 << 1;
        
        i1 += i0 >> 1;
        i2 += i1 >> 1;
        i3 += i2 >> 1;
        i0 += i3 >> 1;
        
        i1 += i0 >> 2;
        i2 += i1 >> 2;
        i3 += i2 >> 2;
        i0 += i3 >> 2;
        
        i1 += i0 >> 3;
        i2 += i1 >> 3;
        i3 += i2 >> 3;
        i0 += i3 >> 3;
    }
    
    pInOut[(index * 4 + 0) * 4] = i0;
    pInOut[(index * 4 + 1) * 4] = i1;
    pInOut[(index * 4 + 2) * 4] = i2;
    pInOut[(index * 4 + 3) * 4] = i3;
}

임시 변수에 대한 kernel 코드 8개를 더 제공합니다.

__kernel void QueryRegisterCount(__global int *pInOut)
{
    int index = get_global_id(0);

    int i0 = pInOut[(index * 8 + 0) * 4];
    int i1 = pInOut[(index * 8 + 1) * 4];
    int i2 = pInOut[(index * 8 + 2) * 4];
    int i3 = pInOut[(index * 8 + 3) * 4];
    int i4 = pInOut[(index * 8 + 4) * 4];
    int i5 = pInOut[(index * 8 + 5) * 4];
    int i6 = pInOut[(index * 8 + 6) * 4];
    int i7 = pInOut[(index * 8 + 7) * 4];
    
    for(int i = 0; i < 100000; i++)
    {
        i1 += i0 << 1;
        i2 += i1 << 1;
        i3 += i2 << 1;
        i4 += i3 << 1;
        i5 += i4 << 1;
        i6 += i5 << 1;
        i7 += i6 << 1;
        i0 += i7 << 1;
        
        i1 += i0 >> 1;
        i2 += i1 >> 1;
        i3 += i2 >> 1;
        i4 += i3 >> 1;
        i5 += i4 >> 1;
        i6 += i5 >> 1;
        i7 += i6 >> 1;
        i0 += i7 >> 1;
        
        i1 += i0 >> 2;
        i2 += i1 >> 2;
        i3 += i2 >> 2;
        i4 += i3 >> 2;
        i5 += i4 >> 2;
        i6 += i5 >> 2;
        i7 += i6 >> 2;
        i0 += i7 >> 2;
        
        i1 += i0 >> 3;
        i2 += i1 >> 3;
        i3 += i2 >> 3;
        i4 += i3 >> 3;
        i5 += i4 >> 3;
        i6 += i5 >> 3;
        i7 += i6 >> 3;
        i0 += i7 >> 3;
    }
    
    pInOut[(index * 8 + 0) * 4] = i0;
    pInOut[(index * 8 + 1) * 4] = i1;
    pInOut[(index * 8 + 2) * 4] = i2;
    pInOut[(index * 8 + 3) * 4] = i3;
    pInOut[(index * 8 + 4) * 4] = i4;
    pInOut[(index * 8 + 5) * 4] = i5;
    pInOut[(index * 8 + 6) * 4] = i6;
    pInOut[(index * 8 + 7) * 4] = i7;
}

16개, 32개의 임시 변수처럼 유추한다~
그런 다음 호스트측 코드를 입력합니다.

    /** Prepare for running an OpenCL kernel program to get register count */
    
    /*Step 4: Creating command queue associate with the context.*/
    commandQueue = clCreateCommandQueue(context, device, CL_QUEUE_PROFILING_ENABLE, NULL);
    
    /*Step 5: Create program object */
    // Read the kernel code to the buffer
    kernelPath = [[NSBundle mainBundle] pathForResource:@"reg" ofType:@"ocl"];
    aSource = [[NSString stringWithContentsOfFile:kernelPath encoding:NSUTF8StringEncoding error:nil] UTF8String];
    kernelLength = strlen(aSource);
    program = clCreateProgramWithSource(context, 1, &aSource, &kernelLength, NULL);
    
    /*Step 6: Build program. */
    status = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
    
    /*Step 7: Initial inputs and output for the host and create memory objects for the kernel*/
    const size_t memSize = global_work_size[0] * 1024 * 4 * 4;
    cl_int *orgBufer = (cl_int*)malloc(memSize);
    memset(orgBufer, 1, memSize);
    outputMemObj = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR, memSize, orgBufer, NULL);
    
    /*Step 8: Create kernel object */
    kernel = clCreateKernel(program, "QueryRegisterCount", NULL);
    
    /*Step 9: Sets Kernel arguments.*/
    status |= clSetKernelArg(kernel, 0, sizeof(outputMemObj), &outputMemObj);
    
    /*Step 10: Running the kernel.*/
    for(int i = 0; i < 5; i++)
    {
        NSTimeInterval beginTime = [[NSProcessInfo processInfo] systemUptime];
        status |= clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
        clFinish(commandQueue);
        NSTimeInterval endTime = [[NSProcessInfo processInfo] systemUptime];
        
        NSLog(@"Time spent: %f", endTime - beginTime);
    }
    
    free(orgBufer);

    if(status != CL_SUCCESS)
    {
        NSLog(@"Program built failed!");
        return;
    }
    
    clReleaseMemObject(outputMemObj);
    clReleaseProgram(program);
    clReleaseKernel(kernel);
    clReleaseCommandQueue(commandQueue);
    
    clReleaseContext(context);

이상은 OS X에서 개발되었기 때문에 Objective-C 파일로 직접 읽기와 쓰기가 편리합니다.하지만 대부분 C 코드여서 쉽게 읽을 수 있다.
그 중에서 마지막 코드에서 우리는 다섯 번의 순환을 하고 시간을 통계했다.우리는 비교할 때 종종 5번의 집행 시간 중 가장 적은 시간을 골라 비교한다.
2013년 MacBook Air의 Intel HD 5000 테스트 결과:
4개의 임시 변수 사용량: 0.061020초
8개의 임시 변수 소모량: 0.121868초
16개의 임시 변수 소모량: 0.243470초
32개의 임시 변수 소모량:0.719506초
Intel HD Graphics 5000은 최소한 모든 워크-item에 16개의 레지스터를 할당할 수 있다는 것은 명백하다.
실제 응용 프로그램에 사용하려면,kernel 문자열을 동적으로 생성하여 순서대로 실행할 수 있으며, 인접한 두 단락의kernel의 실행 시간이 2.2배가 넘을 때까지 검사할 수 있습니다. 그러면 우리는 종료할 수 있습니다.