2012-10-17 2 views
3

공유 메모리를 사용하여 OpenACC로 항목을 캐시하려고합니다. 내가 일하고 있어요 기본적으로 무엇을OpenACC와 공유 메모리 사용

는 행렬 곱셈, 그리고 내가 가지고있는 것은 이것이다 : 내가하고 싶은 무엇

typedef float ff; 

// Multiplies two square row-major matrices a and b, puts the result in c. 
void mmul(const restrict ff* a, 
      const restrict ff* b, 
      restrict ff* c, 
      const int n) { 
#pragma acc data copyin(a[0:n*n], b[0:n*n]) copy(c[0:n*n]) 
{ 

#pragma acc region 
{ 

#pragma acc loop independent vector(16) 
    for (int i = 0; i < n; ++i) { 
#pragma acc loop independent vector(16) 
    for (int j = 0; j < n; ++j) { 
     ff sum = 0; 
     for (int k = 0; k < n; ++k) { 
     sum += a[i + n * k] * b[k + n * j]; 
     } 
     c[i + n * j] = sum; 
    } 
    } 

} 
} 
} 

이 행렬의 캐시 타일 공유 메모리를 사용하다 '는 '및'b '를'c '계산에 사용하려면 CUDA mmul 알고리즘과 비슷한 방식으로 사용해야합니다.

는 기본적으로 CUDA에 내 블록의 정확한 크기를 알 것이며, 할 수있을 것입니다 :

  • 가의 '관련'부분을 복사 블록의 크기와 공유 메모리를 선언 블록
  • 사용 데이터

나는 내가

#pragma acc cached 
사용할 수 있습니다 이해하고이 데이터

지시어를 사용하고 벡터gang 옵션을 사용하여 블록 크기를 지정할 수 있지만 CUDA 아키텍처에 매핑되는 방법을 이해하는 데 어려움이 있습니다.

OpenACC와 비슷한 것을 얻을 수있는 방법이 있습니까? 캐시 된 지시문을 사용하거나 CUDA에서 공유 메모리의 성능을 OpenACC로 매핑하는 방법에 대한 자습서/리소스가 있습니까? 당신이 PGI 가속기 컴파일러를 사용하는 경우

+1

PGI 가속기 컴파일러는 이미 공유 메모리를 사용하고있을 수 있습니다. -Minfo 스위치를 사용하여 출력을 검사 했습니까? 이 [튜토리얼] (http://developer.nvidia.com/cuda/openacc-example-part-1)이 중요 할 수 있습니다. –

+1

예,하지만 Minfo 스위치는 구현에 사용중인 공유 메모리가 얼마나되는지를 알려줍니다. 이 방법이 유용하지만 ** 이러한 메모리를 명시 적으로 ** 조작하는 방법이 있는지 알고 싶습니다. 생성 된 높은 수준의 쿠다를 볼 수 있다는 것은 매우 유용합니다. – leo

+0

@leo 질문에 대한 답변을 찾았습니까? OpenACC에서 명시 적으로 공유 메모리를 정의 할 수 있었습니까? – Millad

답변

4

, 당신은 생성 된 PTX 파일을 덤프 및 실행의 부하에 무슨 일이 일어나고 있는지 볼 수 있습니다

pgcc -acc -fast -Minfo -ta=nvidia,cc13,keepptx matrixMult.c -o matrixMult 

생성 된 PTX

는 현재 디렉토리에 저장됩니다.

편집 : 고급 코드 (C 또는 Fortran 용)를 선호 할 수 있습니다. 따라서 -ta=nvidia,cc13,keepptx,keepgpu을 따르십시오.