【发布时间】：2014-05-19 14:30:14
【问题描述】：
在涉及不同 CUDA 版本（分别为 v5.0 和 v5.5）的两种不同架构（GTX480 和 GTX TITAN）中使用 nppiCopyConstBorder_8u_C1R 函数时性能下降。
在第一种情况下（GTX480 和 CUDA 5.0）函数的执行时间是
T = 0.00005 seconds
在第二种情况下（GTX TITAN 和 CUDA 5.5），执行时间是
T = 0.969831 seconds
我已使用以下代码重现了此行为：
// 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 线连接。对于我测试过的其他内核，我没有发现此配置有任何缺点。
【问题讨论】：
-
您可以尝试使用 nvprof 运行 API 跟踪吗？我猜你的时间可能是进程生命周期早期发生的事情的受害者，现在在内核启动时懒惰地发生。问题是内核函数仍然需要几微秒，但运行它的 cuLuanch 需要数百毫秒。
-
@talonmies 我会检查两台机器上的 API 跟踪。