OpenCL 성능 최적화

Question

OpenCL을 배우기 시작했으며 현재 간단한 뼈대 애니메이션 알고리즘의 성능을 얼마나 향상시킬 수 있는지 테스트하려고합니다. 이를 위해 무작위로 생성 된 버텍스와 변환 행렬을 일반 C++로 SSE 최적화 선형 대수 라이브러리에서 한 번, GPU에서 내 자신의 OpenCL 커널을 사용하여 한 번씩 두 번 골격 애니메이션을 수행하는 프로그램을 작성했습니다. 엔비디아 GTX 460).

각 작업 항목이 모든 값을 전역 메모리에서 읽은 정확히 하나의 꼭지점을 변환하는 간단한 커널로 시작했습니다. 이 커널의 성능에 만족하지 않았기 때문에 조금 최적화하려고했습니다.

inline float4 MultiplyMatrixVector(float16 m, float4 v) 
{ 
    return (float4) (
     dot(m.s048C, v), 
     dot(m.s159D, v), 
     dot(m.s26AE, v), 
     dot(m.s37BF, v) 
    ); 
} 


kernel void skelanim(global const float16* boneMats, global const float4* vertices, global const float4* weights, global const uint4* indices, global float4* resVertices) 
{ 
    int gid = get_global_id(0); 
    int lid = get_local_id(0); 

    local float16 lBoneMats[NUM_BONES]; 
    async_work_group_copy(lBoneMats, boneMats, NUM_BONES, 0); 

    barrier(CLK_LOCAL_MEM_FENCE); 

    for (int i = 0 ; i < NUM_VERTICES_PER_WORK_ITEM ; i++) { 
     int vidx = gid*NUM_VERTICES_PER_WORK_ITEM + i; 

     float4 vertex = vertices[vidx]; 
     float4 w = weights[vidx]; 
     uint4 idx = indices[vidx]; 

     resVertices[vidx] = (MultiplyMatrixVector(lBoneMats[idx.x], vertex * w.x) 
       + MultiplyMatrixVector(lBoneMats[idx.y], vertex * w.y) 
       + MultiplyMatrixVector(lBoneMats[idx.z], vertex * w.z) 
       + MultiplyMatrixVector(lBoneMats[idx.w], vertex * w.w)); 
    } 
} 

가 지금은 작업 항목 당 정점의 상수를 처리하고 난 이어질 것이라고 믿었다 각 작업 항목, 한 번만 로컬 메모리에 모든 뼈 매트릭스를 프리 페치 : 내 현재 커널은 다음과 같습니다 이후 여러 개의 정점에 대한 행렬을 더 빠른 로컬 메모리에서 읽을 수 있기 때문에 성능이 향상됩니다. 불행히도,이 커널은 첫 번째 시도보다 성능이 떨어지며 CPU 전용 구현보다 더 나쁩니다.

성능이 그렇게 좋지 않은 이유는 무엇입니까?

는 여기에 도움이 나는 커널 실행 방법 인 경우 : 나는 어쩌면 다른 세계가 함께 읽는 중 일부를 일괄 처리, 최적화 할 수있는 많은 것들이 추측

#define NUM_BONES 50 
#define NUM_VERTICES 30000 
#define NUM_VERTICES_PER_WORK_ITEM 100 
#define NUM_ANIM_REPEAT 1000 

uint64_t PerformOpenCLSkeletalAnimation(Matrix4* boneMats, Vector4* vertices, float* weights, uint32_t* indices, Vector4* resVertices) 
{ 
    File kernelFile("/home/alemariusnexus/test/skelanim.cl"); 

    char opts[256]; 
    sprintf(opts, "-D NUM_VERTICES=%u -D NUM_REPEAT=%u -D NUM_BONES=%u -D NUM_VERTICES_PER_WORK_ITEM=%u", NUM_VERTICES, NUM_ANIM_REPEAT, NUM_BONES, NUM_VERTICES_PER_WORK_ITEM); 

    cl_program prog = BuildOpenCLProgram(kernelFile, opts); 

    cl_kernel kernel = clCreateKernel(prog, "skelanim", NULL); 

    cl_mem boneMatBuf = clCreateBuffer(ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, NUM_BONES*sizeof(Matrix4), boneMats, NULL); 
    cl_mem vertexBuf = clCreateBuffer(ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, NUM_VERTICES*sizeof(Vector4), vertices, NULL); 
    cl_mem weightBuf = clCreateBuffer(ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, NUM_VERTICES*4*sizeof(float), weights, NULL); 
    cl_mem indexBuf = clCreateBuffer(ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, NUM_VERTICES*4*sizeof(uint32_t), indices, NULL); 
    cl_mem resVertexBuf = clCreateBuffer(ctx, CL_MEM_WRITE_ONLY | CL_MEM_ALLOC_HOST_PTR, NUM_VERTICES*sizeof(Vector4), NULL, NULL); 

    uint64_t s, e; 
    s = GetTickcount(); 

    clSetKernelArg(kernel, 0, sizeof(cl_mem), &boneMatBuf); 
    clSetKernelArg(kernel, 1, sizeof(cl_mem), &vertexBuf); 
    clSetKernelArg(kernel, 2, sizeof(cl_mem), &weightBuf); 
    clSetKernelArg(kernel, 3, sizeof(cl_mem), &indexBuf); 
    clSetKernelArg(kernel, 4, sizeof(cl_mem), &resVertexBuf); 

    size_t globalWorkSize[] = { NUM_VERTICES/NUM_VERTICES_PER_WORK_ITEM }; 
    size_t localWorkSize[] = { NUM_BONES }; 

    for (size_t i = 0 ; i < NUM_ANIM_REPEAT ; i++) { 
     clEnqueueNDRangeKernel(cq, kernel, 1, NULL, globalWorkSize, localWorkSize, 0, NULL, NULL); 
    } 

    clEnqueueReadBuffer(cq, resVertexBuf, CL_TRUE, 0, NUM_VERTICES*sizeof(Vector4), resVertices, 0, NULL, NULL); 

    e = GetTickcount(); 

    return e-s; 
} 

을하지만, 처음에는 정말 싶습니다 이 첫 번째 최적화가 작동하지 않는 이유를 알아야합니다.

Answer 1

작업 그룹의 각 스레드가 계산을 시작하기 전에 동일한 50 개의 부동 소수점을 복사하고있는 것처럼 보입니다. 그러면 전역 메모리 대역폭이 포화 상태가됩니다.

는이

if (lid == 0) 
{ 
    async_work_group_copy(lBoneMats, boneMats, NUM_BONES, 0); 
} 

이 작업 그룹 당 한 번만 복사를 수행하려고합니다.

Answer 2

커널 속도가 느린 이유를 알고 있습니까?

아마 내가 틀렸지 만 동일한 로컬 메모리에 액세스하는 작업 그룹 내의 모든 작업 항목이 병목 현상을 일으킬 수 있다고 생각합니다.

Answer 3

두 가지 작업이 수행 능력에 영향을줍니다.

1) OpenCL은 즉 CLCC 컴파일러 인라인 함수에 대해 아무것도 포함하지 않는 C99 표준에 부합하거나 단지 inline 키워드를 무시하고 일반 전화를합니까, 아니면 자동으로 인라인을 지원합니다. 그러나 해당 기능을 지원하는 것은 아닙니다.

그래서 MultiplyMatrixVector을 사전 처리기 매크로로 정의하는 것이 좋습니다. 이것은 귀하의 경우에 중요한 문제는 아니지만.

2) 로컬 메모리 (LDM)를 잘못 사용했습니다.

대기 시간이 global memory의 액세스 시간보다 짧지 만 local memory은 은행 갈등의 대상이됩니다.

정점 인덱스는 작업 항목 당 100 걸음으로 계산됩니다. 뱅크 수는 사용중인 GPU에 따라 다르지만 대개 16 또는 32입니다.이자형. 모든 뱅크가 서로 다른 뱅크에있는 경우 페널티없이 한 사이클에서 최대 16 (32) 바이트의 4 바이트 LDM 변수에 액세스 할 수 있습니다. 그렇지 않으면 bank conflict (두 개 이상의 스레드가 동일한 뱅크에 액세스하는 경우)이 직렬화됩니다. 작업 그룹의 100 개의 스레드가 LDM의 배열에 액세스하며 은행 충돌에 대해서는 특별한 조치가 필요하지 않습니다. 또한, 배열 요소는 float16, 즉 단일 요소가 16 개의 모든 뱅크 (또는 32 개의 뱅크 중 절반)에 걸쳐있다. 따라서 MultiplyMatrixVector 함수의 각 행에 은행 갈등이 있습니다. 누적 된 degree은 적어도 16x32 (여기서 16은 액세스하는 벡터 요소의 수이고 32는 반 파장의 반 파장 또는 반 파장의 크기 임)와 충돌합니다.

솔루션은 여기에 LDM에 해당 배열을 복사하는 것이 아니라 CL_MEM_READ_ONLY와 호스트에 할당 (이미 않았다) 및 boneMats 인수 __constant 지정자를 사용하여 커널을 선언하지 않는 것입니다. 그런 다음 OpenCL 라이브러리는 GPU 내부의 일정한 영역에 메모리를 할당 할 것이고, 그 배열에 액세스 빠른 것 :

kernel void skelanim(__constant const float16* boneMats, 
        global const float4* vertices, 
        global const float4* weights, 
        global const uint4* indices, 
        global float4* resVertices)