CUDA : 공유 메모리를 언제 사용하고 언제 L1 캐싱을 사용해야합니까?

Question

Compute Capability 2.0 (Fermi)이 출시 된 후 공유 메모리에 남은 유스 케이스가 있는지 궁금해했습니다. 즉, L1이 백그라운드에서 마술을 수행하도록하는 것보다 공유 메모리를 사용하는 것이 더 나은 이유는 무엇입니까?

간단히 공유 메모리가있어 수정하지 않고 효율적으로 CC < 2.0을 실행하도록 설계된 알고리즘을 허용합니까?

공유 메모리를 통해 공동 작업하려면 블록의 스레드가 공유 메모리에 쓰고 __syncthreads()과 동기화하십시오. 단순히 전역 메모리 (L1을 통해)에 쓰기 만하고 __threadfence_block()과 동기화해야하는 이유는 무엇입니까? 후자의 옵션은 두 개의 서로 다른 값 위치와 관련이 없으므로 구현이 더 쉬워야하며 전역 메모리에서 공유 메모리로의 명시적인 복사가 없기 때문에 더 빨라야합니다. 데이터가 L1에서 캐시되기 때문에 스레드는 데이터를 전역 메모리에 실제로 전달할 때까지 기다릴 필요가 없습니다.

공유 메모리를 사용하면 블록 기간 동안 거기에 저장된 값이 그대로 유지됩니다. 이것은 L1에서 값과 반대되는데, 이는 충분히 자주 사용하지 않으면 축출됩니다. 알고리즘이 실제 사용하는 패턴을 기반으로 L1이 공유 메모리를 관리하도록하는 것보다 공유 메모리에서 데이터가 거의 사용되지 않는 캐시가 더 나은 경우가 있습니까?

Answer 1

내가 아는 한, GPU의 L1 캐시는 CPU의 캐시와 매우 유사하게 동작합니다. 그래서 "충분히 자주 사용하지 않으면 퇴출되는 L1의 값과 반대입니다"라는 귀하의 의견은 나에게 많은 의미가 없습니다.

L1 캐시의 데이터는 제거되지 않습니다. 자주 사용되었습니다. 일반적으로 이전에 캐시에 없었던 메모리 영역에 대한 요청이 이루어지고 이미 사용중인 주소로 해석되는 메모리 영역에 대한 요청이있을 때 제거됩니다. NVidia에서 사용하는 정확한 캐싱 알고리즘을 모르지만 일반적인 n 방향 연결을 가정하면 각 메모리 항목은 주소를 기반으로 전체 캐시의 작은 하위 집합에만 캐시 될 수 있습니다.

귀하의 질문에 답변 할 수도 있습니다. 공유 메모리를 사용하면 캐시를 통해 모든 것이 자동으로 수행되는 동안 저장된 내용을 완벽하게 제어 할 수 있습니다. 컴파일러와 GPU가 여전히 메모리 액세스를 최적화하는 데 매우 영리 할지라도, 어떤 입력이 주어질 것인지, 그리고 어떤 스레드가 무엇을 할 것인지를 알고 있기 때문에 더 좋은 방법을 찾을 수 있습니다 (특정 확장 영역)

Answer 2

전역 메모리로드/저장은 데이터가 캐시에 있더라도 병합 규칙의 적용을 받지만 공유 메모리는 임의 액세스면에서 훨씬 융통성이 있습니다. 나는 히스토그램을 저장/컴퓨팅하기 위해 L1 캐싱을 사용해 보았지만 세미 임의 액세스 패턴으로 인해 공유 메모리를 사용하는 것보다 훨씬 느리다.

NVIDIA employee에 따르면 현재 L1 캐시는 write-through (즉시 L2 캐시에 쓰기)되어 프로그램 속도가 느려집니다.

기본적으로 성능을 원하면 기본적으로 캐시가 방해가됩니다.