nppiCopyConstBorder_8u_C1R 的性能下降
Posted
技术标签:
【中文标题】nppiCopyConstBorder_8u_C1R 的性能下降【英文标题】:Performance drop in nppiCopyConstBorder_8u_C1R 【发布时间】: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 跟踪。 【参考方案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
的热身通话。当我运行它时(在 sm_30 设备上使用 linux 的 CUDA 5.5),我看到了:
~$ 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
即。添加预热调用完全改变了函数的定时性能。当我查看来自 nvprof
的 API 跟踪时,我发现这两个 npp 函数调用都需要大约 6 微秒。但是,第一次调用的 CUDA 启动需要数百毫秒,而第二次调用大约需要 12 微秒。
所以,正如我在之前的评论中提到的,有一些惰性过程被包含在 Titan 案例的 CUDA 5.5 的时间安排中,而费米案例的 CUDA 5.0 可能没有。但这不是 npp 的特性,因为我猜实际功能的性能在 Titan 上与在 Fermi 卡上一样快或更快。
【讨论】:
你完全正确。但是,我尝试过相同的方法,但使用规范的方式创建 CUDA 上下文 ()***.com/questions/10415204/… 和 ***.com/questions/13313930/…) 并且行为是相同的。第一次调用 NPP 库中的函数是否可能需要不同的上下文初始化? 我已经用PS
更新了关于 TITAN 配置的问题(尽管我没有看到任何缺点)。
在第一条评论中,我的意思是行为与问题相同,而不是您的答案:)。以防万一。
@pQB:很可能。如果这让您感到困扰,您可能想通过错误报告向 NVIDIA 投诉。就我个人而言,我不喜欢惰性驱动程序 API 行为的想法。在我的代码中,我希望有可预测、可重复的 API 延迟,而这种行为恰恰相反。
@JonathanCohen 使用 cudaFree(0) 不会改变我们正在测试的 NPP 示例的行为。问题仍然存在。这肯定与第一次调用 NPP 库有关以上是关于nppiCopyConstBorder_8u_C1R 的性能下降的主要内容,如果未能解决你的问题,请参考以下文章
33.Python面向对象类的专有方法__iter____getitem____getattr____call____new____init__
33.Python面向对象类的专有方法__iter____getitem____getattr____call____new____init__