CUDA : 통합 된 전역 메모리와 비슷한 텍스처 메모리의 액세스 시간입니까?

Question

내 커널 스레드가 병합 된 방식으로 선형 문자 배열에 액세스합니다. 배열을 텍스처로 매핑하면 속도가 향상되지 않습니다. 실행 시간은 과 거의 같습니다. 저는 컴퓨팅 능력 2.0을 갖춘 Tesla C2050을 개발 중이며 어딘가에 글로벌 액세스가 캐싱되어 있습니다. 그게 사실이야? 아마도 그것이 내가 실행 시간의 차이를 보지 못하는 이유입니다.

메인 프로그램 배열

char *dev_database = NULL; 
cudaMalloc((void**) &dev_database, JOBS * FRAGMENTSIZE * sizeof(char)); 

이고 I는

cudaBindTexture(NULL, texdatabase, dev_database, JOBS * FRAGMENTSIZE * sizeof(char)); 

각 스레드는 다음 ID 가 threadIdx.x + blockIdx.x * blockDim.x 및 p이다 ch = tex1Dfetch(texdatabase, p + id)는 오프셋 문자를 판독하여 texture<char> texdatabase 텍스처 바인딩 .

나는 한 번만 바인딩되며 dev_database은 큰 배열입니다. 사실 크기가 너무 큰 경우 을 찾았습니다. 바인드가 실패합니다. 배열 의 크기에 제한이 있습니까? 매우 감사합니다.

Answer 1

성능에 차이가없는 이유는 여러 가지가있을 수 있지만 대부분이 메모리 액세스가 은 병목 현상이 아닌입니다. 병목 현상이 아닌 경우 속도를 높이면 성능에 아무런 영향을 미치지 않습니다.

캐싱 관련 :이 경우 바이트 만 읽으므로 각 워프는 32 바이트를 읽습니다. 이는 4 개의 워핑 그룹이 각 캐시 라인에 매핑됨을 의미합니다. 따라서 캐시 충돌이 거의 없다고 가정하면 캐시에서 최대 4 배까지 재사용 할 수 있습니다. 따라서이 메모리 액세스가 병목 현상이라면 텍스처 캐시가 범용 캐시보다 더 많은 이점을 얻지 못할 수도 있습니다.

먼저 대역폭을 사용하는지 여부와이 데이터 액세스가 범인인지 확인해야합니다. 일단 그렇게하면 메모리 액세스를 최적화하십시오. 고려해야 할 또 다른 전략은 한 번에 한 번에 한 번에 메모리 트랜잭션을 늘리기 위해 스레드 당 하나씩이 아니라로드 당 스레드 당 4 - 16 개의 문자 (바이트 패킹/압축 풀기가있는 char4 또는 int4 구조체 사용)에 액세스하는 것입니다. 글로벌 메모리 버스를 포화시킨다.

시청할 수도있는 a good presentation by Paulius Micikevicius from GTC 2010이 있습니다. 분석 중심 최적화와 비행 중 메모리 트랜잭션의 특정 개념을 다룹니다.