OpenCL : __constant 대 __local?

Question

많은 양의 값 (여전히 64KB보다 작음)이 있다고 가정 해 봅시다. 은 매우이지만 커널에는 쓰이지는 않습니다. 그러나 외부에서 바뀔 수 있습니다. 배열에는 두 개의 값 집합이 있으며 왼쪽과 오른쪽을 호출 할 수 있습니다. 문제는 큰 배열을 __global로 가져 와서 __local 왼쪽 및 __local 오른쪽 배열에 쓰는 것이 더 빠릅니까? 또는 상수 __constant 큰 그것을 얻고 커널에서 accesing 처리 할?

__kernel void f(__global large, __local left, __local right, __global x, __global y) { 
    for(int i; i < size; i++) { 
     left[i] = large[i]; 
     right[i] = large[i + offset]; 
    } 
    ... 
    x = foo * left[idx]; 
    y = bar * right[idx]; 
} 

대 : 예를 들어

__kernel void f(__constant large, __global x, __global y) { 
    ... 
    x = foo * large[idx]; 
    y = bar * large[idx * offset]; 
} 

(인덱싱은 좀 더 복잡하지만 예를 들어, 매크로를 만들 수 있습니다) 나는 상수 메모리가 글로벌 공간에 살고 있음을 읽을 , 그래서 더 천천히해야합니까? Nvidia 카드에서 실행됩니다.

Answer 1

우선 두 번째 경우에는 CPU에 결과를 사용할 수 있도록해야합니다. 계산 후 globalspace으로 다시 복사한다고 가정합니다.

커널에서하는 일에 달려 있다고 생각합니다. 예를 들어, 커널 계산이 무거 우면 (스레드 당 많은 계산) 첫 번째 옵션은 지불 할 수 있습니다. 왜?

• 당신은 local 공간 left 및 right에 global large 공간에서 시간 복사 데이터를 보낼 - 허용
• 당신은 지역 공간의 데이터에 대한 계산을 많이 할 - 당신은 어떤을 보내고 OK
• 을 localleft 및 right부터 globallarge까지 시간 복사 - 받아 들일 만하다. 당신은 커널 경우

당신은 constant 공간에 데이터를 몇 가지 계산을 수행 한 후,

• 각 스레드가 몇 가지 작은 계산을 할 것입니다 즉 상대적으로 가볍다. 아마도 가장 많이 액세스 할 필요는 없음을 의미합니다.
• 중간 결과를 로컬 공간에 저장합니다.
• local 공간에서 global 공간으로 다시 복사하는 데 약간의 시간이 소요됩니다. - 받아 들일 만하다.

큰 커널을 요약하면 첫 번째 옵션이 더 좋습니다. 작은 커널의 경우 두 번째.

P. 하나 더주의해야 할 점은 당신이 복수의 커널을 가지고 있고 하나씩 다른 커널을 가지고 있다면, 첫 번째 옵션으로 확실히 간다. 왜냐하면 전역 메모리 공간에 데이터를 보관할 수 있고 커널을 시작할 때마다 복사 할 필요가 없기 때문입니다.

편집 : 매우 자주에 액세스했다고 말했기 때문에 자주 첫 번째 옵션을 사용해야한다고 생각합니다.