해시를 사용하여 통합 된 전역 메모리 쓰기

제 질문은 동적으로 변경되는 CUDA의 배열 요소 집합에 대한 전역 기록에 관한 것입니다. 다음 커널을 고려해해시를 사용하여 통합 된 전역 메모리 쓰기

여기

__global__ void 
kernel (int n, int *odata, int *idata, int *hash) 
{ 
    int i = blockIdx.x * blockDim.x + threadIdx.x; 
    if (i < n) 
    odata[hash[i]] = idata[i]; 
}

hash 어레이의 제 n 요소 idata의 제 n 소자로부터 갱신 할 odata의 인덱스를 포함한다. 분명히 이것은 끔찍하고 끔찍한 유착 부족을 초래합니다. 내 코드의 경우 한 커널 호출의 해시는 다른 커널의 해시와 완전히 관련이 없으므로 (다른 커널은 다른 방식으로 데이터를 업데이트합니다.)이 특정 kenrel을 최적화하기 위해 데이터를 다시 정렬하는 것은 옵션이 아닙니다.

이 상황의 성능을 향상시킬 수있는 몇 가지 기능이 CUDA에 있습니까? 텍스처 메모리에 대해 많은 이야기를 들었지만, 읽은 것을이 문제의 해결책으로 번역 할 수 없었습니다.

출처

2012-10-17 coastal

텍스처링은 읽기 전용 메커니즘이므로 GMEM에 흩어져있는 쓰기 성능을 직접 향상시킬 수 없습니다. 이 대신 같은 "해시"된 경우

odata[i] = idata[hash[i]];

(아마 당신의 알고리즘은 변환 할 수 있습니까?)

는 그런 다음 Texture mechanism을 고려 몇 가지 이점이있을 수 있습니다. (당신의 본보기는 본질적으로 1D 인 것처럼 보입니다.)

또한 공유 메모리/L1 분할이 캐시쪽으로 최적화되어 있는지 확인해야 할 수도 있습니다. 이것은 흩어져있는 글에 많은 도움이되지 않습니다.

출처

2012-10-17 16:16:23

답장을 보내 주셔서 감사합니다. 사실 나는 여러분이 묘사하는 상황 (전역 메모리로부터의 흩어져있는 읽기)에 정확히 맞닥뜨린 여러 곳에서 내 코드를 사용한다. 이러한 읽기는 실제로 배열에서 무작위 (일반적으로 스파 스)입니다. 이 경우 텍스쳐 메카니즘을 살펴볼 가치가 있습니까? – coastal

텍스처링이 도움이 될 수 있습니다. 텍스처링은 여전히 데이터 지역성 및 재사용에 따라 어떤 이점도 제공합니다. 흩어져있는 독서의 효과가 각 위치를 한 번만 읽는 것이라면 텍스처링이 도움이되지 않습니다. 그러나 잠재적으로 여러 위치를 한 번 이상 읽는다면 (아마 여러 스레드에서) 질감이 약간 개선 될 수 있습니다. –

해시 결과 범위를 제한 할 수 있습니까? 예를 들어 스레드의 처음 1K 반복은 0 - 8K 범위의 odata 만 액세스 할 수 있습니다.

그렇다면 공유 메모리를 사용할 수 있습니다. 공유 메모리 블럭을 할당하고 공유 메모리에 빠른 쓰기 쓰기를 임시로 수행 할 수 있습니다. 그런 다음 병합 된 트랜잭션에서 공유 메모리 블록을 전역 메모리에 다시 기록하십시오.

출처

2012-10-18 16:28:58 ahmad

해시를 사용하여 통합 된 전역 메모리 쓰기

답변

관련 문제