Code GPU with CUDA

Identifying performance limiters

previous

Timing: What to measure?

• Wall time: user will see this time
• GPU time: specific kernel time
• CPU ⇔ GPU memory transfers time:
  • not considered for GPU time analysis
  • significantly impact wall time
• Data dependent cases timing:
  • worst case time
  • time of single iteration
  • consider probability

Why to profile?

Profiler will not do your work for you,
but profiler helps:

• to verify memory access patterns
• to identify bottlenecks
• to collect statistic in data-dependent workloads
• to check your hypothesis
• to understand how hardware behaves

Think about profiling and benchmarking
as about scientific experiments

Device code profiler

• events are hardware counters, usually reported per SM
  • SM id selected by profiler with assumption that all SMs do approximately the same amount of work
  • Exceptions: L2 and DRAM counters
• metrics computed from number of events and hardware specific properties (e.c. number of SM)
• Single run can collect only a few counters
  • Profiler repeats kernel launches to collect all counters
• Results may vary for repeated runs

instructions/bytes ratio

• Profiler counters:
  • instructions_issued, instructions_executed
  • incremented by warp, but “issued” includes replays
  • global_store_transaction, uncached_global_load_transaction
  • transaction can be 32,64,128 byte. Requires additional analysis to determine average.
• Compute ratio:
  (warpSize X instructions_issued) v.s. (global_store_transaction + l1_global_load_miss) * avgTransactionSize

ROI profiling

#include <cuda_profiler_api.h>

// algorithm setup code
udaProfilerStart();
perf_test_cuda_accelerated_code();
cudaProfilerStop();

• Profile only part that you are optimizing right now
• shorter and simpler profiler log
• Do not significantly overhead your code runtime
• Used with --profile-from-start off nvprof option

Code paths analysis

• The main idea: determine performance limiters through measuring different parts independently
• Simple case: time memory-only and math-only versions of the kernel
• Shows how well memory operations are overlapped with arithmetic: compare the sum of mem-only and math-only times to full-kernel time

template<typename T>
__global__ void
benchmark_contiguous_direct_load(T* s, typename T::value_type* r, bool doStore)
{
   int global_index = threadIdx.x + blockDim.x * blockIdx.x;
   T data = s[global_index];
   asm (""::: "memory");
   if (s && doStore)
       r[global_index] = sum(data);
}

device side timing

• Device timer located on ROP/SM depending on hardware revision
• It's relatively easy to compute per thread values but hard to analyze kernel performance due to grid serialization
• sometimes is suitable for benchmarking

template<typename T, typename D, typename L>__global__
void latency_kernel(T** a, int len, int stride, int inner_its,
    D* latency, L func)
{
    D start_time, end_time;
    volatile D sum_time = 0;
    for (int k = 0; k < inner_its; ++k)
    {
        T *j = ((T*) a) + threadIdx.y * len + threadIdx.x;
        start_time = clock64();
        for (int curr = 0; curr < len / stride; ++curr) j = func(j);
        end_time = clock64(); sum_time += (end_time - start_time);
    }
    if (!threadIdx.x) atomicAdd(latency, sum_time);
}

Final words

• Time
• Profile
• (Micro)benchmark
• Prototype
• Look into SASS

THE END

list of presentations

BY cuda.geek / 2013–2015