0

以下のコードを実行して、NVIDIA Visual Profilerのグローバル メモリに 1 GB を書き込むと、次の結果が得 られ
ます 。 メモリ書き込みは結合されるはずであり、カーネルに発散はありません。そのため、問題は、グローバル メモリ リプレイ オーバーヘッドがどこから発生しているのかということです。nvidia-cuda-toolkit バージョン 5.0.35-4ubuntu1 を使用して、これを Ubuntu 13.04 で実行しています。





#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;
}
4

1 に答える 1

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 より低くできないという印象を受けました。ここで何が起こっているのですか?同じコマンドで両方のイベントを要求することはできないため、2 つの別々のコマンドを実行する必要がありました。これが問題でしょうか?

不思議なことに、API プロファイラーは完全な合体で異なる結果を示します。出力は次のとおりです。適切なカウンターを取得するには、2 回実行する必要がありました。

$ 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 が結合していない書き込みがあると表示するのはなぜですか? まだ重要な命令のリプレイがあり、それらがどこから来ているのかわかりません:(

何か案は?同じマシンに 2 つのボードがあり、どちらも同じ動作を示すため、これはハードウェアの誤動作ではないと思います。

于 2013-06-26T14:19:21.837 に答える