Tối ưu hóa hiệu suất OpenCL

Question

Tôi đã bắt đầu học OpenCL và hiện tại tôi đang thử kiểm tra xem tôi có thể cải thiện hiệu suất cho thuật toán hoạt hình xương đơn giản như thế nào. Để thực hiện điều này, tôi đã viết một chương trình thực hiện hoạt hình xương từ các đỉnh được tạo ngẫu nhiên và các ma trận chuyển đổi hai lần, một lần với thư viện đại số tuyến tính được tối ưu hóa SSE ở đồng bằng C++ và một khi sử dụng hạt nhân OpenCL của riêng tôi trên GPU (tôi đang thử nghiệm trên một Nvidia GTX 460).

Tôi bắt đầu với một hạt nhân đơn giản trong đó mỗi mục công việc biến đổi chính xác một đỉnh, với tất cả các giá trị được đọc từ bộ nhớ chung. Bởi vì tôi không hài lòng với hiệu suất của hạt nhân này, tôi đã cố gắng tối ưu hóa một chút. hạt nhân hiện tại của tôi trông như thế này:

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)); 
    } 
} 

Bây giờ tôi xử lý một hằng số của đỉnh mỗi công việc mục, và tôi prefetch tất cả các ma trận xương vào bộ nhớ địa phương chỉ một lần cho mỗi công việc mục, mà tôi tin rằng sẽ dẫn để đạt hiệu suất tốt hơn bởi vì các ma trận cho nhiều đỉnh có thể được đọc từ bộ nhớ cục bộ nhanh hơn sau đó. Thật không may, hạt nhân này thực hiện tồi tệ hơn nỗ lực đầu tiên của tôi, và thậm chí tệ hơn so với việc thực hiện chỉ CPU.

Tại sao hiệu suất lại xấu như vậy với tối ưu hóa này?

Nếu nó giúp, đây là cách tôi thực hiện các hạt nhân:

#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; 
} 

Tôi đoán có rất nhiều điều mà có thể được tối ưu hóa, có lẽ Trạm trộn một số toàn cầu khác đọc với nhau, nhưng đầu tiên tôi thực sự muốn để biết tại sao tối ưu hóa đầu tiên này không hiệu quả.

Answer 1

Dường như M EI chuỗi trong Nhóm làm việc đang sao chép 50 số nổi tương tự trước khi bắt đầu tính toán. Điều này sẽ làm ướt băng thông bộ nhớ toàn cầu.

thử này

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

này làm bản sao duy nhất một lần cho mỗi nhóm làm việc.

Answer 2

Bạn đã tìm ra nguyên nhân làm chậm hạt nhân của mình chưa?

Có lẽ tôi sai nhưng tôi nghĩ rằng có tất cả các mục công việc trong nhóm làm việc truy cập cùng một bộ nhớ cục bộ có thể dẫn đến tắc nghẽn.

Answer 3

Hai yếu tố ảnh hưởng đến hiệu suất trong bài tập của bạn.

1) OpenCL phù hợp với C99 std không chứa bất kỳ điều gì về chức năng nội tuyến, tức là trình biên dịch clcc chỉ bỏ qua từ khóa inline và thực hiện cuộc gọi thông thường hoặc hỗ trợ âm thầm nội tuyến. Nhưng nó không phải là bắt buộc để hỗ trợ tính năng đó.

Vì vậy, hãy xác định rõ hơn MultiplyMatrixVector của bạn làm macro tiền xử lý. Mặc dù đây không phải là một vấn đề lớn trong trường hợp của bạn.

2) Bạn đe dọa không chính xác bộ nhớ cục bộ (số LDM).

Mặc dù thời gian trễ của nó nhỏ hơn độ trễ của global memory khi được truy cập đúng cách, local memory chịu sự xung đột của ngân hàng.

Chỉ số đỉnh của bạn được tính bằng stride 100 cho mỗi mục công việc. Số lượng ngân hàng phụ thuộc vào GPU được sử dụng nhưng thường là 16 hoặc 32, i.e. bạn có thể truy cập tới 16 (32) bốn byte LDM biến trong một chu kỳ mà không bị phạt nếu tất cả chúng đều ở các ngân hàng khác nhau. Nếu không, bạn nhận được một bank conflict (khi hai hoặc nhiều chủ đề truy cập vào cùng một ngân hàng) được tuần tự hóa. 100 chủ đề của bạn trong một nhóm làm việc truy cập mảng trong LDM mà không có sự sắp xếp đặc biệt nào về xung đột ngân hàng. Hơn nữa, các phần tử mảng là float16, tức là một phần tử đơn lẻ kéo dài tất cả 16 ngân hàng (hoặc một nửa trong số 32 ngân hàng). Do đó, bạn có xung đột ngân hàng ở mỗi hàng của hàm MultiplyMatrixVector. Sốcummulative mà xung đột ít nhất 16x32 (ở đây 16 là số lượng các phần tử vectơ bạn truy cập và 32 là một kích thước của một nửa sóng hoặc halfwarp).

Giải pháp ở đây không phải là để sao chép mảng đó để LDM, nhưng để phân bổ nó trong máy chủ với CL_MEM_READ_ONLY (mà bạn đã làm) và tuyên bố hạt nhân bằng cách sử dụng __constant specifier cho boneMats tranh cãi. Sau đó, thư viện OpenCL sẽ cấp phát bộ nhớ trong khu vực liên tục bên trong GPU và truy cập vào mảng đó sẽ được nhanh chóng:

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