Точный метод вычисления двойной задержки FMA и общей памяти

Question

Я пытаюсь найти точный способ измерения латентности двух операций: 1) Задержка операции двойной точности FMA. 2) Задержка двойной точности загрузки из общей памяти. Я использую K20x и задавался вопросом, даст ли этот код точные измерения.

#include <cuda.h> 

#include <stdlib.h> 
#include <stdio.h> 
#include <iostream> 

using namespace std; 

//Clock rate 
#define MHZ 732e6 
//number of streaming multiprocessors 
#define SMS 14 
// number of double precision units 
#define DP_UNITS 16*4 
//number of shared banks 
#define SHARED_BANKS 32 

#define ITER 100000 
#define NEARONE 1.0000000000000004 

__global__ void fma_latency_kernal(double *in, double *out){ 
    int tid = blockIdx.x*blockDim.x+threadIdx.x; 
    double val = in[tid]; 
#pragma unroll 100 
    for(int i=0; i<ITER; i++){ 
    val+=val*NEARONE; 
    } 
    out[tid]=val; 
} 

__global__ void shared_latency_kernel(double *in, double *out){ 
    volatile extern __shared__ double smem[]; 
    int tid = blockIdx.x*blockDim.x+threadIdx.x; 
    smem[threadIdx.x]=in[tid]; 
#pragma unroll 32 
    for(int i=0; i<ITER; i++){ 
    smem[threadIdx.x]=smem[(threadIdx.x+i)%32]*NEARONE; 
    } 
    out[tid]=smem[threadIdx.x]; 
} 

int main (int argc , char **argv){ 

    float time; 
    cudaEvent_t start, stop, start2, stop2; 

    double *d_A, *d_B; 
    cudaMalloc(&d_A, DP_UNITS*SMS*sizeof(float)); 
    cudaMalloc(&d_B, DP_UNITS*SMS*sizeof(float)); 

    cudaError_t err; 

    cudaEventCreate(&start); 
    cudaEventCreate(&stop); 
    cudaEventRecord(start, 0); 

    fma_latency_kernal<<<SMS, DP_UNITS>>>(d_A, d_B); 

    cudaEventRecord(stop, 0); 
    cudaEventSynchronize(stop); 
    cudaEventElapsedTime(&time, start, stop); 
    time/=1000; 
    err = cudaGetLastError(); 
    if(err!=cudaSuccess) 
    printf("Error FMA: %s\n", cudaGetErrorString(err)); 
    printf("Latency of FMA = %3.1f clock cycles\n", (time/(double)ITER)*(double)MHZ); 


    cudaDeviceSetSharedMemConfig(cudaSharedMemBankSizeFourByte); 
    cudaEventCreate(&start2); 
    cudaEventCreate(&stop2); 
    cudaEventRecord(start2, 0); 

    shared_latency_kernel<<<1, SHARED_BANKS, sizeof(double)>>>(d_A, d_B); 

    cudaEventRecord(stop2, 0); 
    cudaEventSynchronize(stop2); 
    cudaEventElapsedTime(&time, start2, stop2); 
    time/=1000; 
    err = cudaGetLastError(); 
    if(err!=cudaSuccess) 
    printf("Error Shared Memory: %s\n", cudaGetErrorString(err)); 

    printf("Latency of Shared Memory = %3.1f clock cycles\n", time/(double)ITER*(double)MHZ); 

} 

Мои результаты на K20x являются следующие: латентности FMA = 16,4 тактов Задержка разделяемой памяти = 60.7 тактовых циклов Это кажется разумным для меня, но я не уверен, насколько точны это.

Answer 1

Ваши значения латентности выглядят очень высоко для меня - почти вдвое больше, чем я ожидал. Чтобы измерить, сколько циклов требуется на GPU, вы можете вставлять функции clock() до и после соответствующей части функции ядра. Функция часов возвращает текущий цикл как int, поэтому путем вычитания первого значения из второго вы получаете количество циклов, которое передавалось между отправкой первой команды синхронизации и диспетчеризацией второй команды синхронизации.

Обратите внимание, что цифры, полученные вами по этому методу, будут включать дополнительное время из самих инструкций часов; Я считаю, что по умолчанию поток будет блокировать несколько циклов непосредственно перед и после каждой инструкции синхронизации, поэтому вы можете поэкспериментировать с этим, чтобы узнать, сколько циклов оно добавляет, чтобы вы могли вычесть их обратно.