我正在 CUDA C 编程世界迈出我的第一步!
作为第一个测试，我编写了简单的算法来对图像进行灰度转换和阈值处理(我是计算机视觉和 OpenCV 的粉丝!)。 我决定将我的 CUDA 性能结果与 CPU 上的类似算法以及相应的 OpenCV (cpu) 函数进行比较。 这是全高清视频的结果:
Frame Count: 4754
Frame Resolution: 1920x1080
Total time CPU: 67418.6 ms
Frame Avg CPU: 14.1814 ms
Frame Count: 4754
Frame Resolution: 1920x1080
Total time OpenCV: 23805.3 ms
Frame Avg OpenCV: 5.00742 ms
Frame Count: 4754
Frame Resolution: 1920x1080
==6149== NVPROF is profiling process 6149, command: ./OpenCV_test
Total time CUDA: 28018.2 ms
Frame Avg CUDA: 5.89361 ms
==6149== Profiling application: ./OpenCV_test
==6149== Profiling result:
Time(%) Time Calls Avg Min Max Name
55.45% 4.05731s 4754 853.45us 849.54us 1.1141ms doThreshold(unsigned char const *, unsigned char*, unsigned int, unsigned int, unsigned int)
34.03% 2.49028s 4754 523.83us 513.67us 1.3338ms [CUDA memcpy HtoD]
10.52% 769.46ms 4754 161.85us 161.15us 301.06us [CUDA memcpy DtoH]
==6149== API calls:
Time(%) Time Calls Avg Min Max Name
80.11% 8.19501s 9508 861.91us 490.81us 2.7719ms cudaMemcpy
12.82% 1.31106s 9508 137.89us 66.639us 218.56ms cudaMalloc
5.74% 587.05ms 9508 61.742us 39.566us 2.0234ms cudaFree
1.21% 124.16ms 4754 26.116us 16.990us 365.86us cudaLaunch
0.06% 5.7645ms 23770 242ns 97ns 106.27us cudaSetupArgument
0.05% 5.4291ms 4754 1.1410us 602ns 10.150us cudaConfigureCall
0.01% 594.89us 83 7.1670us 249ns 282.44us cuDeviceGetAttribute
0.00% 45.536us 1 45.536us 45.536us 45.536us cuDeviceTotalMem
0.00% 35.649us 1 35.649us 35.649us 35.649us cuDeviceGetName
0.00% 1.8960us 2 948ns 345ns 1.5510us cuDeviceGetCount
0.00% 892ns 2 446ns 255ns 637ns cuDeviceGet
如您所见，OpenCV 比我的 cpu 实现好得多，也比我的 Cuda 算法好得多!诀窍在哪里？我怀疑 OpenCV 使用了一些特殊的 cpu 硬件指令集。 我对 CUDA 有更多的期待:人们谈论原始图像处理的 20x-30x 加速!我错过了什么？
这里有一些关于我的系统配置的细节:
这里是关于我的 OpenCV 3.0 构建的一些信息:
下面是为测试执行的代码:
#include <iostream>
#include <numeric>
#include <string>
#include <stdlib.h>
#include <chrono>
#include <opencv2/opencv.hpp>
using namespace cv;
using namespace std;
using namespace std::chrono;
const char* file = "PATH TO A VIDEO FILE";
__global__ void doThreshold(const uchar* bgrInput, uchar* output, uint inputSize, uint soglia, uint maxVal)
{
uint i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < inputSize)
{
output[i] = 0.5f + ((bgrInput[3 * i] + bgrInput[3 * i + 1] + bgrInput[3 * i + 2]) / 3.0f); // gray conversion
output[i] = output[i] > soglia ? maxVal : 0; // thresholding
}
}
void cudaCvtThreshold(const Mat& mat, Mat& result, uint soglia, uint maxVal)
{
if (mat.type() == CV_8UC3)
{
uint size = mat.rows * mat.cols;
uint blockSize = 128; // no significant result varying this variable
uint gridSize = ceil(size/(float)blockSize);
uchar* d_bgrInput, *d_output;
cudaMalloc((void**)&d_bgrInput, mat.channels() * size);
cudaMalloc((void**)&d_output, size);
cudaMemcpy(d_bgrInput, mat.data, mat.channels() * size, cudaMemcpyHostToDevice);
doThreshold<<<gridSize, blockSize>>>(d_bgrInput, d_output, size, soglia, maxVal);
result = Mat(mat.rows, mat.cols, CV_8UC1);
cudaMemcpy(result.data, d_output, size, cudaMemcpyDeviceToHost);
cudaFree(d_bgrInput);
cudaFree(d_output);
}
else
cerr << "Only CV_8UC3 matrix supported" << endl;
}
void cpuCvtThreshold(const Mat& mat, Mat& result, uint soglia, uint maxVal)
{
if (mat.type() == CV_8UC3)
{
uint size = mat.rows * mat.cols;
result = Mat(mat.rows, mat.cols, CV_8UC1);
uchar* input = mat.data;
uchar* output = result.data;
for (uint i = 0; i < size; ++i)
{
output[i] = 0.5f + ((input[3 * i] + input[3 * i + 1] + input[3 * i + 2]) / 3.0f); // gray conversion
output[i] = output[i] > soglia ? maxVal : 0; // thresholding
}
}
else
cerr << "Only CV_8UC3 matrix supported" << endl;
}
void cudaTest(const string src)
{
VideoCapture cap(src);
Mat frame, result;
uint frameCount = cap.get(CAP_PROP_FRAME_COUNT);
cout << "Frame Count: " << frameCount << endl;
auto startTs = system_clock::now();
cap >> frame;
cout << "Frame Resolution: " << frame.cols << "x" << frame.rows << endl;
while (not frame.empty()) {
cudaCvtThreshold(frame, result, 127, 255);
cap >> frame;
}
auto stopTs = system_clock::now();
auto diff = stopTs - startTs;
auto elapsed = chrono::duration_cast<chrono::microseconds>(diff).count() / (double)1e3;
cout << "Total time CUDA: " << elapsed << " ms" << endl;
cout << "Frame Avg CUDA: " << elapsed / frameCount << " ms" << endl << endl;
}
void naiveCpu(const string src)
{
VideoCapture cap(src);
Mat frame, result;
uint frameCount = cap.get(CAP_PROP_FRAME_COUNT);
cout << "Frame Count: " << frameCount << endl;
auto startTs = system_clock::now();
cap >> frame;
cout << "Frame Resolution: " << frame.cols << "x" << frame.rows << endl;
while (not frame.empty()) {
cpuCvtThreshold(frame, result, 127, 255);
cap >> frame;
}
auto stopTs = system_clock::now();
auto diff = stopTs - startTs;
auto elapsed = chrono::duration_cast<chrono::microseconds>(diff).count() / (double)1e3;
cout << "Total time CPU: " << elapsed << " ms" << endl;
cout << "Frame Avg CPU: " << elapsed / frameCount << " ms" << endl << endl;
}
void opencv(const string src)
{
VideoCapture cap(src);
Mat frame, result;
uint frameCount = cap.get(CAP_PROP_FRAME_COUNT);
cout << "Frame Count: " << frameCount << endl;
auto startTs = system_clock::now();
cap >> frame;
cout << "Frame Resolution: " << frame.cols << "x" << frame.rows << endl;
while (not frame.empty()) {
cv::cvtColor(frame, result, COLOR_BGR2GRAY);
threshold(result, result, 127, 255, THRESH_BINARY);
cap >> frame;
}
auto stopTs = system_clock::now();
auto diff = stopTs - startTs;
auto elapsed = chrono::duration_cast<chrono::microseconds>(diff).count() / (double)1e3;
cout << "Total time OpenCV: " << elapsed << " ms" << endl;
cout << "Frame Avg OpenCV: " << elapsed / frameCount << " ms" << endl << endl;
}
int main(void)
{
naiveCpu(file);
opencv(file);
cudaTest(file);
return 0;
}
编辑:
添加/修改代码
__global__ void doThreshold(const uchar* bgrInput, uchar* output, uint inputSize, uint soglia, uint maxVal)
{
uint i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < inputSize)
{
uchar grayPix = 0.5f + ((bgrInput[3 * i] + bgrInput[3 * i + 1] + bgrInput[3 * i + 2]) / 3.0f); // gray conversion
output[i] = grayPix > soglia ? maxVal : 0; // thresholding
}
}
void cudaCvtThreshold(const Mat& mat, Mat& result, uint soglia, uint maxVal, uchar* d_bgrInput, uchar* d_output)
{
uint size = mat.rows * mat.cols;
uint blockSize = 128; // no significant result varying this variable
uint gridSize = ceil(size/(float)blockSize);
doThreshold<<<gridSize, blockSize>>>(d_bgrInput, d_output, size, soglia, maxVal);
}
void cudaTestOutMallocFree(const string src)
{
VideoCapture cap(src);
Mat frame;
uint frameCount = cap.get(CAP_PROP_FRAME_COUNT);
cout << "Frame Count: " << frameCount << endl;
auto startTs = system_clock::now();
cap >> frame;
cout << "Frame Resolution: " << frame.cols << "x" << frame.rows << endl;
uint size = frame.rows * frame.cols;
Mat result(frame.rows, frame.cols, CV_8UC1);
uchar* d_bgrInput, *d_output;
cudaMalloc((void**)&d_bgrInput, frame.channels() * size);
cudaMalloc((void**)&d_output, size);
while (not frame.empty())
{
cudaMemcpy(d_bgrInput, frame.data, frame.channels() * size, cudaMemcpyHostToDevice);
cudaCvtThreshold(frame, result, 127, 255, d_bgrInput, d_output);
cudaMemcpy(result.data, d_output, size, cudaMemcpyDeviceToHost);
cap >> frame;
}
cudaFree(d_bgrInput);
cudaFree(d_output);
auto stopTs = system_clock::now();
auto diff = stopTs - startTs;
auto elapsed = chrono::duration_cast<chrono::microseconds>(diff).count() / (double)1e3;
cout << "Total time CUDA (out malloc-free): " << elapsed << " ms" << endl;
cout << "Frame Avg CUDA (out malloc-free): " << elapsed / frameCount << " ms" << endl << endl;
}
int main(void)
{
naiveCpu(file);
opencv(file);
cudaTest(file);
cudaTestOutMallocFree(file);
return 0;
}
和结果:
Frame Count: 4754
Frame Resolution: 1920x1080
Total time CPU: 70972.6 ms
Frame Avg CPU: 14.929 ms
Frame Count: 4754
Frame Resolution: 1920x1080
Total time OpenCV: 23475.4 ms
Frame Avg OpenCV: 4.93804 ms
Frame Count: 4754
Frame Resolution: 1920x1080
==4493== NVPROF is profiling process 4493, command: ./OpenCV_test
Total time CUDA: 27451.3 ms
Frame Avg CUDA: 5.77435 ms
Frame Count: 4754
Frame Resolution: 1920x1080
Total time CUDA (out malloc-free): 26137.3 ms
Frame Avg CUDA (out malloc-free): 5.49796 ms
==4493== Profiling application: ./OpenCV_test
==4493== Profiling result:
Time(%) Time Calls Avg Min Max Name
53.74% 7.53280s 9508 792.26us 789.61us 896.17us doThreshold(unsigned char const *, unsigned char*, unsigned int, unsigned int, unsigned int)
35.57% 4.98604s 9508 524.40us 513.54us 979.37us [CUDA memcpy HtoD]
10.69% 1.49876s 9508 157.63us 157.09us 206.24us [CUDA memcpy DtoH]
==4493== API calls:
Time(%) Time Calls Avg Min Max Name
88.22% 15.7392s 19016 827.68us 482.18us 1.7570ms cudaMemcpy
7.07% 1.26081s 9510 132.58us 65.458us 198.86ms cudaMalloc
3.26% 582.24ms 9510 61.223us 39.675us 304.16us cudaFree
1.33% 236.64ms 9508 24.888us 13.497us 277.21us cudaLaunch
0.06% 10.667ms 47540 224ns 96ns 347.09us cudaSetupArgument
0.06% 9.9587ms 9508 1.0470us 504ns 9.4800us cudaConfigureCall
0.00% 428.88us 83 5.1670us 225ns 228.70us cuDeviceGetAttribute
0.00% 43.388us 1 43.388us 43.388us 43.388us cuDeviceTotalMem
0.00% 34.389us 1 34.389us 34.389us 34.389us cuDeviceGetName
0.00% 1.7010us 2 850ns 409ns 1.2920us cuDeviceGetCount
0.00% 821ns 2 410ns 225ns 596ns cuDeviceGet
单个 malloc 和 free 的性能更好，但改进很小......
编辑2:
按照 Jez 的建议，我修改了 Cuda 内核，以便在每个 GPU 线程内处理多个像素(在以下执行中为 8 个):
修改后的代码:
__global__ void doThreshold(const uchar* bgrInput, uchar* output, uint inputSize, uint soglia, uint maxVal, uint pixelPerThread)
{
uint i = pixelPerThread * (blockIdx.x * blockDim.x + threadIdx.x);
if (i < inputSize)
{
for (uint j = 0; j < pixelPerThread; j++) {
uchar grayPix = 0.5f + ( (bgrInput[3 * (i + j)] + bgrInput[3 * (i + j) + 1] + bgrInput[3 * (i + j) + 2]) / 3.0f ); // gray conversion
output[i + j] = grayPix > soglia ? maxVal : 0; // thresholding
}
}
}
void cudaCvtThreshold(const Mat& mat, Mat& result, uint soglia, uint maxVal, uchar* d_bgrInput, uchar* d_output)
{
uint size = mat.rows * mat.cols;
uint pixelPerThread = 8;
uint blockSize = 128; // no significant result varying this variable
uint gridSize = ceil(size/(float)(blockSize * pixelPerThread));
doThreshold<<<gridSize, blockSize>>>(d_bgrInput, d_output, size, soglia, maxVal, pixelPerThread);
}
然后结果:
Frame Count: 4754
Frame Resolution: 1920x1080
Total time OpenCV: 23628.8 ms
Frame Avg OpenCV: 4.97031 ms
Frame Count: 4754
Frame Resolution: 1920x1080
==13441== NVPROF is profiling process 13441, command: ./OpenCV_test
Total time CUDA (out malloc-free): 25655.5 ms
Frame Avg CUDA (out malloc-free): 5.39662 ms
==13441== Profiling application: ./OpenCV_test
==13441== Profiling result:
Time(%) Time Calls Avg Min Max Name
49.30% 3.15853s 4754 664.39us 658.24us 779.04us doThreshold(unsigned char const *, unsigned char*, unsigned int, unsigned int, unsigned int, unsigned int)
38.69% 2.47838s 4754 521.32us 513.35us 870.69us [CUDA memcpy HtoD]
12.01% 769.53ms 4754 161.87us 161.31us 200.58us [CUDA memcpy DtoH]
==13441== API calls:
Time(%) Time Calls Avg Min Max Name
95.78% 7.26387s 9508 763.97us 491.11us 1.6589ms cudaMemcpy
2.51% 190.70ms 2 95.350ms 82.529us 190.62ms cudaMalloc
1.53% 116.31ms 4754 24.465us 16.844us 286.56us cudaLaunch
0.09% 6.7052ms 28524 235ns 98ns 233.19us cudaSetupArgument
0.08% 5.9538ms 4754 1.2520us 642ns 12.039us cudaConfigureCall
0.00% 263.87us 83 3.1790us 225ns 111.03us cuDeviceGetAttribute
0.00% 174.45us 2 87.227us 52.521us 121.93us cudaFree
0.00% 34.612us 1 34.612us 34.612us 34.612us cuDeviceTotalMem
0.00% 29.376us 1 29.376us 29.376us 29.376us cuDeviceGetName
0.00% 1.6950us 2 847ns 343ns 1.3520us cuDeviceGetCount
0.00% 745ns 2 372ns 217ns 528ns cuDeviceGet
请注意，内核执行的平均时间现在是 664,39 us 而不是 792,26 us 不错! :-) 但是 OpenCV(使用英特尔 IPP)仍然更快!
编辑3: 我在没有 IPP 和各种 SSE 指令的情况下重新编译了 OpenCV。 OpenCV 的性能似乎是一样的!!
Frame Count: 4754
Frame Resolution: 1920x1080
Total time OpenCV: 23541.7 ms
Frame Avg OpenCV: 4.95198 ms
最佳答案
这里有两件事。
您花费大约一半的 GPU 时间在 GPU 之间分配和复制内存。 CPU-GPU 连接是一个相对较慢的链接，与数据在 GPU 上开始和结束并且内存分配一次的情况相比，性能会立即减半。您可以在这里做一些事情来提供帮助，例如将分配移到循环之外，并将一帧的数据传输与下一帧的计算重叠，但是复制->执行->复制的模式很少产生很好的结果运行时，除非执行非常复杂。
您的内核应该是内存限制的。您(理想情况下)移动 4 个字节/线程，使用约 200 万个线程(像素)和 853us 的运行时间，您将获得大约 10GB/s。 GTX 970 的峰值为 224GB/s。你还有很长的路要走。
这里的问题是您正在进行 8 位交易。这种情况下的解决方案是使用共享内存。如果您在内核开始时以高性能方式将数据加载到共享内存中(例如，将指针转换为 int4，确保对齐)，则您可以从该内存中读取，然后以 32+ 位每个写回线。这意味着您必须在一个线程中处理多个像素，但这不是问题。
另一种解决方案是找到一个库来执行此操作。 NPP ，例如，涵盖了许多与图像相关的任务，并且可能比手写代码更快。
有了良好的内存访问模式，我希望这个内核能快 10 倍以上。根据 Amdahl 定律，一旦你这样做了，你就会被开销所支配，所以除非你能摆脱它们，否则运行时间只会快约 2 倍。
关于c++ - 计算机视觉算法的 CUDA 性能，我们在Stack Overflow上找到一个类似的问题： https://stackoverflow.com/questions/33777338/
我的瘦服务器配置了nginx，我的ROR应用程序正在它们上运行。在我发布代码更新时运行thinrestart会给我的应用程序带来一些停机时间。我试图弄清楚如何优雅地重启正在运行的Thin实例，但找不到好的解决方案。有没有人能做到这一点？ 最佳答案 #Restartjustthethinserverdescribedbythatconfigsudothin-C/etc/thin/mysite.ymlrestartNginx将继续运行并代理请求。如果您将Nginx设置为使用多个上游服务器，例如server{listen80;server
这里是Ruby新手。完成一些练习后碰壁了。练习:计算一系列成绩的字母等级创建一个方法get_grade来接受测试分数数组。数组中的每个分数应介于0和100之间，其中100是最大分数。计算平均分并将字母等级作为字符串返回，即“A”、“B”、“C”、“D”、“E”或“F”。我一直返回错误:avg.rb:1:syntaxerror,unexpectedtLBRACK,expecting')'defget_grade([100,90,80])^avg.rb:1:syntaxerror,unexpected')',expecting$end这是我目前所拥有的。我想坚持使用下面的方法或.join，
目录一.加解密算法数字签名对称加密DES(DataEncryptionStandard)3DES(TripleDES)AES(AdvancedEncryptionStandard)RSA加密法DSA(DigitalSignatureAlgorithm)ECC(EllipticCurvesCryptography)非对称加密签名与加密过程非对称加密的应用对称加密与非对称加密的结合二.数字证书图解一.加解密算法加密简单而言就是通过一种算法将明文信息转换成密文信息，信息的的接收方能够通过密钥对密文信息进行解密获得明文信息的过程。根据加解密的密钥是否相同，算法可以分为对称加密、非对称加密、对称加密和非
如何将send与+=一起使用？a=20;a.send"+=",10undefinedmethod`+='for20:Fixnuma=20;a+=10=>30 最佳答案 恐怕你不能。+=不是方法，而是语法糖。参见http://www.ruby-doc.org/docs/ProgrammingRuby/html/tut_expressions.html它说Incommonwithmanyotherlanguages,Rubyhasasyntacticshortcut:a=a+2maybewrittenasa+=2.你能做的最好的事情是:
项目介绍随着我国经济迅速发展,人们对手机的需求越来越大,各种手机软件也都在被广泛应用,但是对于手机进行数据信息管理,对于手机的各种软件也是备受用户的喜爱小学生兴趣延时班预约小程序的设计与开发被用户普遍使用,为方便用户能够可以随时进行小学生兴趣延时班预约小程序的设计与开发的数据信息管理,特开发了小程序的设计与开发的管理系统。小学生兴趣延时班预约小程序的设计与开发的开发利用现有的成熟技术参考,以源代码为模板,分析功能调整与小学生兴趣延时班预约小程序的设计与开发的实际需求相结合,讨论了小学生兴趣延时班预约小程序的设计与开发的使用。开发环境开发说明：前端使用微信微信小程序开发工具：后端使用ssm：VU
我对如何计算通过{%assignvar=0%}赋值的变量加一完全感到困惑。这应该是最简单的任务。到目前为止，这是我尝试过的:{%assignamount=0%}{%forvariantinproduct.variants%}{%assignamount=amount+1%}{%endfor%}Amount:{{amount}}结果总是0。也许我忽略了一些明显的东西。也许有更好的方法。我想要存档的只是获取运行的迭代次数。 最佳答案 因为{{incrementamount}}将输出您的变量值并且不会影响{%assign%}定义的变量，我
给定一个nxmbool数组:[[true,true,false],[false,true,true],[false,true,true]]有什么简单的方法可以返回“该列中有多少个true？”结果应该是[1,3,2] 最佳答案 使用转置得到一个数组，其中每个子数组代表一列，然后将每一列映射到其中的true数:arr.transpose.map{|subarr|subarr.count(true)}这是一个带有inject的版本，应该在1.8.6上运行，没有任何依赖:arr.transpose.map{|subarr|subarr.in
给定两个大小相等的数组，如何找到不考虑位置的匹配元素的数量？例如:[0,0,5]和[0,5,5]将返回2的匹配项，因为有一个0和一个5共同；[1,0,0,3]和[0,0,1,4]将返回3的匹配项，因为0有两场，1有一场；[1,2,2,3]和[1,2,3,4]将返回3的匹配项。我尝试了很多想法，但它们都变得相当粗糙和令人费解。我猜想有一些不错的Ruby习惯用法，或者可能是一个正则表达式，可以很好地回答这个解决方案。 最佳答案 您可以使用count完成它:a.count{|e|index=b.index(e)andb.delete_at
我有一个数组数组，想将元素附加到子数组。+=做我想做的，但我想了解为什么push不做。我期望的行为(并与+=一起工作):b=Array.new(3,[])b[0]+=["apple"]b[1]+=["orange"]b[2]+=["frog"]b=>[["苹果"],["橙子"],["Frog"]]通过推送，我将推送的元素附加到每个子数组(为什么？):a=Array.new(3,[])a[0].push("apple")a[1].push("orange")a[2].push("frog")a=>[[“苹果”、“橙子”、“Frog”]、[“苹果”、“橙子”、“Frog”]、[“苹果”、“
我正在使用Ruby解决一些ProjectEuler问题，特别是这里我要讨论的问题25(Fibonacci数列中包含1000位数字的第一项的索引是多少？)。起初，我使用的是Ruby2.2.3，我将问题编码为:number=3a=1b=2whileb.to_s.length但后来我发现2.4.2版本有一个名为digits的方法，这正是我需要的。我转换为代码:whileb.digits.length当我比较这两种方法时，digits慢得多。时间./025/problem025.rb0.13s用户0.02s系统80%cpu0.190总计./025/problem025.rb2.19s用户0.0