CUDA 코드의 상수 메모리 사용량

Question

직접 알 수는 없지만, 내 커널에서 사용되는 메모리를 보장하는 가장 좋은 방법은 상수입니다. http://stackoverflow...r-pleasant-way에도 비슷한 질문이 있습니다. GTX580으로 작업 중이며 2.0 용으로 만 컴파일 중입니다.

대체 방법입니다

cudaMalloc(src, size); 
cudaMemcpy(src, hostSrc, size, cudaMemcpyHostToDevice); 
Foo<<<...>>>(src, result); 

이

__constant__ src[size]; 

파일을 .CU에 추가하려면에서 SRC에게 포인터를 제거 : 내가 호스트에서 다음 코드를 실행

__global__ Foo(const int *src, float *result) {...} 

처럼 내 커널 보인다 커널 및 실행

cudaMemcpyToSymbol("src", hostSrc, size, 0, cudaMemcpyHostToDevice); 
Foo<<<...>>>(result); 

이 두 가지 방법이 동일합니까? 아니면 첫 번째 방법이 전역 메모리 대신 상수 메모리의 사용을 보장하지 않습니까? 크기가 인 경우 동적으로 변경되므로 두 번째 방법은 내 경우에 편리하지 않습니다.

Answer 1

두 번째 방법은 배열이 CUDA 상수 메모리로 컴파일되고 상수 메모리 캐시를 통해 올바르게 액세스되는지 확인하는 유일한 방법입니다. 그러나 배열 블록의 내용이 스레드 블록 내에서 어떻게 액세스 될 것인지 스스로에게 물어야합니다. 모든 스레드가 어레이에 균일하게 액세스하면 상수 메모리 캐시의 브로드 캐스트 메커니즘이 있기 때문에 상수 메모리를 사용할 때 성능상의 이점이 있습니다 (상시 메모리가 오프 칩 DRAM에 저장되고 캐시가 전역 메모리 대역폭을 저장하기 때문에 캐시 DRAM 트랜잭션 수를 줄입니다. 그러나 액세스가 무작위 인 경우 로컬 메모리에 대한 액세스가 직렬화되어 성능에 부정적인 영향을 미칠 수 있습니다.

__constant__ 메모리에 적합 할 수있는 전형적인 것들은 모델 계수, 가중치 및 런타임에 설정해야하는 기타 상수 값입니다. 페르미 GPU에서 커널 인수 목록은 예를 들어 상수 메모리에 저장됩니다. 그러나 내용이 불균일 한 액세스이고 구성원 유형이나 크기가 호출에서 호출까지 일정하지 않으면 일반 전역 메모리가 바람직합니다.

또한 GPU 컨텍스트 당 상수 메모리의 한계는 64kb이므로 상수 메모리에 매우 많은 양의 데이터를 저장하는 것은 실용적이지 않습니다. 캐시가있는 읽기 전용 저장소가 많이 필요한 경우 데이터를 텍스처에 바인딩하고 성능이 어떤지 확인하는 것이 좋습니다. Pre-Fermi 카드에서는 대개 편리한 성능 이득을 얻습니다. Fermi에서 아키텍처의 캐시 레이아웃이 향상 되었기 때문에 결과는 전역 메모리에 비해 예측이 어려울 수 있습니다.

Answer 2

첫 번째 방법은 메모리가 함수 Foo 내부에서 일정 함을 보장합니다. 2 개는 동등하지 않고, 2 번째의 것은 초기화 후에 거기에 계속되는 것을 보증합니다. 당신보다 더 역동적 인 것이 필요하다면 첫 번째 방법과 비슷한 것을 사용하십시오.