OpenCL의 로컬 메모리 펜스와 글로벌 메모리 펜스의 차이점은 무엇입니까?

Question

예제와 비슷한 PyOpenCL을 사용하여 감소 합계를 만들려고했습니다 : https://dournac.org/info/gpu_sum_reduction. 모든 값이 1 인 벡터를 합계하려고합니다. 결과는 첫 번째 요소에서 16384 여야합니다. 그러나 일부 지점 만 수집되고있는 것처럼 보입니다. 로컬 색인이 필요합니까? 경쟁 조건이 있습니까 (두 번 실행하면 결과가 동일하지 않습니다)? 다음 코드는 무엇이 잘못 되었습니까?

import numpy as np 
import pyopencl as cl 

def readKernel(kernelFile): 
    with open(kernelFile, 'r') as f: 
     data=f.read() 
    return data 

a_np = np.random.rand(128*128).astype(np.float32) 
a_np=a_np.reshape((128,128)) 
print(a_np.shape) 

device = cl.get_platforms()[0].get_devices(cl.device_type.GPU)[0] 
print(device) 
ctx=cl.Context(devices=[device]) 
#ctx = cl.create_some_context() #ask which context to use 
queue = cl.CommandQueue(ctx) 
mf = cl.mem_flags 

a_g = cl.Buffer(ctx, mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=a_np) 

prg = cl.Program(ctx,readKernel("kernel2.cl")).build() 

prg.test(queue, a_np.shape, None, a_g) 

cl.enqueue_copy(queue, a_np, a_g).wait() 
np.savetxt("teste2.txt",a_np,fmt="%i") 

커널은 다음과 같습니다

__kernel void test(__global float *count){ 
    int id = get_global_id(0)+get_global_id(1)*get_global_size(0); 
    int nelements = get_global_size(0)*get_global_size(1); 

    count[id] = 1; 
    barrier(CLK_GLOBAL_MEM_FENCE); 

    for (int stride = nelements/2; stride>0; stride = stride/2){ 
     barrier(CLK_GLOBAL_MEM_FENCE); //wait everyone update 
     if (id < stride){ 
      int s1 = count[id]; 
      int s2 = count[id+stride]; 
      count[id] = s1+s2; 
     } 
    } 
    barrier(CLK_GLOBAL_MEM_FENCE); //wait everyone update 
} 

Answer 1

문제는 커널이 하나 개의 작업 그룹 내에서 감소 할 구현되고 묵시적으로 많은 작업 그룹이 schedulled 점이다.

GPU에 따라 작업 그룹 당 최대 작업 항목 수가 다릅니다. Nvidia의 경우 1024, AMD 및 Intel 256 (구형 GPU의 Intel은 512)

이 예에서는 GPU의 작업 그룹 당 최대 작업 항목이 256 인 것으로 가정합니다.이 경우 최대 2d worgroup 크기는 16x16 일 수 있으므로이 크기의 행렬을 사용하면 커널이 올바른 결과를 반환합니다. 원래 크기 인 128x128을 사용하고 커널을 예약 할 때 로컬 크기를 지정하지 않으면 구현시 계산 된 전역 크기 128x128과 로컬 크기 (가능성이 높음) 16x16이되므로 8 개의 그룹이 예약됩니다. 현재 커널에서 각 작업 그룹은 다른 id의 계산을 시작하지만 인덱스는 0까지 줄어들 기 때문에 경쟁 조건이 있으므로 각각 다른 결과가 실행됩니다. (16 × 16), (16, 16) 또는 무엇이든 최대 작업 :

하나 개의 작업 그룹 내에서 모든 것을 계산하고 글로벌, 지역의 크기를 일정 커널을 다시 작성
1. :

   는 당신이 해결하는 2 가지 옵션이 있습니다 작업 그룹 장치 당 항목 수는

2. 입니다. 전역, 로컬 크기 : (128x128), (16x16)을 사용하고 각 작업 그룹은 결과를 계산 한 다음 최종 결과를 얻으려면 각 작업 그룹에 대해 합계해야합니다.

128x128의 경우 첫 번째 옵션이 더 빨리 수행되어야하며 구현이 더 간단해야하므로 첫 번째 옵션이 선호됩니다.