性能下降nppiCopyConstBorder_8u_C1R

Question

在兩個不同的體系結構（GTX480和GTX TITAN）中，使用nppiCopyConstBorder_8u_C1R函數的性能下降，也涉及到不同的CUDA版本（分別爲v5.0和v5.5）。

在第一種情況（GTX480和CUDA 5.0）的功能的執行時間是

T = 0.00005 seconds 

在第二種情況下（GTX TITAN和CUDA 5.5）的執行時間是

​​

我用以下代碼複製了此行爲：

// GTX480 nvcc -lnpp -m64 -O3 --ptxas-options=-v -gencode arch=compute_20,code=sm_20 --compiler-options -use_fast_math 
// GTXTITAN nvcc -lnppi -m64 -O3 --ptxas-options=-v -gencode arch=compute_35,code=sm_35 --compiler-options -use_fast_math 
#include <stdlib.h> 
#include <stdio.h> 
// CUDA 
#include <cuda.h> 
#include <cuda_runtime_api.h> 
// CUDA Nvidia Performance Primitives 
#include <npp.h> 

#include <assert.h> 

#define w 256 // width 
#define h 256 // height 
#define b 16 // extra border 

#define BORDER_TYPE 0 

int main(int argc, char *argv[]) 
{ 
    // input data 
    Npp8u* h_idata[w*h]; 
    // output data 
    Npp8u* h_odata[(w+b)*(h+b)]; 

    /* MEMORY ALLOCTION AND INITIAL COPY OF DATA FROM CPU TO GPU */ 

    Npp8u *i_devPtr, *i_devPtr_Border; 

    // size of input the data 
    int d_Size = w * h * sizeof(Npp8u); 
    // allocate input data 
    CUDA_CHECK_RETURN(cudaMalloc((void**) &i_devPtr, d_Size)); 
    // copy initial data to GPU 
    CUDA_CHECK_RETURN(cudaMemcpy(i_devPtr, h_idata, d_Size, cudaMemcpyHostToDevice)); 

    // size of output the data 
    int d_Size_o = (w+b) * (h+b) * sizeof(Npp8u);  
    // allocation for input data with extended border 
    CUDA_CHECK_RETURN(cudaMalloc((void**) &i_devPtr_Border, d_Size_o)); 

    // create struct with ROI size given the current mask 
    NppiSize SizeROI = {w, h}; 

    NppiSize SizeROI_Border = { SizeROI.width + b, SizeROI.height + b }; 

    // create events 
    cudaEvent_t start, stop; 
    cudaEventCreate(&start); 
    cudaEventCreate(&stop); 

    // NPP Library Copy Constant Border 
    cudaEventRecord(start, 0); 
    NppStatus eStatusNPP = nppiCopyConstBorder_8u_C1R(i_devPtr,SizeROI.width, SizeROI, 
        i_devPtr_Border, SizeROI_Border.width, SizeROI_Border, 
        b, b, BORDER_TYPE); 

    cudaDeviceSynchronize(); 
    assert(NPP_NO_ERROR == eStatusNPP); 
    cudaEventRecord(stop, 0); 
    cudaEventSynchronize(stop); 

    float milliseconds = 0; 
    cudaEventElapsedTime(&milliseconds, start, stop); 
    printf("T= %1.5f sg\n", milliseconds/1000.0f); 


    // copy output data from GPU 
    CUDA_CHECK_RETURN(cudaMemcpy(h_odata, i_devPtr_Border, d_Size_o, cudaMemcpyDeviceToHost)); 

    /* free resources */ 
    cudaFree(i_devPtr); 
    cudaFree(i_devPtr_Border); 

    CUDA_CHECK_RETURN(cudaDeviceReset()); 

    return 0; 
} 

問：任何人都知道這個問題？

這使我問以下問題：

問：如何nppiCopyConstBorder_8u_C1R實施？該功能是否涉及將數據從設備複製到主機，擴展主機中的邊界並將結果複製到設備？

PS：帶有TITAN的機器將GPU安裝在分離的主板上，專門設計用於多個PCIe連接，並通過PCIe線連接。在我已經測試的其他內核的配置中，我沒有看到任何缺陷。

Answer 1

我想你會發現唯一的區別是在程序執行過程中何時/何地計算API延遲，並且底層npp函數本身在兩個CUDA版本和GPU之間的性能差別不大架構。

我對這個假設的證據是這個版本發佈的代碼的：

#include <stdlib.h> 
#include <stdio.h> 
#include <cuda.h> 
#include <cuda_runtime_api.h> 
#include <npp.h> 

#include <assert.h> 

#define w 256 // width 
#define h 256 // height 
#define b 16 // extra border 

#define BORDER_TYPE 0 

#define CUDA_CHECK_RETURN(ans) { gpuAssert((ans), __FILE__, __LINE__); } 
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true) 
{ 
    if (code != cudaSuccess) 
    { 
     fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); 
     if (abort) exit(code); 
    } 
} 

int main(int argc, char *argv[]) 
{ 
    Npp8u* h_idata[w*h]; 
    Npp8u* h_odata[(w+b)*(h+b)]; 
    Npp8u *i_devPtr, *i_devPtr_Border; 

    int d_Size = w * h * sizeof(Npp8u); 
    CUDA_CHECK_RETURN(cudaMalloc((void**) &i_devPtr, d_Size)); 
    CUDA_CHECK_RETURN(cudaMemcpy(i_devPtr, h_idata, d_Size, cudaMemcpyHostToDevice)); 

    int d_Size_o = (w+b) * (h+b) * sizeof(Npp8u);  
    CUDA_CHECK_RETURN(cudaMalloc((void**) &i_devPtr_Border, d_Size_o)); 

    NppiSize SizeROI = {w, h}; 
    NppiSize SizeROI_Border = { SizeROI.width + b, SizeROI.height + b }; 
    NppStatus eStatusNPP; 

#ifdef __WARMUP_CALL__ 
    // Warm up call to nppi function 
    eStatusNPP = nppiCopyConstBorder_8u_C1R(i_devPtr,SizeROI.width, SizeROI, 
        i_devPtr_Border, SizeROI_Border.width, SizeROI_Border, 
        b, b, BORDER_TYPE); 

    assert(NPP_NO_ERROR == eStatusNPP); 
    CUDA_CHECK_RETURN(cudaDeviceSynchronize()); 
#endif 

    // Call for timing 
    cudaEvent_t start, stop; 
    CUDA_CHECK_RETURN(cudaEventCreate(&start)); 
    CUDA_CHECK_RETURN(cudaEventCreate(&stop)); 

    CUDA_CHECK_RETURN(cudaEventRecord(start, 0)); 
    eStatusNPP = nppiCopyConstBorder_8u_C1R(i_devPtr,SizeROI.width, SizeROI, 
        i_devPtr_Border, SizeROI_Border.width, SizeROI_Border, 
        b, b, BORDER_TYPE); 

    assert(NPP_NO_ERROR == eStatusNPP); 
    CUDA_CHECK_RETURN(cudaEventRecord(stop, 0)); 
    CUDA_CHECK_RETURN(cudaEventSynchronize(stop)); 

    float milliseconds = 0; 
    cudaEventElapsedTime(&milliseconds, start, stop); 
    printf("T= %1.5f sg\n", milliseconds/1000.0f); 

    CUDA_CHECK_RETURN(cudaMemcpy(h_odata, i_devPtr_Border, d_Size_o, cudaMemcpyDeviceToHost)); 

    cudaFree(i_devPtr); 
    cudaFree(i_devPtr_Border); 

    CUDA_CHECK_RETURN(cudaDeviceReset()); 

    return 0; 
} 

注意熱身調用nppiCopyConstBorder_8u_C1R定時調用之前。當我運行它（CUDA 5.5與sm_30設備上的Linux），我看到這個：

~$ nvcc -lnppi -m64 -O3 --ptxas-options=-v -gencode arch=compute_30,code=sm_30 --compiler-options -use_fast_math pqb.cc 
~$ ./a.out 
T= 0.39670 sg 

~$ nvcc -lnppi -m64 -O3 --ptxas-options=-v -gencode arch=compute_30,code=sm_30 --compiler-options -use_fast_math -D__WARMUP_CALL__ pqb.cc 
~$ ./a.out 
T= 0.00002 sg 

ie。添加熱身呼叫完全改變了該功能的計時性能。當我查看來自nvprof的API跟蹤時，我發現兩個npp函數調用都需要大約6微秒。但是，當第二次調用需要大約12微秒時，第一次調用的CUDA啓動需要幾百毫秒。

因此，正如我在前面的評論中提到的那樣，有一些懶惰的過程被納入CUDA 5.5關於Titan案例的時間，可能不在CUDA 5.0上的費米案例中。但這不是npp的一個特性，因爲我猜測Titan上的實際函數的性能與Fermi卡相比速度更快或更快。