c++ - OpenCL:动态内存分配，使用空闲工作项好还是同时写入好

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 上工作。但我不知道哪种方法最好。

编辑:

基于 kanna 答案，我添加了更多代码以获取更多信息(因为我目前正在处理它，所以它会不断更新)。

我的目标是跟踪每个内核的 head/next_head，并在工作组之间保持内存块指针的一致性。

第一种方法，我直接修改了全局内存ptr中的head，当work-group数较高时会出现问题，出现 block 位置不同步，用下面的代码，似乎一切都是按预期运行并且每个工作组访问相同的 block ptr，尽管代码执行之后基于它们的 get_global_id 使用这些 block 。

因此，我正在寻找 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 的优点是提前终止所有其他 warp 和较少的内存流量。

OpenCL 中没有锁，甚至不能保证使用原子操作构造您自己的锁。