cuda 상수 메모리 참조

Question

상수 메모리 (글로벌 변수)에 배열이 있고 함수 호출 cudaGetSymbolAddress에 의해 참조가 확보되었습니다. 이 참조를 사용하여 전역 변수를 사용하는 대신 상수 데이터를 가져올 때 커널이 느리게 실행됩니다. 이것에 대한 이유는 무엇입니까?

__constant__ int g[2] = {1,2}; 
// __device__ int g[2] = {1,2}; 

// kernel: use by reference 
__global__ void add_1(int *a, int *b, int *c, int *f) 
{ 
    int tid = blockIdx.x * blockDim.x + threadIdx.x; 
    c[tid] = f[0] * a[tid] + f[1] * b[tid]; 
} 

// kernel: use global variable 
__global__ void add_2(int *a, int *b, int *c, int *f) 
{ 
    int tid = blockIdx.x * blockDim.x + threadIdx.x; 
    c[tid] = g[0] * a[tid] + f[1] * b[tid]; 
} 

int main() 
{ 
    ...... 
    // a,b,c are large arrays in device memory of size 40960. 

    int *f; 
    cudaGetSymbolAddress((void **)&f, (char *)&g); 

    add_1 <<< 160, 256 >>> (a, b, c, f); 

    ...... 
} 

이 샘플 코드는 warp의 모든 스레드가 동시에 같은 위치를로드합니다. 주석으로 코드를 직접 상수 메모리 캐시가

이유는 일정 캐시의 부족 (talonmies에 의해)를 사용하지 않는 이유에 대해 일정한 메모리를

설명에 접근하는 것입니다. 캐시 된 액세스는 컴파일러가 상수 상태 공간에있는 것으로 표시된 변수 명시에 특정 PTX 명령어 (ld.const)를 내보낼 때만 발생합니다. 그리고 컴파일러가 이것을 알고있는 방법은 변수가 선언되었을 때입니다. __constant__ - 코드 생성에 영향을주는 정적 인 컴파일 타임 속성입니다. 동일한 프로세스가 런타임에 발생할 수 없습니다.

전역 메모리에 포인터를 전달하고 컴파일러가 상수 상태 공간의 포인터를 확인할 수없는 경우 상수 캐시를 통해 해당 메모리에 액세스하기위한 올바른 PTX를 생성하지 않습니다. 결과적으로 액세스 속도가 느려집니다.

질문 답변되지

배열 g__device__는 변수로 선언 된 경우에도, 코드가 참조가 사용될 때 느린 이유. 레지스터에 전역 메모리로드를 들어, PTX 코드를보고 작성자 : ld.global.s32의

• 2 명령어는 레지스터에 4 바이트를로드하는 데 사용됩니다. (글로벌 변수를 사용하는 코드) 2 개 레지스터에 8 개의 바이트를로드하는 ld.global.v2.s32의 명령어가 사용될
• 1 (코드 사용 참조)

차이 및 문서 참조 무슨 감상 할 수 있는가? 글로벌 메모리와는 달리

Answer 1

, 그들은 균일 한 (계산 능력 1.x에서의 (절반의 모든 스레드가 아닌 경우 워프 액세스) 같은 주소를 (여러 트랜잭션으로 분리) 직렬화 얻을 것이다 일정한 메모리에 액세스합니다.

액세스가 일정 할 경우 상수 메모리 만 사용하십시오.