OpenCL : 가변 길이의 배열 출력

Question

우리는 GPGPU 코스에 대한 과제를 진행하고 있습니다. 우리는 알고리즘을 선택하여 CPU에서 구현했으며 이제 OpenCL로 변환합니다.

선택한 알고리즘은 모델을 삼각형 집합으로로드하고이를 복셀로 래스터 화합니다. 복셀은 점 데이터의 VBO로 정의됩니다. 그런 다음 기하학 쉐이더를 사용하여이 점들을 삼각형 모양의 보셀로 변환합니다.

따라서 OpenCL 프로그램은 삼각형 목록을 가져와 점의 변수 목록을 출력해야합니다.

가변 길이 배열을 출력하는 것이 문제인 것처럼 보입니다.

우리가 찾은 해결책은 카운터를 원자 적으로 증가시키고 해당 카운터를 출력 배열의 인덱스와 배열의 최종 크기로 사용하는 것입니다. 예외는 ... 두 GPU 모두 원자 연산 확장을 지원하지 않습니다.

이것은 우리가 지금까지 무엇을 가지고 :

#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable 
#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable 
#pragma OPENCL EXTENSION cl_khr_int64_extended_atomics : enable 

#define POS1  i0 * 3 + 0 
#define POS2  i0 * 3 + 1 
#define POS3  i0 * 3 + 2 

void WritePosition(__global float* OutBuffer, uint inIndex, __global float* inPosition) 
{ 
    OutBuffer[ inIndex * 3 ] = inPosition[0]; 
    OutBuffer[ inIndex * 3 + 1] = inPosition[1]; 
    OutBuffer[ inIndex * 3 + 2] = inPosition[2]; 
} 

__kernel void Voxelize( 
    __global float* outPointcloudBuffer, 
    __global float* inTriangleBuffer, 
    __global uint* inoutIndex 
) 
{ 
    size_t i0 = get_global_id(0); 
    size_t i1 = get_local_id(0); 

    WritePosition(outPointcloudBuffer, inIndex[0], &inTriangleBuffer[ i0 ]); 

    //atomic_inc(inoutIndex[0]); 
    inoutIndex[0] = max(inoutIndex[0], i0); 
} 

그리고 이것의 출력이 매우 이상하다. 우리는 매우 작은 모델 (삼각형 12 개, 위치 36 개, 부동 소수점 108 개)을 테스트하고 있으며, 결과는 31, 63 또는 95입니다. 항상 16에서 1을 뺀 것입니다.

어떻게 길이를 구할 수 있습니까? 우리의 가변 길이 출력 배열?

미리 감사드립니다.

Answer 1

I는 다음과 같이 정상적으로 부딪히는 것을 추측 할 :

• 먼저 패스 : 프리미티브 scan (병렬 프리픽스 합계)를 사용하여 GPU에 어레이의 필요한 크기를 계산한다. 위의 링크에는 Apple의 예제 구현이 포함되어 있습니다.
• 검색 알고리즘의 결과를 사용하여 호스트 측에서 필요한 리소스를 할당하십시오. 스캔 알고리즘의 결과는 종종 개별 작업 항목의 결과에 대한 인덱스 힌트로 사용될 수 있습니다.
• 두 번째 패스 (선택 사항) : 세 번째 패스에서 고려해야 할 요소로 배열을 압축합니다.
• 세 번째 패스 : 대상 인덱스와 할당 된 배열을 전달하는 알고리즘을 다시 실행합니다.

위에서 언급 한 세 가지 단계가 모두 구현 된 NVIDIA의 OpenCL 행진 큐브 implementation을 살펴볼 수 있습니다.

최고, Christoph