오픈 CL : 동적 메모리 할당, 더 나은 동시에 유휴 작업 항목 또는 서면을 사용하는 것입니다

Question

나는 다음과 같은 문제가, 내가 가장 좋은 방법이다 궁금 해요 :

__kernel void test(__global int* output){ 

    // ... Code execution to define myValue. 

    // V1 : Threads are idle and wait for the first Work-item to write 
    // the output value. 
    if(get_global_id(0) == 0) output[0] = myValue; 

    // V2 : All work-items perform the same action and try to write in 
    // same global memory. (Is there any lock ?) 
    output[0] = myValue;  
} 

모두가 노력하고 있습니다 내 AMD GPU에. 그러나 나는 어느 것이 최선의 접근인지는 모른다.

는 편집 : (IM은 현재 자사가 간다 최신지고, 그 작업 이후) 칸나의 답변에 따라

, 나는 더 많은 정보를 더 많은 코드를 추가했다.

내 목표는 각 커널에서 head/next_head를 추적하고 작업 그룹 사이에 메모리 블록 포인터의 일관성을 유지하는 것입니다.

첫 번째 접근법에서는 전역 메모리 ptr에서 헤드를 직접 수정합니다. 작업 그룹 번호가 높을 때 블록 위치의 디 - 싱크가 다음 코드와 함께 나타나면 문제가 발생합니다. 코드 실행이 나중에 get_global_id를 기반으로하는 블록을 사용하고 있음에도 불구하고 예상대로 실행되며 각 작업 그룹은 동일한 블록 ptr에 액세스합니다.

그래서 코드를 향상시키고 향후 '병목 현상'이 없는지 확인하기 위해 OpenCL 우수 사례를 찾고 있습니다. 다음 코드에 대해 조언을 구하십시오.

__global void* malloc(size_t sizePtr, __global uchar* heap, ulong* head){ 
    // Get the new ptr inside the heap 
    __global void* ptr = heap + head[0]; 

    // Increment the head. 
    head[0] = head[0] + sizePtr; 

    return ptr; 
} 

__kernel void test(__global uchar* heap, 
        __global ulong* head, 
        __global ulong* next){ 

    // Each work-item set its own local head based on the 
    // global variable. So every thread in any work-group 
    // will start at the same head in the heap. 
    ulong local_head = head[0]; 

    // If get_global_size(0) is 1000. We allocate 1000 + 4000. 
    const uint g_size = get_global_size(0); 

    // Get pointers in a Huge memory block (heap) which allows 
    // to have less memory transfer in-between kernel. 
    // Just need to keep track of them (work in-progess). 
    __global uchar* block1 = malloc(sizeof(uchar) * g_size , heap, &local_head); 
    __global int* block2 = malloc(sizeof(int) * g_size , heap, &local_head); 

    // Process the blocks in here, access them via the get_global_id(0) 
    // as index. 

    // V1 
    if(get_global_id(0) == 0) next[0] = local_head; 

    // V2 
    next[0] = local_head; 

    // If head was 0, the next is now 5000 for all the work-items, 
    // whenever the work-group they are in. 
} 

Answer 1

warp 기반 GPU에서는 확실히 V1이 좋습니다.

V1의 장점은 다른 모든 워프의 조기 종료와 메모리 트래픽 감소입니다.

OpenCL에는 잠금 장치가 없으므로 원자 적 조작을 사용하는 자체 잠금 구성조차 작동하지 않을 수 있습니다.