- c - 在位数组中找到第一个零
- linux - Unix 显示有关匹配两种模式之一的文件的信息
- 正则表达式替换多个文件
- linux - 隐藏来自 xtrace 的命令
我的问题:我正在寻找有人指出我尝试在 CUDA 中使用零复制的方式中的错误,或者揭示一个更“幕后”的视角来解释为什么零复制方法不会比 memcpy 方法更快.顺便说一句,我正在使用 Ubuntu 在 NVidia 的 TK1 处理器上执行我的测试。
我的问题与有效使用 NVIDIA TK1 的(物理上)统一内存架构与 CUDA 有关。 NVIDIA 为 GPU/CPU 内存传输抽象提供了两种方法。
我的测试代码的简短描述:我使用方法 1 和方法 2 测试同一个 cuda 内核。我预计方法 1 会更快,因为没有将源数据复制到设备或从设备复制结果数据.然而,结果倒退到我的假设(方法#1 慢了 50%)。下面是我的测试代码:
#include <libfreenect/libfreenect.hpp>
#include <iostream>
#include <vector>
#include <cmath>
#include <pthread.h>
#include <cxcore.h>
#include <time.h>
#include <sys/time.h>
#include <memory.h>
///CUDA///
#include <cuda.h>
#include <cuda_runtime.h>
///OpenCV 2.4
#include <highgui.h>
#include <cv.h>
#include <opencv2/gpu/gpu.hpp>
using namespace cv;
using namespace std;
///The Test Kernel///
__global__ void cudaCalcXYZ( float *dst, float *src, float *M, int height, int width, float scaleFactor, int minDistance)
{
float nx,ny,nz, nzpminD, jFactor;
int heightCenter = height / 2;
int widthCenter = width / 2;
//int j = blockIdx.x; //Represents which row we are in
int index = blockIdx.x*width;
jFactor = (blockIdx.x - heightCenter)*scaleFactor;
for(int i= 0; i < width; i++)
{
nz = src[index];
nzpminD = nz + minDistance;
nx = (i - widthCenter )*(nzpminD)*scaleFactor;
ny = (jFactor)*(nzpminD);
//Solve for only Y matrix (height vlaues)
dst[index++] = nx*M[4] + ny*M[5] + nz*M[6];
//dst[index++] = 1 + 2 + 3;
}
}
//Function fwd declarations
double getMillis();
double getMicros();
void runCudaTestZeroCopy(int iter, int cols, int rows);
void runCudaTestDeviceCopy(int iter, int cols, int rows);
int main(int argc, char **argv) {
//ZERO COPY FLAG (allows runCudaTestZeroCopy to run without fail)
cudaSetDeviceFlags(cudaDeviceMapHost);
//Runs kernel using explicit data copy to 'device' and back from 'device'
runCudaTestDeviceCopy(20, 640,480);
//Uses 'unified memory' cuda abstraction so device can directly work from host data
runCudaTestZeroCopy(20,640, 480);
std::cout << "Stopping test" << std::endl;
return 0;
}
void runCudaTestZeroCopy(int iter, int cols, int rows)
{
cout << "CUDA Test::ZEROCOPY" << endl;
int src_rows = rows;
int src_cols = cols;
int m_rows = 4;
int m_cols = 4;
int dst_rows = src_rows;
int dst_cols = src_cols;
//Create and allocate memory for host mats pointers
float *psrcMat;
float *pmMat;
float *pdstMat;
cudaHostAlloc((void **)&psrcMat, src_rows*src_cols*sizeof(float), cudaHostAllocMapped);
cudaHostAlloc((void **)&pmMat, m_rows*m_cols*sizeof(float), cudaHostAllocMapped);
cudaHostAlloc((void **)&pdstMat, dst_rows*dst_cols*sizeof(float), cudaHostAllocMapped);
//Create mats using host pointers
Mat src_mat = Mat(cvSize(src_cols, src_rows), CV_32FC1, psrcMat);
Mat m_mat = Mat(cvSize(m_cols, m_rows), CV_32FC1, pmMat);
Mat dst_mat = Mat(cvSize(dst_cols, dst_rows), CV_32FC1, pdstMat);
//configure src and m mats
for(int i = 0; i < src_rows*src_cols; i++)
{
psrcMat[i] = (float)i;
}
for(int i = 0; i < m_rows*m_cols; i++)
{
pmMat[i] = 0.1234;
}
//Create pointers to dev mats
float *d_psrcMat;
float *d_pmMat;
float *d_pdstMat;
//Map device to host pointers
cudaHostGetDevicePointer((void **)&d_psrcMat, (void *)psrcMat, 0);
//cudaHostGetDevicePointer((void **)&d_pmMat, (void *)pmMat, 0);
cudaHostGetDevicePointer((void **)&d_pdstMat, (void *)pdstMat, 0);
//Copy matrix M to device
cudaMalloc( (void **)&d_pmMat, sizeof(float)*4*4 ); //4x4 matrix
cudaMemcpy( d_pmMat, pmMat, sizeof(float)*m_rows*m_cols, cudaMemcpyHostToDevice);
//Additional Variables for kernels
float scaleFactor = 0.0021;
int minDistance = -10;
//Run kernel! //cudaSimpleMult( float *dst, float *src, float *M, int width, int height)
int blocks = src_rows;
const int numTests = iter;
double perfStart = getMillis();
for(int i = 0; i < numTests; i++)
{
//cudaSimpleMult<<<blocks,1>>>(d_pdstMat, d_psrcMat, d_pmMat, src_cols, src_rows);
cudaCalcXYZ<<<blocks,1>>>(d_pdstMat, d_psrcMat, d_pmMat, src_rows, src_cols, scaleFactor, minDistance);
cudaDeviceSynchronize();
}
double perfStop = getMillis();
double perfDelta = perfStop - perfStart;
cout << "Ran " << numTests << " iterations totaling " << perfDelta << "ms" << endl;
cout << " Average time per iteration: " << (perfDelta/(float)numTests) << "ms" << endl;
//Copy result back to host
//cudaMemcpy(pdstMat, d_pdstMat, sizeof(float)*src_rows*src_cols, cudaMemcpyDeviceToHost);
//cout << "Printing results" << endl;
//for(int i = 0; i < 16*16; i++)
//{
// cout << "src[" << i << "]= " << psrcMat[i] << " dst[" << i << "]= " << pdstMat[i] << endl;
//}
cudaFree(d_psrcMat);
cudaFree(d_pmMat);
cudaFree(d_pdstMat);
cudaFreeHost(psrcMat);
cudaFreeHost(pmMat);
cudaFreeHost(pdstMat);
}
void runCudaTestDeviceCopy(int iter, int cols, int rows)
{
cout << "CUDA Test::DEVICE COPY" << endl;
int src_rows = rows;
int src_cols = cols;
int m_rows = 4;
int m_cols = 4;
int dst_rows = src_rows;
int dst_cols = src_cols;
//Create and allocate memory for host mats pointers
float *psrcMat;
float *pmMat;
float *pdstMat;
cudaHostAlloc((void **)&psrcMat, src_rows*src_cols*sizeof(float), cudaHostAllocMapped);
cudaHostAlloc((void **)&pmMat, m_rows*m_cols*sizeof(float), cudaHostAllocMapped);
cudaHostAlloc((void **)&pdstMat, dst_rows*dst_cols*sizeof(float), cudaHostAllocMapped);
//Create pointers to dev mats
float *d_psrcMat;
float *d_pmMat;
float *d_pdstMat;
cudaMalloc( (void **)&d_psrcMat, sizeof(float)*src_rows*src_cols );
cudaMalloc( (void **)&d_pdstMat, sizeof(float)*src_rows*src_cols );
cudaMalloc( (void **)&d_pmMat, sizeof(float)*4*4 ); //4x4 matrix
//Create mats using host pointers
Mat src_mat = Mat(cvSize(src_cols, src_rows), CV_32FC1, psrcMat);
Mat m_mat = Mat(cvSize(m_cols, m_rows), CV_32FC1, pmMat);
Mat dst_mat = Mat(cvSize(dst_cols, dst_rows), CV_32FC1, pdstMat);
//configure src and m mats
for(int i = 0; i < src_rows*src_cols; i++)
{
psrcMat[i] = (float)i;
}
for(int i = 0; i < m_rows*m_cols; i++)
{
pmMat[i] = 0.1234;
}
//Additional Variables for kernels
float scaleFactor = 0.0021;
int minDistance = -10;
//Run kernel! //cudaSimpleMult( float *dst, float *src, float *M, int width, int height)
int blocks = src_rows;
double perfStart = getMillis();
for(int i = 0; i < iter; i++)
{
//Copty from host to device
cudaMemcpy( d_psrcMat, psrcMat, sizeof(float)*src_rows*src_cols, cudaMemcpyHostToDevice);
cudaMemcpy( d_pmMat, pmMat, sizeof(float)*m_rows*m_cols, cudaMemcpyHostToDevice);
//Run Kernel
//cudaSimpleMult<<<blocks,1>>>(d_pdstMat, d_psrcMat, d_pmMat, src_cols, src_rows);
cudaCalcXYZ<<<blocks,1>>>(d_pdstMat, d_psrcMat, d_pmMat, src_rows, src_cols, scaleFactor, minDistance);
//Copy from device to host
cudaMemcpy( pdstMat, d_pdstMat, sizeof(float)*src_rows*src_cols, cudaMemcpyDeviceToHost);
}
double perfStop = getMillis();
double perfDelta = perfStop - perfStart;
cout << "Ran " << iter << " iterations totaling " << perfDelta << "ms" << endl;
cout << " Average time per iteration: " << (perfDelta/(float)iter) << "ms" << endl;
cudaFree(d_psrcMat);
cudaFree(d_pmMat);
cudaFree(d_pdstMat);
cudaFreeHost(psrcMat);
cudaFreeHost(pmMat);
cudaFreeHost(pdstMat);
}
//Timing functions for performance measurements
double getMicros()
{
timespec ts;
//double t_ns, t_s;
long t_ns;
double t_s;
clock_gettime(CLOCK_MONOTONIC, &ts);
t_s = (double)ts.tv_sec;
t_ns = ts.tv_nsec;
//return( (t_s *1000.0 * 1000.0) + (double)(t_ns / 1000.0) );
return ((double)t_ns / 1000.0);
}
double getMillis()
{
timespec ts;
double t_ns, t_s;
clock_gettime(CLOCK_MONOTONIC, &ts);
t_s = (double)ts.tv_sec;
t_ns = (double)ts.tv_nsec;
return( (t_s * 1000.0) + (t_ns / 1000000.0) );
}
我已经看过帖子Cuda zero-copy performance ,但我觉得这不相关,原因如下:GPU 和 CPU 具有物理上统一的内存架构。
谢谢
最佳答案
当您使用 ZeroCopy 时,对内存的读取会通过一些路径查询内存单元以从系统内存中获取数据。此操作有一些延迟。
当使用直接访问内存时,内存单元从全局内存中收集数据,并且具有不同的访问模式和延迟。
要真正看到这种差异,需要进行某种形式的分析。
尽管如此,您对全局函数的调用使用了单个线程
cudaCalcXYZ<<< blocks,1 >>> (...
在这种情况下,当从系统内存(或全局内存)收集内存时,GPU 几乎没有办法隐藏延迟。我建议您使用更多线程(64 的倍数,总共至少 128 个),并在其上运行探查器以获得内存访问成本。您的算法似乎是可分离的,并且从
修改代码for(int i= 0; i < width; i++)
到
for (int i = threadIdx.x ; i < width ; i += blockDim.x)
可能会提高整体性能。图像大小为 640 宽,这将变成 128 个线程的 5 次迭代。
cudaCalcXYZ<<< blocks,128 >>> (...
我相信这会带来一些性能提升。
关于c++ - CUDA 零拷贝与 Jetson TK1 上的 CudaMemcpy,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/36758620/
我想向 tk.Tk 和 tk.Toplevel 的实例添加 2 个方法。前者作为应用程序根窗口存在,后者由用户创建任意次数。每个 tk.Toplevel 代表应用程序的不同功能,目前有 13 种变体。
尝试创建 Listbox 的子类,以便我可以创建一个新的 KeyListbox from tkinter import * class KeyListbox(Listbox): def __
以下是错误: $ perl ftper.plUse of uninitialized value $id in hash element at /usr/lib/perl5/vendor_perl/5
我正在尝试使用 tk.Scale 来更改 tk.IntVar。我可以在第一个 tk 窗口上执行此操作,但不能在第二个窗口上执行。为什么? 看看这个可怜的样本: import tkinter as tk
我输入了我的框架比例,但我不确定如何在我的标签中显示这个比例的值。每次秤移动时我都需要更新它。我怎样才能做到这一点? self.options_settings.framepripojeni6 = F
这个video向我介绍了 X 的问题以及 Wayland 协议(protocol)的替代方案。多年来,Wayland 协议(protocol)的采用似乎在不断增长。 我的问题: tkinter 和 t
这些天我经常看到这一行:tk.Tk.__init__(self,*args,**kwargs)并且不明白它的用途。我的问题不是 *args 和 **kwargs)。 例如,在此处的这些代码行中: cl
这个问题涉及Python和Tkinter。 我想同时使用两个不同的 ttk 主题,例如一个 Tkinter 窗口中的“clam”和“vista”。所以我写了这个: import tkinter as
从 Tk 8.5 开始,Tk 开始使用基于原生的 UI 组件——按钮、滚动条等。我想知道 C++/Tk 是否支持这种原生 GUI 方法? 最佳答案 据我所知,C++/Tk 是基于 Tk 8.4(如果是
我无法让 ruby 成功地 require 'tk'。我正在使用 rvm、ruby 2.0.0、ActiveTcl-8.6 和 Ubuntu 12.04 LTS。我已经运行了随 ActiveTcl
我想用最有效的方式来限制用户不输入数字以外的任何内容。例如:当他们在条目中输入字母时,该条目将被清除。有没有什么方法可以用最少的结构改变来做到这一点? 这是我的代码:[??????标志是我被困的地方]
我真的很喜欢Perl/Tk ,但已经得出的意见认为它是 DOA。我认为 Tcl::Tk和 Tkx是更好的解决方案。假设我放弃 Perl/Tk . Tcl::Tk 是“最受支持”的路线吗? (自 200
我正在尝试使用 perl:latest 和 activestate/circleci-activeperl:latest 图像在 ubuntu:16.04 容器上运行使用 Tk 模块制作的简单 per
我曾经使用 tk.Scale 的 digits 属性来确保 Label 或 Spinbox 中的数字> 在 slider 移动时显示固定的小数位数。比如 3.456, 4444.567, 555555
问题就在标题中,本质上是:如何在设置 Entry 的 textvariable 后继续调用 validatecommand 回调?这是最低工作示例 (MWE): import tkinter as t
在Perl/Tk应用程序中,我想将带有任意Unicode文件名的文件拖放到小工具上。。有一个正在运行的DropSite示例:Windows资源管理器中的perl tk拖放文件夹,但它不能处理包含Uni
我在 macOS BigSur 上通过 VMWare 使用 Ubuntu 20.04.2 LTS。我安装了最新版本的 tcl、tcl-dev、tk 和 tk-dev - 版本 8.6。我想编译 Arc
我组装 Tk 窗口的方式有问题(在 Win XP 下使用 R tcltk 和 tcltk2) library(tcltk) library(tcltk2) expandTk <- function()
我想在 Textdocument 的特定行中写入,但我的代码有问题,我不知道错误在哪里。 set fp [open C:/Users/user/Desktop/tst/settings.txt w]
当在 TK 中创建新的顶层或按钮时,需要输入路径名。我看过一个基本代码,如下所示: toplevel .a ... button .a.b ... 我的问题是:点与字母的处理方式是否不同?它们是某种创
我是一名优秀的程序员,十分优秀!