Cuda 프로그래밍의 액티브 워프

Nsight IDE를 사용하여 코드를 성능 분석하려고합니다.Cuda 프로그래밍의 액티브 워프

나는 행렬을 더하는 간단한 예제를 사용했다.

내가 좋아하는 내 커널을 호출 오전 :

VecAdd < < < 1, BLOCK_SIZE BLOCK_SIZE >>> DA (, dB, DC, BLOCK_SIZE BLOCK_SIZE);

여기 BLOCK_SIZE는 16

__global__ void VecAdd(float *dA, float *dB, float *dC, int N) 
{ 
    int i = threadIdx.x; 
    if (i < N) 
     dC[i] = dA[i] + dB[i]; 
}

점유율 분석을하는 동안,

내가 0.97으로 달성 액티브 휘어을 얻고있다.

이유에 대해 알 수 없습니다.

보고서를 첨부했습니다. 누군가 이것이 왜 그런지 설명해 주시겠습니까? 점유 달성

enter image description here

출처

2014-05-21 silver surfer

* 하나 * 16 * 16 블록을 실행 중입니까? 나는 그것이 합리적인 프로필을위한 충분한 지 모르겠습니다. –

성능 분석을보고 싶기 때문에 16 * 16 스레드와 1 블록으로 시작했습니다. 이제 나의 목표는 점유를 늘리는 것입니다. –

는 active_warps/elapse_cycles/MAX_WARPS_PER_SM * 100

커널 출시 8 휘어 1 개 블록의 비율입니다. 달성 된 점유 통계는 평균 1 워프 활성이 매우 낮은 것을 보여줍니다. 분명한 이유는 이것이 8이 아닌 이유입니다.

소스를 제공하지 않았으므로 상수 읽기 5 개, 2 개의 32 비트 전역로드 1 개, 32 비트 쓰기 1 개, 32 비트 쓰기 1 개를 수행하는 VecAdd CUDA SDK 샘플을 수정했다고 가정합니다. 색인 생성 및 주소 계산을위한 몇 가지 기본 수학. 이것은 L2에서 모든 메모리 연산이 성공했다고 가정 할 때 워프 당 약 300 사이클을 필요로합니다. 이는 실행하기 전에 배열을 호스트에서 장치로 복사했을 가능성이 높습니다. 커널 지속 시간 자체는 2-3 μs입니다. 8 * 300 사이클/2500 사이클 = 1 SM에서 사이클 당 ~ 1 활성 워핑.

실행 오버 헤드, 작업 분산 오버 헤드 및 각 워프 저장소가 쓰기 데이터 버퍼를 지우는 데 걸리는 시간은 8 개의 워프가 활성화 된 시간으로 계산되지 않습니다. 워프 당 작업량을 늘리면 값이 8에 가까워지고 실행 된 스레드 수가 주어진 경우 달성 가능한 최대 값이됩니다. 장치를 포화시키기 위해 그리드 크기를 늘리면 SM 당 64 개의 평균 활성 워프에 가까워 야합니다.

출처

2014-05-21 07:05:06

커널 코드를 추가했습니다. –

Cuda 프로그래밍의 액티브 워프

답변

관련 문제