전역 메모리 재생 오버 헤드는 어디에서 발생합니까? 엔비디아 비주얼 프로파일에서 글로벌 메모리 1 GB를 작성하는 아래의 코드를 실행

Question

, 내가 얻을 :
- 100 %의 스토리지 효율성
- 69.4 % (128.6 GB/s)의 DRAM 활용
- 18.3 % 총 재생 오버 헤드
- 전체 메모리 재생 오버 헤드가 18.3 %입니다.

메모리 쓰기가 병합되고 커널에 분기가 없으므로 전역 메모리 재생 오버 헤드는 어디에서 오는가? 나는 이것을 Ubuntu 13.04에서 실행하고 있으며, nvidia-cuda-toolkit 버전 5.0.35-4ubuntu1을 사용하고 있습니다.

#include <cuda.h> 
#include <unistd.h> 
#include <getopt.h> 
#include <errno.h> 
#include <stdio.h> 
#include <stdlib.h> 
#include <time.h> 
#include <stdint.h> 
#include <ctype.h> 
#include <sched.h> 
#include <assert.h> 

static void 
HandleError(cudaError_t err, const char *file, int line) 
{ 
    if (err != cudaSuccess) { 
     printf("%s in %s at line %d\n", cudaGetErrorString(err), file, line); 
     exit(EXIT_FAILURE); 
    } 
} 
#define HANDLE_ERROR(err) (HandleError(err, __FILE__, __LINE__)) 

// Global memory writes 
__global__ void 
kernel_write(uint32_t *start, uint32_t entries) 
{ 
    uint32_t tid = threadIdx.x + blockIdx.x*blockDim.x; 

    while (tid < entries) { 
     start[tid] = tid; 
     tid += blockDim.x*gridDim.x; 
    } 
} 

int main(int argc, char *argv[]) 
{ 
    uint32_t *gpu_mem;    // Memory pointer 
    uint32_t n_blocks = 256;  // Blocks per grid 
    uint32_t n_threads = 192;  // Threads per block 
    uint32_t n_bytes = 1073741824; // Transfer size (1 GB) 
    float elapsedTime;    // Elapsed write time 

    // Allocate 1 GB of memory on the device 
    HANDLE_ERROR(cudaMalloc((void **)&gpu_mem, n_bytes)); 

    // Create events 
    cudaEvent_t start, stop; 
    HANDLE_ERROR(cudaEventCreate(&start)); 
    HANDLE_ERROR(cudaEventCreate(&stop)); 

    // Write to global memory 
    HANDLE_ERROR(cudaEventRecord(start, 0)); 
    kernel_write<<<n_blocks, n_threads>>>(gpu_mem, n_bytes/4); 
    HANDLE_ERROR(cudaGetLastError()); 
    HANDLE_ERROR(cudaEventRecord(stop, 0)); 
    HANDLE_ERROR(cudaEventSynchronize(stop)); 
    HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime, start, stop)); 

    // Report exchange time 
    printf("#Delay(ms) BW(GB/s)\n"); 
    printf("%10.6f %10.6f\n", elapsedTime, 1e-6*n_bytes/elapsedTime); 

    // Destroy events 
    HANDLE_ERROR(cudaEventDestroy(start)); 
    HANDLE_ERROR(cudaEventDestroy(stop)); 

    // Free memory 
    HANDLE_ERROR(cudaFree(gpu_mem)); 

    return 0; 
} 

Answer 1

nvprof 프로파일과 API 프로파일
는 다른 결과를주고있다 :

$ nvprof --events gst_request ./app 
======== NVPROF is profiling app... 
======== Command: app 
#Delay(ms) BW(GB/s) 
13.345920 80.454690 
======== Profiling result: 
      Invocations  Avg  Min  Max Event Name 
Device 0 
    Kernel: kernel_write(unsigned int*, unsigned int) 
        1 8388608 8388608 8388608 gst_request 

$ nvprof --events global_store_transaction ./app 
======== NVPROF is profiling app... 
======== Command: app 
#Delay(ms) BW(GB/s) 
    9.469216 113.392892 
======== Profiling result: 
      Invocations  Avg  Min  Max Event Name 
Device 0 
    Kernel: kernel_write(unsigned int*, unsigned int) 
        1 8257560 8257560 8257560 global_store_transaction 

나는 global_store_transation이 gst_request보다 낮은 수 없다는 인상을 가지고 있었다. 여기서 무슨 일이 일어나고있는거야? 동일한 명령에서 두 이벤트를 모두 요청할 수는 없으므로 두 개의 별도 명령을 실행해야했습니다. 이것이 문제가 될 수 있습니까?

이상하게도 API 프로파일 러는 완전한 통합으로 다른 결과를 보여줍니다. 여기에 출력이, 나는 적절한 카운터를 얻기 위해 두 번 실행했다 :

$ cat config.txt 
inst_issued 
inst_executed 
gst_request 

$ COMPUTE_PROFILE=1 COMPUTE_PROFILE_CSV=1 COMPUTE_PROFILE_LOG=log.csv COMPUTE_PROFILE_CONFIG=config.txt ./app 

$ cat log.csv 
# CUDA_PROFILE_LOG_VERSION 2.0 
# CUDA_DEVICE 0 GeForce GTX 580 
# CUDA_CONTEXT 1 
# CUDA_PROFILE_CSV 1 
# TIMESTAMPFACTOR fffff67eaca946b8 
method,gputime,cputime,occupancy,inst_issued,inst_executed,gst_request,gld_request 
_Z12kernel_writePjj,7771.776,7806.000,1.000,4737053,3900426,557058,0 

$ cat config2.txt 
global_store_transaction 

$ COMPUTE_PROFILE=1 COMPUTE_PROFILE_CSV=1 COMPUTE_PROFILE_LOG=log2.csv COMPUTE_PROFILE_CONFIG=config2.txt ./app 

$ cat log2.csv 
# CUDA_PROFILE_LOG_VERSION 2.0 
# CUDA_DEVICE 0 GeForce GTX 580 
# CUDA_CONTEXT 1 
# CUDA_PROFILE_CSV 1 
# TIMESTAMPFACTOR fffff67eea92d0e8 
method,gputime,cputime,occupancy,global_store_transaction 
_Z12kernel_writePjj,7807.584,7831.000,1.000,557058 

여기 gst_request 및 global_store_transactions 완벽 병합을 보여주는 정확히 동일합니다. 어떤 것이 올바른지 (nvprof 또는 API 프로파일 러)? 왜 NVIDIA Visual Profiler에서 필자와 공동 작성을 한 적이 없다고 말합니까? 여전히 중요한 지침 재생이 있으며 어디에서 왔는지 전혀 알 수 없습니다. (

아이디어가 있습니까? 동일한 기계에 두 개의 보드가 있고 둘 다 같은 것을 보여주기 때문에 하드웨어가 고장 나는 것 같지는 않습니다. 동작.