2017-01-02 3 views
1

CUDA로 이미지를 처리하고 싶습니다. 각 픽셀의 새 값은 한 행의 두 인접 픽셀을 기반으로 계산됩니다. 각 값이 두 번만 사용되므로 픽셀 값에 __shared__ 메모리를 사용하는 것이 맞습니까? 타일이 문제의 구조에 맞지 않기 때문에 타일을 사용하는 것이 잘못된 방법이 아닙니까? 내 접근 방식은 각 픽셀마다 스레드를 실행하고 각 스레드에 대해 매번 인접 픽셀 값을로드하는 것입니다.인접 배열 요소에 공유 메모리를 사용 하시겠습니까?

+0

누가이 질문을 종료하겠습니까? 이 질문은 너무 광범위하지 않으며 간단하고 간결하게 대답 할 수 있습니다. 쿠다 (CUDA)에 관한 많은 질문들이 요즘에는 질이 낮다는 것에 동의하지만,이 질문은 그렇지 않습니다. 질문을 즉시 닫는 습관이 너무 평범한 것으로 생각됩니다 (나는 스스로 이것을 포함하도록 기꺼이 인정할 것입니다). – tera

답변

4

현재 지원되는 모든 CUDA 아키텍처에는 캐시가 있습니다. 계산 기능 3.5부터 읽기 전용 데이터 (읽기 - 쓰기 데이터는 L2에만 캐시되므로 L1 캐시는 읽기 전용 데이터로 제한됨)에 특히 효율적입니다. 입력 데이터에 대한 포인터를 const __restrict__으로 표시하면 컴파일러에서 via the L1 texture cache을로드 할 가능성이 높습니다. __ldg() builtin을 명시 적으로 사용하여이를 강제로 수행 할 수도 있습니다.

공유 메모리를 통해 인접 픽셀의 데이터 재사용을 명시 적으로 관리 할 수는 있지만 캐시를 사용하는 것 이상의 이점은없는 것으로 나타났습니다.

물론 공유 메모리를 사용하는지 여부에 관계없이 x 방향의 블록 크기를 최대화하고 최적의 액세스 지역성을 위해 blockSize.y를 1로 사용하려고합니다.

1

병합 된 메모리 액세스를 활용하여 공유 메모리를 사용하여 결합하십시오. 이미지가 행 단위로 저장되는지 확인하기 만하면됩니다. 각 블록은 선형 배열의 덩어리를 처리합니다. 데이터 재사용 (첫 번째와 마지막 픽셀을 제외한 모든 픽셀은 세 번 처리에 참여할 것입니다) 때문에 커널의 시작 부분에서 처리 할 모든 픽셀의 값을 공유 메모리로 복사하면 유용합니다.