전역 쓰기에서 메모리 병합

Question

CUDA 디바이스에서 전역 메모리 읽기에서 병합하는 것만 큼 전역 메모리 쓰기가 병합되고 있습니까? 그렇다면 어떻게 설명 할 수 있습니까? 또한 초기 세대의 CUDA 장치와이 문제에 관한 가장 최근의 장치 사이에는 차이점이 있습니까?

Answer 1

통합 읽기 (또는 일치하지 않음)는 병합 읽기 (또는 해당 없음)와 마찬가지로 성능에 영향을 줄 수 있습니다.

int i = my_int_data[threadIdx.x+blockDim.x*blockIdx.x]; 

단일 하여 만족스러운 본질적으로 모든 개별 스레드를 말하는 메모리 컨트롤러 트랜잭션 (판독 할 수

는 합체 판독 휘어짐 명령에 의해 트리거 판독 트랜잭션, 예를 들어,이 때 발생 . 읽기) 하나의 캐시 라인에서 오는

A는 쓰기 병합하는 것은 발생하면 워프 명령에 의해 트리거 쓰기 트랜잭션 예 :

my_int_data[threadIdx.x+blockDim.x*blockIdx.x] = i; 

은 메모리 컨트롤러에서 쓰기 트랜잭션으로 단일 충족 될 수 있습니다.

위의 예에서 나는 세대별로 차이가 없다는 것을 보여 주었다.

그러나 이전 장치가 아닌 이후 장치에서 병합 (즉, 단일 메모리 컨트롤러 트랜잭션으로 축소) 할 수있는 다른 유형의 읽기 또는 쓰기가 있습니다. 위의 예에서

int i = my_int_data[0]; 

같은 글로벌 위치에서 읽어들이는 모든 스레드 : 하나의 예는 "방송 읽기"입니다. 최신 장치에서는 이러한 읽기가 단일 트랜잭션의 모든 스레드에 "브로드 캐스트"됩니다. 이전의 일부 장치에서는 스레드의 직렬화 된 서비스가 발생합니다. 하나의 위치에 쓰는 다중 스레드는 정의되지 않은 동작을하기 때문에 이러한 예는 아마도 쓰기 작업에서 결과가 없습니다. 그러나 쓰기는 새로운 장치에 병합 할 수 있지만 "스크램블"오래되지 :

my_int_data[(threadIdx.x+5)%32] = i; 

위의 모든 쓰기가 (경사 이내) 고유 주 및 개별 캐시 라인에 속하는하지만 그들은 병합을 만족하지 않는 1.0 또는 1.1 장치에 대한 요구 사항이지만 최신 장치에 적용해야합니다.

global memory access description for devices of cc 1.0 and 1.1을 읽고 나중 장치와 비교하면 나중 장치에서 완화 된 이전 장치에서 병합하기위한 요구 사항 중 일부가 표시됩니다.

Answer 2

내가 수행 한 코스에서이 실험을 수행했습니다. 합병은 적당히 이 읽기보다 쓰기에서 중요하다는 것이 L1과 L2 캐시가 나중에 사용하기 위해 일부 사용하지 않는 데이터를 저장하기 때문에 나타났습니다.