性能下降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卡相比速度更快或更快。