Nsight Compute로 Warp 성능 측정하기

References

• Professional CUDA C Programming
• https://docs.nvidia.com/nsight-compute/NsightComputeCli/index.html

Contents

• Matrix Addition 예제
• Active Warp 측정
• Memory Operations 측정

---

WARP Execution

WARP Execution

지난 포스팅에 이어서, 이번에 warp execution의 동작을 조금 더 잘 이해하기 위해서 행렬 덧셈 커널을 다양한 execution configurations를 사용하여 비교해보도록 하겠습니다. 그리고, Nsight Compute를 사용해 커널의 여러 가지 성능을 측정도 수행할 예정입니다.

 

전체 코드는 아래 링크를 참조해주시기 바랍니다.

https://github.com/junstar92/parallel_programming_study/blob/master/CUDA/matrixAdd/matrixAdd2.cu

GitHub - junstar92/parallel_programming_study: Study parallel programming - CUDA, OpenMP, MPI, Pthread

 

사용되는 커널은 아래의 sumMatrixOnGPU2D 입니다.

__global__
void sumMatrixOnGPU2D(float* A, float* B, float* C, int NX, int NY)
{
    unsigned int ix = blockDim.x*blockIdx.x + threadIdx.x;
    unsigned int iy = blockDim.y*blockIdx.y + threadIdx.y;
    unsigned int idx = iy*NX + ix;

    if (ix < NX && iy < NY) {
        C[idx] = A[idx] + B[idx];
    }
}

사용되는 행렬의 크기는 \(2^{14} \times 2^{14}\) 입니다.

그리고 커널을 실행할 때 사용할 block dimension은 커맨드 입력이 주어지지 않는다면 (32, 32)입니다. 변경하려면 실행할 때 커맨드 라인 argument로 지정할 수 있습니다. 전체 코드는 위 링크를 참조해주세요.

int main(int argc, char** argv)
{
	...
    // setup data size of matrix
    int nx = 1 << 14;
    int ny = 1 << 14;
    
    ...
    // invoke kernel at host
    int dimx = 32;
    int dimy = 32;

    if (argc > 2) {
        dimx = atoi(argv[1]);
        dimy = atoi(argv[2]);
    }

    dim3 block(dimx, dimy);
    dim3 grid((nx + block.x - 1) / block.x, (ny + block.y - 1) / block.y);
    
    ...
}

 

이제 다음의 커맨드로 컴파일합니다.

nvcc -O3 -arch=sm_75 -o matrixAdd2 matrixAdd2.cu -I..

컴파일이 완료되면, 실행 가능한 파일이 생성됩니다.

 

---

Checking Active Warps

먼저, 기본적인 특정 스레드 configuration으로 성능을 평가해보겠습니다. 사용할 스레드 블록의 크기는 (32, 32) / (32, 16) / (16, 32) / (16, 16) 입니다.

 

제 노트북으로 실행한 결과입니다.

교재에서는 스레드 블록 별로 성능의 차이가 약간 발생하는데, 수치상으로 성능의 차이는 확인할 수 없는 것 같습니다. 

 

이전 포스팅에서 active warp의 수를 최대한으로 유지해야 GPU를 최대한으로 활용할 수 있다고 했습니다.

nvprof를 사용하면 '--metrics achieved_occupancy' 옵션으로 사이클 당 평균 active warp의 수와 SM에서 지원하는 최대 warp 수의 비율을 측정할 수 있습니다. 저의 경우에는 nvprof를 사용할 수 없었기 때문에 Nsight Compute를 통해 이를 측정했습니다.

 

Nsight Compute를 사용하면, '--metrics sm__warps_active.avg.pct_of_peak_sustained_active'를 사용하면 동일하게 측정할 수 있습니다.

ncu.bat --metrics sm__warps_active.avg.pct_of_peak_sustained_active ./matrixAdd2.exe 32 32

 

성능의 차이가 거의 나지 않기 때문에 active warp의 비율 또한 큰 차이는 없습니다. 다만, 블록의 수가 증가할수록 이 비율이 커질 줄 알았습니다만, (32, 32)일 때의 비율이 (16, 32)일 때의 비율보다 크게 나와서 의외였습니다. 그래도 (16, 16)일 때 블록의 수가 더 많기 때문에 비율이 조금 더 큰 것을 볼 수 있는 것 같습니다.

 

하지만 가장 높은 achieved occupancy라고, 가장 빠르지는 않습니다. 그러므로 높은 occupancy가 항상 높은 성능을 가지는 것도 아닙니다.

 

---

Checking Memory Operations

sumMatrixOnGPU2D 커널은 두 번의 메모리 load와 한 번의 메모리 store, 총 3개의 메모리 연산(C[idx] = A[idx] + B[idx])이 있습니다. 이 메모리 연산의 효율성도 측정할 수 있습니다.

nvprof의 경우에는 '--metrics gld_throughput'으로 측정할 수 있으며, Nsight Compute의 경우에는 l1tex__t_bytes_pipe_lsu_mem_global_op_ld.sum.per_second metrics을 사용하여 측정할 수 있습니다.

ncu.bat --metrics l1tex__t_bytes_pipe_lsu_mem_global_op_ld.sum.per_second ./matrixAdd2.exe 32 32

위 커맨드로 Global Memory Load 효율을 측정한 결과는 다음과 같습니다.

큰 차이는 없지만, 수치상으로 (16, 16)이 가장 높은 load 처리량을 보여주고 있습니다. 마찬가지로 load 처리량이 높다고, 높은 성능이라는 것은 아닙니다.

 

다음으로 Global load efficiency를 측정해보겠습니다. nvprof로는 'gld_efficiency' metrics으로 측정할 수 있고, Nsight Compute로는 'smsp__sass_average_data_bytes_per_sector_mem_global_op_ld.pct'로 측정할 수 있습니다.

이는 required global memory load 처리량 대비 requested globa memory load 처리량의 비율입니다.

ncu.bat --metrics smsp__sass_average_data_bytes_per_sector_mem_global_op_ld.pct ./matrixAdd2.exe 32 32

모두 100%로 측정됩니다. 메모리를 요구하는대로 바로바로 요청했다고 해석할 수 있을 것 같습니다.

 

---

그렇다면 block.x를 더 조절하면 결과는 어떨까요 ?

더 다양한 범위의 thread configuration으로 테스트를 해보도록 하겠습니다.

여러 크기의 thread block으로 테스트한 결과입니다. 대부분 유사한 성능을 보여주고 있습니다. 마지막의 (256, 8)로는 커널을 실행할 수 없는데, 이는 제 GPU에서 블록당 최대 스레드의 개수가 1024개이기 때문입니다.

유의미한 성능의 차이는 보여주지 않고 있습니다.

 

다음으로 각 크기에서 achieved_occupancy를 측정해보도록 하겠습니다.

(64, 2) / (64, 4) / (64, 8)일 때의 결과입니다.

다음은 (128, 2) / (128, 4) / (128, 8)일 때의 결과입니다.

다음은 (256, 2) / (256, 4) 일 때의 결과입니다.

(128, 8)일 때 가장 좋은 결과를 보여주고 있으며, (256, 4)일 때도 높은 occupancy를 보여줍니다.

 

 

여러 가지 조건으로 성능을 측정했지만, 유의미한 성능의 차이는 크게 보이지 않고 그 결과 또한 비교할만큼의 큰 차이를 보여주지 않았습니다. 최근 디바이스들의 성능이 다들 좋아져서 그런 것도 있겠지만, 다음에는 조금 더 복잡한 커널을 통해서 분석하는 기회를 가지도록 하겠습니다. 이번 포스팅에서는 Nsight Compute에 대한 사용법을 조금이나마 익히는 것으로 만족합니다.. ! ㅠ