原文：
annas-archive.org/md5/b5d2ddddf00cdfdea66355a7259934ba译者：飞龙
第九章：在 Jetson TX1 上部署计算机视觉应用
上一章介绍了在 Jetson TX1 开发板上安装 OpenCV 和 CUDA 的过程。本章将描述如何使用这些功能。将详细描述 Jetson TX1 GPU 的属性，这些属性使其适用于并行处理。本章还将描述如何在我们这本书中之前看到的 CUDA 和 C++代码在 Jetson TX1 上执行。它还将展示 Jetson TX1 GPU 在执行 CUDA 代码时的性能。本章的主要动机将是展示如何使用 Jetson TX1 部署图像和视频处理应用。以基本的图像处理应用，如图像读取、显示、加法、阈值和滤波为例，来展示如何使用 Jetson TX1 进行计算机视觉应用。此外，摄像头接口对于在实际场景中部署该板非常重要。本章将描述使用板载摄像头或 USB 摄像头进行视频捕获和处理应用的步骤。本章的最后部分将解释如何部署一些高级应用，如人脸检测和背景减法。
本章将涵盖以下主题：
-
Jetson TX1 板的设备属性
-
在 Jetson TX1 板上运行 CUDA 程序
-
在 Jetson TX1 板上进行图像处理
-
将摄像头与 Jetson TX1 开发板连接
-
在 Jetson TX1 开发板上执行高级应用，如人脸检测、眼检测和背景减法
技术要求
本章需要具备对 OpenCV、CUDA 和任何编程语言的良好理解。它还需要任何 Nvidia GPU 开发板，如 Jetson TK1、TX1 或 TX2。本章使用的代码文件可以从以下 GitHub 链接下载：github.com/PacktPublishing/Hands-On-GPU-Accelerated-Computer-Vision-with-OpenCV-and-CUDA。
观看以下视频以查看代码的实际运行情况：
Jetson TX1 GPU 的设备属性
CUDA 提供了一个简单的接口来确定 GPU 设备的性能，这是 Jetson TX1 板上存在的 Tegra X1。了解设备的属性对于编写针对它的优化程序非常重要。查找设备属性的程序包含在 JetPack 安装的 CUDA 示例程序中，位于主文件夹中。您还可以运行我们在第二章中开发的程序来查找设备属性。
程序在 Nvidia Tegra X1 GPU 上的输出如下：
JetPack 3.3 安装了 CUDA 9.0 运行时版本。GPU 设备的全局内存大约为 4 GB，GPU 时钟速度约为 1 GHz。这个时钟速度比本书前面提到的 GeForce 940 GPU 慢。内存时钟速度仅为 13 MHz，而 GeForce 940 为 2.505 GHz，这使得 Jetson TX1 较慢。与 GeForce 940 的 1 MB 相比，L2 缓存为 256 KB。大多数其他属性与 GeForce 940 相似。
在 X、Y 和 Z 方向上，每个块可以启动的最大线程数分别为 1,024、1,024 和 64。在确定从程序中启动的并行线程数量时，应使用这些数字。在启动每个网格的并行块数量时，也应采取相同的谨慎措施。
总结来说，我们已经看到了 Jetson TX1 开发板上可用的 Tegra X1 GPU 的设备属性。它是一个嵌入式板，因此内存可用，时钟速度相对于笔记本电脑中配备的 GPU 设备（如 GeForce 940）较慢。然而，它比 Arduino 和 Raspberry Pi 等嵌入式平台快得多。它可以很容易地用于部署需要高性能计算能力的计算机视觉应用。现在我们已经看到了设备属性，我们将从在 Jetson TX1 上使用 CUDA 开发第一个程序开始。
基于 Jetson TX1 的 CUDA 基础程序
在本节中，我们通过添加两个大数组的示例来展示使用 Jetson TX1 开发板执行 CUDA 程序。程序的性能也使用 CUDA 事件进行了测量。
添加具有 50,000 个元素的两个大数组的内核函数如下：
#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>
//Defining number of elements in Array
#define N 50000
//Defining Kernel function for vector addition
__global__ void gpuAdd(int *d_a, int *d_b, int *d_c) {
//Getting Thread index of current kernel
int tid = threadIdx.x + blockIdx.x * blockDim.x;
while (tid < N)
{
d_c[tid] = d_a[tid] + d_b[tid];
tid += blockDim.x * gridDim.x;
}
}
内核函数接收两个设备指针，这些指针指向输入数组作为输入，以及一个设备指针，该指针指向设备内存中输出数组的参数。计算当前内核执行的线程 ID，并由内核将线程索引索引的数组元素相加。如果启动的内核数量少于数组元素数量，则相同的内核将在while循环中添加由块维度偏移的Array元素。添加两个数组的main函数如下：
int main(void)
{
//Defining host arrays
int h_a[N], h_b[N], h_c[N];
//Defining device pointers
int *d_a, *d_b, *d_c;
cudaEvent_t e_start, e_stop;
cudaEventCreate(&e_start);
cudaEventCreate(&e_stop);
cudaEventRecord(e_start, 0);
// allocate the memory
cudaMalloc((void**)&d_a, N * sizeof(int));
cudaMalloc((void**)&d_b, N * sizeof(int));
cudaMalloc((void**)&d_c, N * sizeof(int));
//Initializing Arrays
for (int i = 0; i < N; i++) {
h_a[i] = 2 * i*i;
h_b[i] = i;
}
// Copy input arrays from host to device memory
cudaMemcpy(d_a, h_a, N * sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(d_b, h_b, N * sizeof(int), cudaMemcpyHostToDevice);
//Calling kernels passing device pointers as parameters
gpuAdd << <1024, 1024 >> >(d_a, d_b, d_c);
//Copy result back to host memory from device memory
cudaMemcpy(h_c, d_c, N * sizeof(int), cudaMemcpyDeviceToHost);
cudaDeviceSynchronize();
cudaEventRecord(e_stop, 0);
cudaEventSynchronize(e_stop);
float elapsedTime;
cudaEventElapsedTime(&elapsedTime, e_start, e_stop);
printf("Time to add %d numbers: %3.1f ms\n",N, elapsedTime);
定义了两个主机数组，并使用cudaMalloc函数为它们分配内存。它们被初始化为一些随机值，并上传到设备内存。创建了两个 CUDA 事件来测量 CUDA 程序的性能。内核以并行方式启动了 1,024 个块，每个块有 1,024 个线程。这些数字来自设备属性，如上一节所述。内核函数的结果被传输到主机内存。内核函数的时间由e_start和e_stop事件在内核启动前后记录。函数所花费的时间显示在控制台上。
以下代码被添加以验证由 GPU 计算的结果的正确性，并清理程序使用的内存：
int Correct = 1;
printf("Vector addition on GPU \n");
//Printing result on console
for (int i = 0; i < N; i++) {
if ((h_a[i] + h_b[i] != h_c[i]))
{
Correct = 0;
}
}
if (Correct == 1)
{
printf("GPU has computed Sum Correctly\n");
}
else
{
printf("There is an Error in GPU Computation\n");
}
//Free up memory
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
return 0;
}
在 CPU 上执行相同的数组加法操作，并与从 GPU 获得的结果进行比较，以验证 GPU 是否正确计算了结果。这也在控制台上显示。通过使用cudaFree函数释放程序使用的所有内存。
需要从终端运行以下两个命令来执行程序。程序应位于当前工作目录中：
$ nvcc 01_performance_cuda_events.cu -o gpu_add
$ ./gpu_add
使用nvcc命令通过 Nvidia CUDA 编译器编译 CUDA 代码。文件名作为命令的参数传递。编译器将创建的目标文件名通过-o选项指定。此文件名将用于执行程序。这是通过第二个命令完成的。程序输出如下：
从结果可以看出，Jetson TX1 计算包含 50,000 个元素的两个数组的和需要3.4ms，这比本书第三章中使用的 GeForce 940 慢，但仍然比 CPU 上的顺序执行快。
总结来说，本节展示了在执行 CUDA 程序中使用 Jetson TX1 开发板的方法。语法与我们在这本书中之前看到的相同。因此，书中之前开发的全部 CUDA 程序都可以在 Jetson TX1 上执行，无需太多修改。程序执行的步骤也进行了描述。下一节将描述使用 Jetson TX1 进行图像处理应用的方法。
在 Jetson TX1 上进行的图像处理
本节将展示在部署图像处理应用中使用 Jetson TX1 的方法。我们还将再次使用 OpenCV 和 CUDA 来加速 Jetson TX1 上的计算机视觉应用。在上一章中，我们看到了包含 OpenCV 和 CUDA 的 JetPack 3.3 的安装过程。但在最新的 JetPack 中，OpenCV 没有编译带有 CUDA 支持，也没有 GStreamer 支持，这是从代码中访问摄像头所必需的。因此，删除 JetPack 附带的 OpenCV 安装，并编译带有 CUDA 和 GStreamer 支持的 OpenCV 新版本是一个好主意。下一节将展示如何执行此过程。
（编译带有 CUDA 支持的 OpenCV，如果需要）
尽管 JetPack 附带的 OpenCV 可以与新安装的 OpenCV 一起工作，但先删除旧安装然后开始新的安装是一个好主意。这将避免不必要的混淆。为了完成这个任务，必须执行以下步骤：
- 从终端运行以下命令：
$ sudo apt-get purge libopencv*
- 确保安装的所有软件包都是最新版本。如果不是这样，可以通过运行以下两个命令来更新它们：
$ sudo apt-get update
$ sudo apt-get dist-upgrade
- 编译从源代码编译 OpenCV 需要最新的 cmake 和 gcc 编译器版本，因此可以通过运行以下两个命令来安装：
$ sudo apt-get install --only-upgrade gcc-5 cpp-5 g++-5
$ sudo apt-get install build-essential make cmake cmake-curses-gui libglew-dev libgtk2.0-dev
- 需要安装一些依赖项以编译支持 GStreamer 的 OpenCV。这可以通过以下命令完成：
sudo apt-get install libdc1394-22-dev libxine2-dev libgstreamer1.0-dev libgstreamer-plugins-base1.0-dev
- 通过执行以下命令下载 OpenCV 最新版本的源代码，并将其提取到一个文件夹中：
$ wget https://github.com/opencv/opencv/archive/3.4.0.zip -O opencv.zip
$ unzip opencv.zip
- 现在，进入
opencv文件夹，创建build目录。然后进入这个新创建的build目录。这些可以通过从命令提示符执行以下命令来完成。
$ cd opencv
$ mkdir build
$ cd build
cmake命令用于编译具有 CUDA 支持的opencv。确保在此命令中将WITH_CUDA标志设置为ON。注意，对于 Jetson TX1 开发板，CUDA_ARCH_BIN应设置为5.3，对于 Jetson TX2，应设置为6.2。示例没有构建以节省时间和空间。整个cmake命令如下：
cmake -D CMAKE_BUILD_TYPE=RELEASE -D CMAKE_INSTALL_PREFIX=/usr/local \
-D WITH_CUDA=ON -D CUDA_ARCH_BIN="5.3" -D CUDA_ARCH_PTX="" \
-D WITH_CUBLAS=ON -D ENABLE_FAST_MATH=ON -D CUDA_FAST_MATH=ON \
-D ENABLE_NEON=ON -D WITH_LIBV4L=ON -D BUILD_TESTS=OFF \
-D BUILD_PERF_TESTS=OFF -D BUILD_EXAMPLES=OFF \
-D WITH_QT=ON -D WITH_OPENGL=ON ..
-
它将启动配置和创建
makefile的过程。在配置成功后，cmake命令将在build目录中创建makefile。 -
要使用
makefile编译 OpenCV，请在命令窗口中执行make -j4命令。 -
编译成功后，您必须从命令行执行
sudo make install命令来安装 OpenCV。
如果这些步骤执行成功，则 OpenCV 3.4.0 将在 Jetson TX1 上安装，并支持 CUDA 和 GStreamer，使用 OpenCV 制作的任何计算机视觉应用都可以部署到它上面。下一节将演示在板上进行简单的图像处理操作。
读取和显示图像
对于任何计算机视觉应用，所需的基本操作之一是读取和显示存储在磁盘上的图像。本节将演示一个简单的代码，在 Jetson TX1 上执行此操作。当我们从计算机上的 GPU 移动到 Jetson TX1 开发板时，OpenCV 语法不会改变太多。将有一些小的变化。在 Jetson TX1 上读取和显示图像的代码如下：
#include <opencv2/opencv.hpp>
#include <iostream>
using namespace cv;
using namespace std;
int main()
{
Mat img = imread("images/cameraman.tif",0);
if (img.empty())
{
cout << "Could not open an image" << endl;
return -1;
}
imshow("Image Read on Jetson TX1"; , img);
waitKey(0);
return 0;
}
必要的 OpenCV 库包含在代码中。图像是通过 Main 函数内的 imread 函数读取的。由于 imread 命令的第二个参数指定为 0，因此图像被读取为灰度图像。要将图像读取为彩色图像，可以指定为 1。if 语句检查图像是否被读取，如果没有，则在控制台上显示错误后终止代码。当图像名称不正确或图像未存储在指定的路径时，可能会发生读取图像的错误。这个错误由 if 语句处理。图像是通过 imshow 命令显示的。waitKey 函数用于显示图像，直到按下键盘上的任何键。
所示的前面代码可以保存为image_read.cpp文件，并使用以下命令从终端执行。请确保程序文件存储在终端的当前工作目录中：
For compilation:
$ g++ -std = c++11 image_read.cpp 'pkg_config --libs --cflags opencv' -o image_read
For execution:
$./image_read
程序的输出如下：
本节演示了在 Jetson TX1 上读取和显示图像的步骤。在下节中，我们将看到一些更多的图像处理操作，并尝试测量它们在 Jetson TX1 上的性能。
图像加法
本节将演示使用 Jetson TX1 进行简单的图像处理应用，如图像加法。在加法后，相同位置的像素强度被相加以构建新的图像。假设在两个图像中，(0,0)位置的像素强度值分别为 50 和 150，那么结果图像中的强度值将是 200，这是两个强度值的和。OpenCV 的加法操作是饱和操作，这意味着如果加法的结果超过 255，则将饱和在 255。在 Jetson TX1 上执行加法的代码如下：
#include <iostream>
#include "opencv2/opencv.hpp"
#include "opencv2/core/cuda.hpp"
int main (int argc, char* argv[])
{
//Read Two Images
cv::Mat h_img1 = cv::imread("images/cameraman.tif");
cv::Mat h_img2 = cv::imread("images/circles.png");
int64 work_begin = cv::getTickCount();
//Create Memory for storing Images on device
cv::cuda::GpuMat d_result1,d_img1, d_img2;
cv::Mat h_result1;
//Upload Images to device
d_img1.upload(h_img1);
d_img2.upload(h_img2);
cv::cuda::add(d_img1,d_img2, d_result1);
//Download Result back to host
d_result1.download(h_result1);
cv::imshow("Image1 ", h_img1);
cv::imshow("Image2 ", h_img2);
cv::imshow("Result addition ", h_result1);
int64 delta = cv::getTickCount() - work_begin;
//Frequency of timer
double freq = cv::getTickFrequency();
double work_fps = freq / delta;
std::cout<<"Performance of Addition on Jetson TX1: " <<std::endl;
std::cout <<"Time: " << (1/work_fps) <<std::endl;
std::cout <<"FPS: " <<work_fps <<std::endl;
cv::imshow("result_add.png", h_result1);
cv::waitKey();
return 0;
}
在进行图像加法时需要注意的一点是，两个图像应该具有相同的大小。如果不是这种情况，那么在加法之前应该将它们调整大小。在前面代码中，从磁盘读取了两个相同大小的图像，并将它们上传到设备内存中进行 GPU 上的加法。使用cv::cuda模块中的add函数在设备上执行图像加法。结果图像被下载到主机并在控制台上显示。
程序的输出如下：
使用cv::getTickCount()和cv::getTickFrequency()函数也测量了图像加法的性能。加法操作所需的时间显示在下面的屏幕截图上：
如前述屏幕截图所示，在 Jetson TX1 上添加两个大小为 256 x 256 的图像大约需要0.26ms。这对于嵌入式平台来说是一个非常不错的性能。需要注意的是，在测量加法操作的准确时间之前，应该测量imshow函数。imshow函数显示图像需要更多的时间，因此测量的时间不会是加法操作所需时间的准确估计。
图像阈值化
本节将演示使用 Jetson TX1 进行更复杂的计算机视觉应用，如图像阈值化。图像阈值化是一种非常简单的图像分割技术，用于根据某些强度值从灰度图像中提取重要区域。在这种技术中，如果像素值大于某个阈值值，则分配一个值，否则分配另一个值。
OpenCV 提供了不同类型的阈值技术，这由函数的最后一个参数决定。这些阈值类型包括：
-
cv:.THRES H_BINARY: 如果像素的强度大于阈值，则将像素强度设置为等于maxVal常量，否则将像素强度设置为等于零。 -
cv::THRESH_BINARY_INV: 如果像素的强度大于阈值，则将像素强度设置为等于零，否则将像素强度设置为maxVal常量。 -
cv::THRESH_TRUNC: 这基本上是一个截断操作。如果像素的强度大于阈值，则将像素强度设置为等于阈值，否则保持强度值不变。 -
cv::THRESH_TOZERO: 如果像素的强度大于阈值，则保持像素强度不变，否则将像素强度设置为等于零。 -
cv::THRESH_TOZERO_INV: 如果像素的强度大于阈值，则将该像素强度设置为等于零，否则保持像素强度不变。
在 Jetson TX1 上使用 OpenCV 和 CUDA 实现所有这些阈值技术的程序如下所示：
#include <iostream>
#include "opencv2/opencv.hpp"
using namespace cv;
int main (int argc, char* argv[])
{
cv::Mat h_img1 = cv::imread("images/cameraman.tif", 0);
cv::cuda::GpuMat d_result1,d_result2,d_result3,d_result4,d_result5, d_img1;
//Measure initial time ticks
int64 work_begin = getTickCount();
d_img1.upload(h_img1);
cv::cuda::threshold(d_img1, d_result1, 128.0, 255.0, cv::THRESH_BINARY);
cv::cuda::threshold(d_img1, d_result2, 128.0, 255.0, cv::THRESH_BINARY_INV);
cv::cuda::threshold(d_img1, d_result3, 128.0, 255.0, cv::THRESH_TRUNC);
cv::cuda::threshold(d_img1, d_result4, 128.0, 255.0, cv::THRESH_TOZERO);
cv::cuda::threshold(d_img1, d_result5, 128.0, 255.0, cv::THRESH_TOZERO_INV);
cv::Mat h_result1,h_result2,h_result3,h_result4,h_result5;
d_result1.download(h_result1);
d_result2.download(h_result2);
d_result3.download(h_result3);
d_result4.download(h_result4);
d_result5.download(h_result5);
//Measure difference in time ticks
int64 delta = getTickCount() - work_begin;
double freq = getTickFrequency();
//Measure frames per second
double work_fps = freq / delta;
std::cout <<"Performance of Thresholding on GPU: " <<std::endl;
std::cout <<"Time: " << (1/work_fps) <<std::endl;
std::cout <<"FPS: " <<work_fps <<std::endl;
return 0;
}
在 OpenCV 和 CUDA 的 GPU 上用于图像阈值化的函数是cv::cuda::threshold。此函数有许多参数。第一个参数是源图像，它应该是一个灰度图像。第二个参数是结果要存储的目标位置。第三个参数是阈值值，用于分割像素值。第四个参数是maxVal常量，表示如果像素值超过阈值值时赋予的值。最后一个参数是前面讨论过的阈值方法。以下程序显示了原始图像和五种阈值技术输出的输出：
使用cv::getTickCount()和cv::getTickFrequency()函数测量图像阈值化的性能。五个阈值操作所需的时间显示在控制台上，如下面的截图所示：
在 Jetson TX1 上执行五个阈值操作需要0.32ms，这对于嵌入式平台上的图像分割任务来说，性能非常好。下一节将描述 Jetson TX1 上的滤波操作。
Jetson TX1 上的图像滤波
图像滤波是图像预处理和特征提取中的一个非常重要的步骤。低通滤波器，如平均、高斯和中值滤波器，用于去除图像中的不同类型的噪声，而高通滤波器，如 Sobel、Scharr 和 Laplacian，用于检测图像中的边缘。边缘是重要的特征，可用于计算机视觉任务，如目标检测和分类。本书中已详细解释了图像滤波。
本节描述了在 Jetson TX1 上对图像应用低通和高通滤波器的步骤。相应的代码如下：
#include <iostream>
#include <string>
#include "opencv2/opencv.hpp"
using namespace std;
using namespace cv;
using namespace cv::cuda;
int main()
{
Mat h_img1;
cv::cuda::GpuMat d_img1,d_blur,d_result3x3;
h_img1 = imread("images/blobs.png",1);
int64 start = cv::getTickCount();
d_img1.upload(h_img1);
cv::cuda::cvtColor(d_img1,d_img1,cv::COLOR_BGR2GRAY);
cv::Ptr<cv::cuda::Filter> filter3x3;
filter3x3 = cv::cuda::createGaussianFilter(CV_8UC1,CV_8UC1,cv::Size(3,3),1);
filter3x3->apply(d_img1, d_blur);
cv::Ptr<cv::cuda::Filter> filter1;
filter1 = cv::cuda::createLaplacianFilter(CV_8UC1,CV_8UC1,1);
filter1->apply(d_blur, d_result3x3);
cv::Mat h_result3x3,h_blur;
d_result3x3.download(h_result3x3);
d_blur.download(h_blur);
double fps = cv::getTickFrequency() / (cv::getTickCount() - start);
std::cout << "FPS : " << fps << std::endl;
imshow("Laplacian", h_result3x3);
imshow("Blurred", h_blur);
cv::waitKey();
return 0;
}
拉普拉斯算子是一种二阶导数，用于从图像中提取垂直和水平图像。它对噪声非常敏感，因此有时需要使用低通滤波器（如高斯模糊）来去除噪声，然后再应用拉普拉斯滤波器。因此，在代码中，使用标准差等于1的 3x3 高斯滤波器对输入图像进行处理。该滤波器是通过 OpenCV 的cv::cuda::createGaussianFilter函数创建的。然后，将拉普拉斯滤波器应用于高斯模糊后的图像。拉普拉斯滤波器是通过 OpenCV 的cv::cuda::createLaplacianFilter函数创建的。高斯模糊和拉普拉斯滤波器的输出被下载回主机内存，以便在控制台上显示。代码中还测量了滤波操作的性能。程序的输出如下截图所示：
从输出中可以看出，对模糊图像应用拉普拉斯滤波器将去除图像中的虚假边缘。它还将去除输入图像中存在的高斯噪声。如果输入图像被盐和胡椒噪声扭曲，则应在拉普拉斯滤波器进行边缘检测之前使用中值滤波器作为预处理步骤。
总结来说，我们已经在 Jetson TX1 上看到了不同的图像处理函数，如图像加法、图像阈值和图像滤波。我们还看到，这些操作在 Jetson TX1 上的性能比在 CPU 上执行相同的代码要好得多。下一节将描述如何将摄像头与 Jetson TX1 连接，以便在现实生活中的场景中使用。
与 Jetson TX1 连接摄像头
Jetson TX1 可以与 USB 摄像头或 CSI 摄像头连接。开发板已预装了一款 5 兆像素的摄像头，并与 Jetson TX1 连接。这款摄像头可以像笔记本电脑上的网络摄像头一样捕获视频。摄像头连接是 Jetson TX1 开发板在实时应用中的重要特性，它支持多达六通道的摄像头。Jetson TX1 支持的摄像头详细列表可以在以下链接中找到：elinux.org/Jetson_TX1
本节将演示使用与 Jetson TX1 连接的摄像头捕获视频的步骤，以及如何使用这些视频开发计算机视觉应用，如人脸检测和背景减法。
从机载摄像头读取和显示视频
本节将描述从 USB 摄像头或与 Jetson TX1 连接的机载摄像头捕获视频的方法。为此，OpenCV 应编译为支持 GStreamer；否则，OpenCV 将不支持捕获视频的格式。
以下代码可以用来从摄像头捕获视频并在屏幕上显示：
#include <opencv2/opencv.hpp>
#include <iostream>
#include <stdio.h>
using namespace cv;
using namespace std;
int main(int, char**)
{
Mat frame;
// open the default camera using default API
VideoCapture cap("nvcamerasrc ! video/x-raw(memory:NVMM), width=(int)1280, height=(int)720, format=(string)I420, framerate=(fraction)24/1 ! nvvidconv flip-method=0 ! video/x-raw, format=(string)I420 ! videoconvert ! video/x-raw, format=(string)BGR ! appsink");
if (!cap.isOpened()) {
cout << "Unable to open camera\n";
return -1;
}
while (1)
{
int64 start = cv::getTickCount();
cap.read(frame);
// check if we succeeded
if (frame.empty()) {
cout << "Can not read frame\n";
break;
}
double fps = cv::getTickFrequency() / (cv::getTickCount() - start);
std::cout << "FPS : " << fps << std::endl;
imshow("Live", frame);
if (waitKey(30) == 'q')
break;
}
return 0;
}
代码与用于在台式机上从网络摄像头捕获视频的代码大致相似。而不是使用设备 ID 作为参数来捕获对象，使用指定 GStreamer 管道的字符串。如下所示：
VideoCapture cap("nvcamerasrc ! video/x-raw(memory:NVMM), width=(int)1280, height=(int)720, format=(string)I420, framerate=(fraction)24/1 ! nvvidconv flip-method=0 ! video/x-raw, format=(string)I420 ! videoconvert ! video/x-raw, format=(string)BGR ! appsink");
捕获的视频的宽度和高度被指定为 1,280 和 720 像素。帧率也被指定。这些值将根据接口摄像头的支持的格式而变化。使用nvvidconv将视频转换为 OpenCV 支持的 BGR 格式。它还用于图像缩放和翻转。要翻转捕获的视频，可以将 flip 方法指定为一个非零的整数值。
使用cap.isOpened属性来检查是否已从摄像头开始捕获。然后使用 read 方法逐个读取帧并在屏幕上显示，直到用户按下q键。代码中也测量了帧捕获的速率。
摄像头捕获了实时视频的两个不同帧，帧率显示在下面的屏幕截图上：
总结来说，在本节中，我们看到了从与 Jetson TX1 开发板相连的摄像头捕获视频的步骤。这个捕获的视频可以用于开发下一节所述的有用的实时计算机视觉应用。
Jetson TX1 上的高级应用
本节将描述在部署高级计算机视觉应用（如人脸检测、眼检测和背景减法）中使用 Jetson TX1 嵌入式平台。
使用 Haar 级联进行人脸检测
Haar 级联使用矩形特征来检测对象。它使用不同大小的矩形来计算不同的线和边缘特征。Haar-like 特征检测算法背后的思想是计算矩形内白色像素总和与黑色像素总和之间的差异。
这种方法的主要优势是使用积分图方法进行快速的求和计算。这使得 Haar 级联非常适合实时目标检测。它处理图像所需的时间比其他用于目标检测的算法要少。由于 Haar 级联具有低计算复杂性和低内存占用，因此非常适合部署在嵌入式系统如 Jetson TX1 上。因此，在本节中，使用此算法在 Jetson TX1 上部署人脸检测应用。
从与 Jetson TX1 接口的摄像头捕获的视频进行人脸检测的代码如下：
#include <iostream>
#include <opencv2/opencv.hpp>
using namespace cv;
using namespace std;
int main()
{
VideoCapture cap("images/output.avi");
//cv::VideoCapture cap("nvcamerasrc ! video/x-raw(memory:NVMM), width=(int)1280, height=(int)720, format=(string)I420, framerate=(fraction)24/1 ! nvvidconv flip-method=0 ! video/x-raw, format=(string)I420 ! videoconvert ! video/x-raw, format=(string)BGR ! appsink");
if (!cap.isOpened()) {
cout << "Can not open video source";
return -1;
}
std::vector<cv::Rect> h_found;
cv::Ptr<cv::cuda::CascadeClassifier> cascade = cv::cuda::CascadeClassifier::create("haarcascade_frontalface_alt2.xml");
cv::cuda::GpuMat d_frame, d_gray, d_found;
while(1)
{
Mat frame;
if ( !cap.read(frame) ) {
cout << "Can not read frame from webcam";
return -1;
}
int64 start = cv::getTickCount();
d_frame.upload(frame);
cv::cuda::cvtColor(d_frame, d_gray, cv::COLOR_BGR2GRAY);
cascade->detectMultiScale(d_gray, d_found);
cascade->convert(d_found, h_found);
for(int i = 0; i < h_found.size(); ++i)
{
rectangle(frame, h_found[i], Scalar(0,255,255), 5);
}
double fps = cv::getTickFrequency() / (cv::getTickCount() - start);
std::cout << "FPS : " << fps << std::endl;
imshow("Result", frame);
if (waitKey(1) == 'q') {
break;
}
}
return 0;
}
Haar 级联是一种需要训练以执行特定任务的算法。从头开始训练特定应用的 Haar 级联是困难的，因此 OpenCV 提供了一些训练好的 XML 文件，可用于检测对象。这些 XML 文件位于 OpenCV 和 CUDA 安装的\usr\local\opencv\data\haarcascades_cuda目录中。
初始化网络摄像头，并逐个捕获网络摄像头的帧。将帧上传到设备内存以在 GPU 上处理。OpenCV 和 CUDA 提供了CascadeClassifier类，可用于实现 Haar 级联。使用create方法创建该类的对象。它需要加载训练好的 XML 文件的文件名。
在while循环内部，将detectMultiscale方法应用于每一帧，以便在每一帧中检测不同大小的人脸。使用convert方法将检测到的位置转换为矩形向量。然后，使用for循环迭代此向量，以便在所有检测到的人脸上使用矩形函数绘制边界框。此过程对从网络摄像头捕获的每一帧重复进行。算法的性能也以每秒帧数来衡量。
程序的输出如下：
从输出中可以看出，人脸在两个不同位置的两个不同的网络摄像头帧中被正确定位。第二帧有点模糊，但这不会影响算法。Jetson TX1 上算法的性能也在右图显示。算法在大约每秒五帧的速度下工作。
总结来说，本节展示了使用 Jetson TX1 从网络摄像头捕获的实时视频中检测人脸。此应用程序可用于人员识别、人脸锁定、考勤监控等。
使用 Haar 级联进行眼检测
本节将描述使用 Haar 级联检测人类眼睛的使用。用于眼检测的训练好的 Haar 级联的 XML 文件位于 OpenCV 安装目录中。此文件用于检测眼睛。其代码如下：
#include "opencv2/objdetect/objdetect.hpp"
#include "opencv2/highgui/highgui.hpp"
#include "opencv2/imgproc/imgproc.hpp"
#include "opencv2/cudaobjdetect.hpp"
#include <iostream>
#include <stdio.h>
using namespace std;
using namespace cv;
int main( )
{
Mat h_image;
h_image = imread("images/lena_color_512.tif", 0);
Ptr<cuda::CascadeClassifier> cascade = cuda::CascadeClassifier::create("haarcascade_eye.xml");
cuda::GpuMat d_image;
cuda::GpuMat d_buf;
int64 start = cv::getTickCount();
d_image.upload(h_image);
cascadeGPU->setMinNeighbors(0);
cascadeGPU->setScaleFactor(1.02);
cascade->detectMultiScale(d_image, d_buf);
std::vector<Rect> detections;
cascade->convert(d_buf, detections);
if (detections.empty())
std::cout << "No detection." << std::endl;
cvtColor(h_image,h_image,COLOR_GRAY2BGR);
for(int i = 0; i < detections.size(); ++i)
{
rectangle(h_image, detections[i], Scalar(0,255,255), 5);
}
double fps = cv::getTickFrequency() / (cv::getTickCount() - start);
std::cout << "FPS : " << fps << std::endl;
imshow("Result image on Jetson TX1", h_image);
waitKey(0);
return 0;
}
代码与面部检测的代码类似。这是使用 Haar 级联的优势。如果有一个给定对象的训练好的 Haar 级联的 XML 文件可用，那么相同的代码将在所有应用程序中工作。只需在创建CascadeClassifier类的对象时更改 XML 文件的名称。在前面的代码中，使用了用于眼检测的训练 XML 文件haarcascade_eye.xml。其他代码是自解释的。缩放因子设置为1.02，以便在每次缩放时图像大小将减少1.02。眼检测程序的输出如下：
现在我们已经使用 Haar 级联从视频和图像中检测到对象，因此捕获的视频也可以使用下一节中描述的背景减法方法来检测和跟踪对象。
使用高斯混合（MoG）的背景减法
背景减法是目标检测和跟踪应用的重要预处理步骤。它也可以用于从监控录像中检测异常活动。本节展示了在背景减法应用中使用 Jetson TX1。与 Jetson TX1 接口的摄像头被安装在一个房间内，用于检测房间内的活动。房间的背景在第一帧中被初始化。
MoG 是一种广泛使用的背景减法方法，用于根据高斯混合将前景与背景分离，用于活动检测。背景从帧序列中持续更新。使用 K 个高斯分布的混合来将像素分类为前景或背景。帧的时间序列也被加权以改进背景建模。持续变化的强度被分类为前景，而静态的强度被分类为背景。
使用 MoG 进行活动监控的代码如下：
#include <iostream>
#include <string>
#include "opencv2/opencv.hpp"
using namespace std;
using namespace cv;
using namespace cv::cuda;
int main()
{
VideoCapture cap("nvcamerasrc ! video/x-raw(memory:NVMM), width=(int)1280, height=(int)720, format=(string)I420, framerate=(fraction)24/1 ! nvvidconv flip-method=0 ! video/x-raw, format=(string)I420 ! videoconvert ! video/x-raw, format=(string)BGR ! appsink");
if (!cap.isOpened())
{
cout << "Can not open camera or video file" << endl;
return -1;
}
Mat frame;
cap.read(frame);
GpuMat d_frame;
d_frame.upload(frame);
Ptr<BackgroundSubtractor> mog = cuda::createBackgroundSubtractorMOG();
GpuMat d_fgmask,d_fgimage,d_bgimage;
Mat h_fgmask,h_fgimage,h_bgimage;
mog->apply(d_frame, d_fgmask, 0.01);
namedWindow("image", WINDOW_NORMAL);
namedWindow("foreground mask", WINDOW_NORMAL);
namedWindow("foreground image", WINDOW_NORMAL);
namedWindow("mean background image", WINDOW_NORMAL);
while(1)
{
cap.read(frame);
if (frame.empty())
break;
d_frame.upload(frame);
int64 start = cv::getTickCount();
mog->apply(d_frame, d_fgmask, 0.01);
mog->getBackgroundImage(d_bgimage);
double fps = cv::getTickFrequency() / (cv::getTickCount() - start);
std::cout << "FPS : " << fps << std::endl;
d_fgimage.create(d_frame.size(), d_frame.type());
d_fgimage.setTo(Scalar::all(0));
d_frame.copyTo(d_fgimage, d_fgmask);
d_fgmask.download(h_fgmask);
d_fgimage.download(h_fgimage);
d_bgimage.download(h_bgimage);
imshow("image", frame);
imshow("foreground mask", h_fgmask);
imshow("foreground image", h_fgimage);
imshow("mean background image", h_bgimage);
if (waitKey(1) == 'q')
break;
}
return 0;
}
与 Jetson TX1 接口的摄像头使用 GStreamer 管道初始化。createBackgroundSubtractorMOG类用于创建 MoG 实现的对象。创建的对象的apply方法用于从第一帧创建前景掩码。它需要一个输入图像、一个image数组来存储前景掩码，以及学习率作为输入。没有活动的房间图像被初始化为 MoG 的背景。因此，任何发生的活动都将被算法归类为前景。
在while循环的每一帧之后，都会持续更新前景掩码和背景图像。getBackgroundImage函数用于获取当前的背景模型。
前景掩码用于创建前景图像，指示哪些对象当前正在移动。它基本上是逻辑操作，在原始帧和前景掩码之间进行。在每一帧之后，前景掩码、前景图像和建模的背景都会下载到主机内存中，以便在屏幕上显示。
以下截图显示了从视频中提取的两个不同帧的输出：
第一行表示房间内没有任何活动时的背景。当有人将手放在摄像头前时，它将被检测为前景，如第二帧结果所示。同样，如果有人将手机放在摄像头前，它也将被归类为前景，如第三帧所示。以下截图显示了代码在每秒帧数方面的性能：
该技术在每秒大约 60-70 帧的速度下工作，可以很容易地用于实时决策。尽管本节中的演示非常简单，但这种应用可以在许多实际情况下使用。房间内的活动可以用来控制房间内现有的设备。这有助于在没有人的情况下节省电力。此应用还可以用于 ATM 机内活动的监控。它还可以用于公共场所的其他视频监控应用。Python 也可以用作 Jetson TX1 上的编程语言，这将在下一节中解释。
在 Jetson TX1 上使用 Python 和 OpenCV 进行计算机视觉
到目前为止，我们使用 C/C++、OpenCV 和 CUDA 开发了所有计算机视觉应用。Jetson TX1 也支持 Python 编程语言用于计算机视觉应用。当在 Jetson TX1 上编译 OpenCV 时，它也会安装 OpenCV 的 Python 二进制文件。因此，熟悉 Python 编程语言的程序员可以使用 Python 接口开发 OpenCV 计算机视觉应用，并在 Jetson TX1 上部署它们。Python 也像所有 Linux 操作系统一样预安装在 Jetson TX1 上。Windows 用户可以单独安装 Python。Python 的安装过程和优势将在下一章中解释。
使用 Python 的一个缺点是，OpenCV Python 接口尚未从 CUDA 加速中受益很大。尽管如此，Python 学习的简便性和其广泛的应用范围已经鼓励了许多软件开发者使用 Python 进行计算机视觉应用。使用 Python 和 OpenCV 读取和显示图像的示例代码如下：
import numpy as np
import cv2
img = cv2.imread('images/cameraman.tif',0)
cv2.imshow("Image read in Python", img)
k = cv2.waitKey(0) & 0xFF
if k == 27: # wait for ESC key to exit
cv2.destroyAllWindows()
在 Python 中，import命令用于在文件中包含一个库。因此，使用import cv2命令包含cv2库。图像以numpy数组的形式存储，所以numpy也被导入到文件中。imread函数用于以与 C++相同的方式读取图像。所有 OpenCV 函数在 Python 中都必须以cv2.为前缀。imshow函数用于显示图像。在 Python 中，所有 OpenCV 函数都具有与 C++类似的签名和功能。
可以使用以下命令在终端中执行代码：
# For Python2.7
$ python image_read.py
# For Python 3
$ python image_read.py
程序的输出如下所示：
这一节只是为了让您知道 Python 也可以用作编程语言，通过 OpenCV 开发计算机视觉应用，并在 Jetson TX1 上部署它。
摘要
本章描述了在部署 CUDA 和 OpenCV 代码时使用 Jetson TX1。本章详细解释了 TX1 板上 GPU 设备的特性，使其非常适合部署计算复杂的应用。本章测量并比较了 Jetson TX1 在执行如添加两个大型数组这样的 CUDA 应用时的性能，并与书中之前提到的笔记本电脑上的 GPU 进行了比较。本章详细解释了在 Jetson TX1 上处理图像的流程。图像处理应用，如图像相加、图像阈值和图像滤波，在 Jetson TX1 上部署，并对它们的性能进行了测量。
Jetson TX1 的最佳之处在于，可以在嵌入式环境中与多个摄像头进行接口连接，并且可以从该摄像头处理视频以设计复杂的计算机视觉应用。从 Jetson TX1 上连接的板载或 USB 摄像头捕获视频的流程在本文中详细解释。
本章还描述了在 Jetson TX1 上部署高级计算机视觉应用，如人脸检测、眼睛检测和背景减法。Python 语言也可以用于在 Jetson TX1 上部署计算机视觉应用。这一概念在章节的最后部分进行了解释。到目前为止，我们已经看到了如何利用 C/C++ 语言的优势来利用 CUDA 和 GPU 加速。
接下来的几章将演示使用 PyCUDA 模块在 Python 语言中使用 CUDA 和 GPU 加速。
问题
-
将 Jetson TX1 上的 GPU 设备性能与书中之前提到的 GeForce 940 GPU 进行比较。
-
判断正误：书中之前提到的所有 CUDA 程序都可以在 Jetson TX1 上执行，无需修改。
-
在 Jetson TX1 上重新编译 OpenCV 的需要是什么？
-
判断正误：OpenCV 无法从连接到 USB 端口的摄像头捕获视频。
-
判断正误：对于计算密集型应用，使用 CSI 摄像头比使用 USB 摄像头更好。
-
如果你正在使用 OpenCV 开发计算密集型的计算机视觉应用，你更倾向于哪种语言以获得更快的性能？
-
在 Jetson TX1 上是否有必要安装单独的 OpenCV Python 绑定或 Python 解释器？
第十章：PyCUDA 入门
我们已经看到如何使用 OpenCV 和 CUDA 加速各种应用。我们使用了 C 或 C++ 作为编程语言。如今，Python 在许多领域都非常流行，因此如果我们能够使用 CUDA 加速 Python 应用程序，将会非常有用。Python 提供了一个 PyCUDA 模块，它正是为此而设计的。
它使用 Nvidia CUDA 工具包，这反过来又需要在计算机上安装 Nvidia 显卡。本章将介绍 Python 语言和 PyCUDA 模块，特别是。它将讨论在 Windows 和 Linux 操作系统上安装 PyCUDA 模块的安装过程。尽管本章需要一些对 Python 语言的熟悉，但新来者也能跟随大多数步骤。
本章将涵盖以下主题：
-
Python 编程语言简介
-
PyCUDA 模块简介
-
在 Windows 上安装 PyCUDA
-
在 Ubuntu 上安装 PyCUDA
技术要求
本章需要良好的 Python 编程语言理解。它还需要任何配备 Nvidia GPU 的计算机或笔记本电脑。本章中使用的 Windows PyCUDA 安装文件可以从以下 GitHub 链接下载：github.com/PacktPublishing/Hands-On-GPU-Accelerated-Computer-Vision-with-OpenCV-and-CUDA。
Python 编程语言简介
Python 正在持续增加其受欢迎程度，因为它可以在许多领域使用，具有广泛的应用。它是一种高级编程语言，可以帮助用几行代码表达复杂的系统。Python 语法易于学习，比其他语言如 C++ 和 Java 更易于阅读，这使得它对新手程序员来说更容易学习。
Python 是一种轻量级的脚本语言，可以轻松用于嵌入式应用。此外，它是一种解释型语言，需要解释器而不是编译器，这与其他编程语言不同。这允许程序员逐行执行代码。它需要一个可以在所有操作系统上轻松安装的 Python 解释器。由于 Python 是开源的，因此一个庞大的社区选择与之合作。他们开发了一系列库，并将其开源，因此它可以无成本地用于应用程序。
Python 可以用于各种领域，如数据科学、机器学习、深度学习、数据分析、图像处理、计算机视觉、数据挖掘和网页开发。它几乎为所有提到的操作系统领域提供了现成的模块，有助于快速开发应用程序。本书前面解释过的 OpenCV 库也具有 Python 接口。因此，它可以轻松地与 Python 代码集成，用于计算机视觉应用程序。Python 还具有用于机器学习和深度学习的库，可以与 OpenCV 一起用于计算机视觉应用程序。
类似于 Python 这样的解释语言的一个缺点是，它比编译语言如 C 或 C++ 慢得多。Python 有一个特性，可以在 Python 脚本中集成 C 或 C++ 代码。这允许你使用 Python 包装器编写计算密集型的 C 或 C++ 代码。
PyCUDA 模块简介
在最后一节中，我们看到了使用 Python 编程语言的优势。还提到 Python 比 C 或 C++ 慢得多。因此，如果它能利用 GPU 的并行处理能力，将会很有益。Python 提供了一个 PyCUDA 包装器，可以通过使用 Nvidia CUDA API 来利用 GPU 的并行计算能力。Python 还有一个 PyOpenCL 模块，可以用于任何 GPU 上的并行计算。
然后，你可能会有一个疑问，为什么你必须使用 PyCUDA，它是专门针对 Nvidia GPU 的。使用 PyCUDA 相比其他类似模块有许多优势；以下是一些原因：
-
它为 Python 开发者提供了一个与 CUDA API 交互的简单接口，并且有良好的文档，这使得学习变得容易。
-
可以使用 PyCUDA 模块在 Python 代码中利用 Nvidia 提供的 CUDA API 的全部功能。
-
PyCUDA 的基础层是用 C++ 编写的，这使得它更快。
-
它具有更高的抽象级别，与基于 Nvidia C 的运行时 API 相比，使用起来更简单。
-
它具有非常高效的内存管理机制，与对象的生存期相关联的对象清理功能。这个特性帮助它编写正确的代码，没有内存泄漏或崩溃。
-
CUDA 代码中的错误也可以通过 Python 异常来处理，这有助于代码中的错误处理机制。
本节描述了使用 PyCUDA 加速 Python 应用程序的优势。在下一节中，我们将看到如何在 Windows 和 Ubuntu 操作系统上安装 PyCUDA 的步骤。
在 Windows 上安装 PyCUDA
本节将描述在 Windows 操作系统上安装 PyCUDA 的步骤。演示使用的是 Windows 10，但该过程适用于任何较新的 Windows 版本。以下是步骤描述：
- 如果您尚未安装 CUDA 工具包，如第一章所述，请从
developer.nvidia.com/cuda-downloads下载最新的 CUDA 工具包。它将询问您的操作系统、CPU 架构以及是否通过互联网安装或先下载整个安装程序。如下面的截图所示，我们选择了带有本地安装程序的 Windows 10。您可以根据您的设置选择相应的值：
-
双击下载的安装程序并按照屏幕提示安装 CUDA 工具包。
-
安装具有可视化 C++版本的最新 Visual Studio 版本。我们使用的是免费的 Visual Studio 2017 社区版，其可视化 C++路径应添加到路径环境变量中。可以通过右键单击我的电脑（此电脑）|属性|高级系统设置|环境变量|系统变量来访问环境变量。在路径环境变量中添加可视化 C++安装和 CUDA 工具包安装的 bin 文件夹路径，如下面的截图所示：
- 将 Anaconda 分布用作 Python 解释器，因此可以从以下网站下载：
www.anaconda.com/download/。我们使用的是 Python 3.6 版本的 Anaconda 5.2，如下面的截图所示：
-
双击下载的安装程序并按照屏幕提示安装 Anaconda。确保勾选将安装路径添加到路径环境变量的复选框。
-
根据您的系统设置，从以下链接下载最新的 PyCUDA 二进制文件：
www.lfd.uci.edu/~gohlke/Pythonlibs/#pycuda。我们使用 CUDA 9.2148 和 Python 3.6，因此选择了相应的 PyCUDA 版本，如下所示：
- 打开命令提示符，转到 PyCUDA 二进制文件下载的文件夹，并执行以下截图所示的命令：
命令将在 Python 分布中完成 PyCUDA 的安装。
检查 PyCUDA 安装步骤
以下步骤用于检查 PyCUDA 是否已正确安装：
-
打开 Spyder，这是一个随 Anaconda 安装一起提供的 Python IDE。您可以在开始菜单中输入
Spyder来打开它。 -
在 Spyder IDE 中，如以下截图所示，在 IPython 控制台中输入
import pycuda。如果没有错误报告，则表示 PyCUDA 已正确安装。
在 Ubuntu 上安装 PyCUDA
本节将描述在 Linux 操作系统上安装 PyCUDA 的步骤。以 Ubuntu 为演示，但此过程适用于任何最新的 Linux 发行版。步骤如下：
- 如果您尚未安装 CUDA 工具包，如第一章所述，请从
developer.nvidia.com/cuda-downloads下载最新的 CUDA 工具包。它将询问您的操作系统、CPU 架构以及是否通过互联网安装或首先下载整个安装程序。如下面的截图所示，我们选择了带有运行文件（本地）安装程序的 Ubuntu。您可以根据您的设置选择值：
-
在命令提示符中运行 sudo sh cuda_9.2.148_396.37_linux.run 命令以安装 CUDA 工具包。
-
将使用 Anaconda 发行版作为 Python 解释器，因此可以从网站：
www.anaconda.com/download/下载和安装。我们使用的是带有 Python 3.6 版本的 Anaconda 5.2，如下面的截图所示：
- 安装 Anaconda 后，在终端中执行以下命令，如下面的截图所示以安装 PyCUDA：
命令将在 Python 发行版中完成 PyCUDA 的安装。
检查 PyCUDA 安装的步骤
以下步骤用于检查 PyCUDA 是否已正确安装：
-
打开 Spyder，这是 Anaconda 安装附带的一个 Python IDE。您可以在终端中输入
Spyder来打开它。 -
在 Spyder IDE 中，在 IPython 控制台中输入
import pycuda，如下面的截图所示。如果没有错误报告，则表示 PyCUDA 已正确安装。
摘要
总结来说，本章介绍了 Python 编程语言及其在各种领域用于广泛应用的用法。与 C 或 C++ 语言相比，它轻量但较慢。因此，如果它能利用 GPU 的并行计算能力，那么它将非常有用。PyCUDA 是一个 Python 包装器，允许 Python 代码利用 Nvidia CUDA API。PyCUDA 相比于 Python 中可用的其他并行处理模块的优势被详细解释。PyCUDA 使用 Nvidia CUDA 运行时 API 和 Python 解释器。Anaconda 是一个著名的 Python 发行版，它附带了许多有用的 Python 库和 IDE，以及 CUDA 工具包。本章讨论了在 Windows 和 Ubuntu 操作系统上安装 PyCUDA 的详细步骤。
在接下来的两章中，我们将详细介绍如何使用 PyCUDA 加速 Python 应用程序。
问题
-
Python 相比于 C 或 C++ 等编程语言有哪些优势？
-
编译型语言和解释型语言之间的区别是什么？
-
判断对错：Python 比 C 或 C++ 更快。
-
PyOpenCL 相比 PyCUDA 的优势是什么？
-
判断对错：Python 允许在 Python 脚本中使用 C 或 C++ 代码。
第十一章：使用 PyCUDA 进行工作
在上一章中，我们看到了为 Windows 和 Linux 操作系统安装 PyCUDA 的步骤。在本章中，我们将首先开发第一个在控制台上显示字符串的 PyCUDA 程序。了解和访问 PyCUDA 运行的 GPU 的设备属性非常重要；这一方法将在本章中详细讨论。我们还将查看 PyCUDA 中内核的线程和块执行。任何 CUDA 编程的重要编程概念，如分配和释放设备上的内存、从主机到设备以及相反的数据传输，以及内核调用，将使用向量加法程序示例进行详细讨论。还将讨论使用 CUDA 事件测量 PyCUDA 程序性能的方法，并将其与 CPU 程序进行比较。这些编程概念将被用于开发一些复杂的 PyCUDA 程序，例如数组元素的平方和矩阵乘法。本章的最后部分描述了在 PyCUDA 中定义内核函数的一些高级方法。
本章将涵盖以下主题：
-
在 PyCUDA 中编写第一个“Hello, PyCUDA!”程序
-
从 PyCUDA 程序访问设备属性
-
PyCUDA 中的线程和块执行
-
使用向量加法程序的基本 PyCUDA 编程概念
-
使用 CUDA 事件测量 PyCUDA 程序的性能
-
PyCUDA 中的一些复杂程序
-
PyCUDA 中的高级内核函数
技术要求
本章需要良好的 Python 编程语言理解。它还需要任何带有 Nvidia GPU 的计算机或笔记本电脑。本章中使用的所有代码都可以从以下 GitHub 链接下载：github.com/PacktPublishing/Hands-On-GPU-Accelerated-Computer-Vision-with-OpenCV-and-CUDA。
查看以下视频以查看代码的实际运行情况：
在 PyCUDA 中编写第一个程序
本节描述了使用 PyCUDA 编写简单“Hello, PyCUDA!”程序的步骤。它将演示编写任何 PyCUDA 程序的工作流程。由于 Python 是一种解释型语言，代码也可以从 Python 终端逐行运行，或者保存为.py扩展名并作为文件执行。
使用 PyCUDA 显示从内核中简单字符串的程序如下所示：
import pycuda.driver as drv
import pycuda.autoinit
from pycuda.compiler import SourceModule
mod = SourceModule("""
#include <stdio.h>
__global__ void myfirst_kernel()
{
printf("Hello,PyCUDA!!!");
}
""")
function = mod.get_function("myfirst_kernel")
function(block=(1,1,1))
在开发 PyCUDA 代码时的第一步是包含代码所需的所有库。使用import指令来包含一个库、模块、类或函数。这与在 C 或 C++中包含指令类似，并且可以通过以下三种不同的方式完成，如下面的步骤所示。以下也展示了使用三个导入模块的示例：
-
导入
pycuda.driver为drv这表示导入了 pymodule 的驱动子模块，并给它一个简写符号
drv，所以当需要使用pycuda.driver模块中的函数时，可以使用drv.functionname。此模块包含内存管理函数、设备属性、数据方向函数等。 -
导入
pycuda.autoinit这个命令表示从
pycuda的autoint模块导入了。没有给出任何缩写符号。autoint模块用于设备初始化、上下文创建和内存清理。此模块不是必需的，上述所有功能也可以手动完成。 -
从
pycuda.compiler导入SourceModule这个命令表示只从
pycuda.compiler模块导入了SourceModule类。当你只想使用一个大模块中的一个类时，这是很重要的。SourceModule类用于在 PyCUDA 中定义类似 C 的内核函数。
C 或 C++内核代码作为构造函数传递给Sourcemodule类，并创建 mod 对象。内核代码非常简单，因为它只是在控制台上打印一个Hello, PyCUDA!字符串。由于内核代码中使用了printf函数，因此包含stdio.h头文件非常重要。myfirst_kernel函数在内核代码中使用__global__指令定义，表示该函数将在 GPU 上执行。该函数不接受任何参数。它只是在控制台上打印一个字符串。此内核函数将由nvcc编译器编译。
这个函数可以通过使用mod对象的get_function方法创建一个指向该函数的指针在 Python 代码内部使用。内核函数的名称作为引号内的参数给出。指针变量可以取任何名字。这个指针变量用于在代码的最后一行调用内核。内核函数的参数也可以在这里指定，但由于myfirst_kernel函数没有参数，所以没有指定参数。内核可以提供的线程数和每个网格要启动的块数也可以通过使用可选的块和网格参数作为参数提供。块参数被赋予值（1,1,1），这是一个 1 x 3 的 Python 元组，表示块大小为 1 x 1 x 1。因此，将启动一个线程，该线程将在控制台上打印字符串。
程序的输出如下所示：
总结来说，本节逐步展示了开发一个简单的 PyCUDA 程序的步骤。
内核调用
使用 ANSI C 关键字以及 CUDA 扩展关键字编写的设备代码被称为 内核。它通过一种名为 内核调用 的方法从 Python 代码中启动。基本上，内核调用的意义是我们从主机代码中启动设备代码。内核代码与正常 C 函数非常相似；只是这段代码是由多个线程并行执行的。它在 Python 中的语法非常简单，如下所示：
kernel (parameters for kernel,block=(tx,ty,tz) , grid=(bx,by,bz))
这从我们要启动的内核函数的指针开始。你应该确保这个内核指针是通过 get_function 方法创建的。然后，它可以包括用逗号分隔的内核函数参数。块参数表示要启动的线程数，而网格参数表示网格中的块数。块和网格参数使用一个 1 x 3 的 Python 元组指定，表示三维空间中的块和线程。内核启动启动的线程总数将是这些数字的乘积。
从 PyCUDA 程序访问 GPU 设备属性
PyCUDA 提供了一个简单的 API 来查找信息，例如，哪些 CUDA 兼容的 GPU 设备（如果有）存在，以及每个设备支持哪些功能。在编写 PyCUDA 程序之前，了解正在使用的 GPU 设备的属性非常重要，这样就可以使用设备的最佳资源。
使用 PyCUDA 显示系统上所有 CUDA 兼容设备属性的程序如下所示：
import pycuda.driver as drv
import pycuda.autoinit
drv.init()
print("%d device(s) found." % drv.Device.count())
for i in range(drv.Device.count()):
dev = drv.Device(i)
print("Device #%d: %s" % (i, dev.name()))
print(" Compute Capability: %d.%d" % dev.compute_capability())
print(" Total Memory: %s GB" % (dev.total_memory()//(1024*1024*1024)))
attributes = [(str(prop), value)
for prop, value in list(dev.get_attributes().items())]
attributes.sort()
n=0
for prop, value in attributes:
print(" %s: %s " % (prop, value),end=" ")
n = n+1
if(n%2 == 0):
print(" ")
首先，重要的是要获取系统上存在的 CUDA 兼容设备数量，因为一个系统可能包含多个启用 GPU 的设备。这个数量可以通过 PyCUDA 中驱动类 drv.Device.count() 函数确定。系统上所有设备都会被迭代以确定每个设备的属性。使用 drv.Device 函数为每个设备创建一个指针对象。这个指针用于确定特定设备的所有属性。
name 函数将给出特定设备的名称，而 total_memory 将给出设备上可用的 GPU 全局内存的大小。其他属性存储为 Python 字典，可以通过 get_attributes().items() 函数检索。这通过 Python 中的列表推导式转换为元组列表。这个列表的所有行都包含一个 2 x 1 的元组，其中包含属性的名称和其值。
使用 for 循环迭代此列表以在控制台上显示所有属性。此程序在配备 GeForce 940 GPU 和 CUDA 9 的笔记本电脑上执行。程序输出如下：
这些属性在本书的早期章节中已经详细讨论过，所以我们不再重复讨论；然而，为了总结，本节展示了从 PyCUDA 程序中访问 GPU 设备属性的方法。
PyCUDA 中的线程和块执行
我们在 内核调用 部分看到，我们可以并行启动多个块和多个线程。那么，这些块和线程以什么顺序开始和结束它们的执行？如果我们想在其他线程中使用一个线程的输出，了解这一点很重要。为了理解这一点，我们修改了前面章节中看到的 hello,PyCUDA! 程序中的内核，通过在内核调用中包含一个打印语句来打印块号。修改后的代码如下所示：
import pycuda.driver as drv
import pycuda.autoinit
from pycuda.compiler import SourceModule
mod = SourceModule("""
#include <stdio.h>
__global__ void myfirst_kernel()
{
printf("I am in block no: %d \\n", blockIdx.x);
}
""")
function = mod.get_function("myfirst_kernel")
function(grid=(4,1),block=(1,1,1))
从代码中可以看出，我们以并行方式启动了 10 个块，每个块有一个线程。在内核代码中，我们正在打印内核执行的块 ID。我们可以将其视为 10 个相同的 myfirstkernel 并行开始执行。这些副本中的每一个都将有一个唯一的块 ID，可以通过 blockIdx.x 指令访问，以及唯一的线程 ID，可以通过 threadIdx.x 访问。这些 ID 将告诉我们哪个块和线程正在执行内核。当你多次运行程序时，你会发现每次块执行的顺序都不同。一个示例输出如下所示：
它可以产生 n 的阶乘数量的不同输出，其中 n 表示并行启动的块的数量。因此，每次你在 PyCUDA 中编写程序时，都应该小心，因为块会以随机顺序执行。
PyCUDA 中的基本编程概念
在本节中，我们将开始使用 PyCUDA 开发一些有用的功能。本节还将通过一个简单的加法示例展示 PyCUDA 的一些有用的函数和指令。
PyCUDA 中的加法
Python 提供了一个用于数值运算的非常快速的库，称为 numpy (Numeric Python)。它是用 C 或 C++ 开发的，并且对于 Python 中的数组操作非常有用。它在 PyCUDA 程序中经常被用作 PyCUDA 内核函数的参数，这些参数作为 numpy 数组传递。本节解释了如何使用 PyCUDA 添加两个数字。添加两个数字的基本内核代码如下所示：
import pycuda.autoinit
import pycuda.driver as drv
import numpy
from pycuda.compiler import SourceModule
mod = SourceModule("""
__global__ void add_num(float *d_result, float *d_a, float *d_b)
{
const int i = threadIdx.x;
d_result[i] = d_a[i] + d_b[i];
}
""")
如前所述，导入 SourceModule 类和驱动类。numpy 库也被导入，因为它将用于将参数传递给内核代码。add_num 内核函数被定义为 SourceModule 类的构造函数。该函数接受两个设备指针作为输入，一个设备指针指向加法的结果作为输出。需要注意的是，尽管我们在添加两个数字，但内核函数被定义为可以同时处理两个数组加法。两个单个数字不过是每个只有一个元素的数组。如果没有错误，此代码将被编译并加载到设备上。从 Python 调用此内核代码的代码如下所示：
add_num = mod.get_function("add_num")
h_a = numpy.random.randn(1).astype(numpy.float32)
h_b = numpy.random.randn(1).astype(numpy.float32)
h_result = numpy.zeros_like(h_a)
d_a = drv.mem_alloc(h_a.nbytes)
d_b = drv.mem_alloc(h_b.nbytes)
d_result = drv.mem_alloc(h_result.nbytes)
drv.memcpy_htod(d_a,h_a)
drv.memcpy_htod(d_b,h_b)
add_num(
d_result, d_a, d_b,
block=(1,1,1), grid=(1,1))
drv.memcpy_dtoh(h_result,d_result)
print("Addition on GPU:")
print(h_a[0],"+", h_b[0] , "=" , h_result[0])
使用 get_function 创建内核函数的指针引用。使用 numpy.random.randn(1) 函数创建两个随机数，该函数用于创建正态分布中的随机数。这些数字使用 astype(numpy.float32) 方法转换为单精度浮点数。用于在主机上存储结果的 numpy 数组被初始化为零。
可以使用 PyCUDA 中驱动类的 mem_alloc 函数在设备上分配内存。内存的大小作为函数的参数传递。使用 h_a.nbytes 函数找到输入的大小。PyCUDA 在驱动类中提供了一个 memcpy 函数，用于从主机内存到设备内存以及相反方向的复制数据。
drv.memcpy_htod 函数用于将数据从主机内存复制到设备内存。设备内存的指针作为第一个参数传递，主机内存指针作为第二个参数传递。通过传递设备指针以及指定要启动的块和线程数量的数字作为参数，调用 add_num 内核。在前面给出的代码中，使用一个线程启动了一个块。内核计算的结果通过使用 drv.memcpy_dtoh 函数复制回主机。结果在控制台上显示，如下所示：
总结来说，本节展示了 PyCUDA 程序的结构。它从内核定义代码开始。然后在 Python 中定义输入。在设备上分配内存并将输入传输到设备内存。接着是内核调用，它将计算结果。然后将结果传输到主机进行进一步处理。PyCUDA 提供了更简单的 API 来执行此操作，这将在下一节中解释。
使用驱动类简化加法程序
PyCUDA 提供了一个更简单的内核调用 API，它不需要内存分配和内存复制。这是通过 API 隐式完成的。这可以通过使用 PyCUDA 中驱动类中的 In 和 Out 函数来实现。修改后的数组加法代码如下所示：
import pycuda.autoinit
import pycuda.driver as drv
import numpy
N = 10
from pycuda.compiler import SourceModule
mod = SourceModule("""
__global__ void add_num(float *d_result, float *d_a, float *d_b)
{
const int i = threadIdx.x;
d_result[i] = d_a[i] + d_b[i];
}
""")
add_num = mod.get_function("add_num")
h_a = numpy.random.randn(N).astype(numpy.float32)
h_b = numpy.random.randn(N).astype(numpy.float32)
h_result = numpy.zeros_like(h_a)
add_num(
drv.Out(h_result), drv.In(h_a), drv.In(h_b),
block=(N,1,1), grid=(1,1))
print("Addition on GPU:")
for i in range(0,N):
print(h_a[i],"+", h_b[i] , "=" , h_result[i])
在前面的代码中，数组中的十个元素被添加，而不是单个元素。内核函数与之前看到的代码完全相同。在主机上创建了两个包含十个随机数的数组。现在，不是创建它们的内存并将它们传输到设备，而是直接调用内核。通过指定数据方向使用 drv.Out 或 drv.In 来修改内核调用。这简化了 PyCUDA 代码并减少了代码的大小。
内核函数使用一个块和每个块 N 个线程。这 N 个线程并行地添加数组中的 N 个元素，从而加速了加法操作。使用 drv.out 指令，内核的结果会自动下载到主机内存，因此这个结果可以直接通过 for 循环打印到控制台。使用 PyCUDA 进行十个元素加法的结果如下所示：
总结来说，本节通过一个简单的数组加法程序介绍了 PyCUDA 的重要概念和函数。使用 PyCUDA 的性能提升可以通过下一节中解释的 CUDA 事件来量化。
使用 CUDA 事件测量 PyCUDA 程序的性能
到目前为止，我们还没有明确确定 PyCUDA 程序的性能。在本节中，我们将看到如何使用 CUDA 事件测量程序的性能。这在 PyCUDA 中是一个非常重要的概念，因为它将允许你从许多选项中选择特定应用的性能最佳算法。
CUDA 事件
我们可以使用 Python 时间测量选项来测量 CUDA 程序的性能，但它不会给出准确的结果。它将包括许多其他因素中的线程延迟在操作系统中的时间开销和调度。使用 CPU 测量的时间也将取决于高精度 CPU 计时器的可用性。很多时候，当 GPU 内核运行时，主机正在执行异步计算，因此 Python 的 CPU 计时器可能不会给出内核执行的正确时间。因此，为了测量 GPU 内核计算的时间，PyCUDA 提供了一个事件 API。
CUDA 事件是在 PyCUDA 程序中指定点记录的 GPU 时间戳。在这个 API 中，GPU 记录时间戳，消除了使用 CPU 计时器测量性能时存在的问题。使用 CUDA 事件测量时间有两个步骤：创建事件和记录事件。我们可以记录两个事件，一个在代码的开始处，一个在结束处。然后我们将尝试计算两个事件之间的时间差，这将给出代码的整体性能。
在 PyCUDA 代码中，可以包含以下行来使用 CUDA 事件 API 测量性能：
import pycuda.driver as drv
start = drv.Event()
end=drv.Event()
#Start Time
start.record()
#The kernel code for which time is to be measured
#End Time
end.record()
end.synchronize()
#Measure time difference
secs = start.time_till(end)*1e-3
使用record方法来测量当前时间戳。在内核代码前后测量时间戳以测量内核执行时间。可以使用time_till方法测量时间戳之间的差异，如前述代码所示。它将以毫秒为单位给出时间，然后将其转换为秒。在下一节中，我们将尝试使用 CUDA 事件来测量代码的性能。
使用大数组加法来测量 PyCUDA 的性能
本节将演示如何使用 CUDA 事件来测量 PyCUDA 程序的性能。同时，还描述了 PyCUDA 代码与简单 Python 代码性能的比较。为了准确比较性能，选取了包含一百万个元素的数组。下面展示了用于大数组加法的内核代码：
import pycuda.autoinit
import pycuda.driver as drv
import numpy
import time
import math
from pycuda.compiler import SourceModule
N = 1000000
mod = SourceModule("""
__global__ void add_num(float *d_result, float *d_a, float *d_b,int N)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
while (tid < N)
{
d_result[tid] = d_a[tid] + d_b[tid];
tid = tid + blockDim.x * gridDim.x;
}
}
""")
由于元素数量较多，因此会启动多个块和线程。所以，既使用线程 ID 又使用块 ID 来计算线程索引。如果启动的总线程数不等于元素数，则同一线程会添加多个元素。这是通过内核函数内部的while循环实现的。同时，它也会确保线程索引不会超出数组元素。除了输入数组和输出数组外，数组的大小也被作为内核函数的参数，因为在SourceModule中内核代码无法访问 Python 的全局变量。下面展示了用于添加大数组的 Python 代码：
start = drv.Event()end=drv.Event()
add_num = mod.get_function("add_num")
h_a = numpy.random.randn(N).astype(numpy.float32)
h_b = numpy.random.randn(N).astype(numpy.float32)
h_result = numpy.zeros_like(h_a)
h_result1 = numpy.zeros_like(h_a)
n_blocks = math.ceil((N/1024))
start.record()
add_num(
drv.Out(h_result), drv.In(h_a), drv.In(h_b),numpy.uint32(N),
block=(1024,1,1), grid=(n_blocks,1))
end.record()
end.synchronize()
secs = start.time_till(end)*1e-3
print("Addition of %d element of GPU"%N)
print("%fs" % (secs))
创建了两个事件start和stop来测量 GPU 代码的执行时间。使用驱动类中的Event()函数来定义事件对象。然后，使用get_function创建内核函数的指针引用。使用numpy库的randn函数初始化两个包含一百万个元素的数组，并使用随机数。由于它们是浮点数，因此将它们转换为单精度数以加快设备上的计算速度。
每个块支持 1,024 个线程，正如我们在设备属性部分所看到的。因此，根据这一点，通过将N除以 1,024 来计算总块数。它可能是一个浮点值，因此使用numpy库的ceil函数将其转换为下一个最高整数值。然后，使用计算出的块数和每个块的 1,024 个线程来启动内核。使用numpy.uint32数据类型传递数组的大小。
使用记录函数在调用内核函数前后记录时间，并计算时间差以测量内核函数的执行时间。计算出的时间将打印在控制台上。为了将此性能与 CPU 时间进行比较，程序中添加了以下代码：
start = time.time()
for i in range(0,N):
h_result1[i] = h_a[i] +h_b[i]
end = time.time()
print("Addition of %d element of CPU"%N)
print(end-start,"s")
使用 Python 的时间库来测量 CPU 时间。使用for循环遍历数组中的每个元素。（注意：由于两个数组都是 numpy 数组，你也可以使用h_result1 = h_a + h_b。）使用time.time()函数在for循环前后测量时间，并将这两个时间之间的差值打印到控制台。程序的输出如下所示：
如输出所示，GPU 添加一百万个元素需要 9.4 毫秒，而 CPU 需要 415.15 毫秒，因此使用 GPU 可以实现大约 50 倍的性能提升。
总结来说，本节展示了使用事件来测量 GPU 代码的计时。将 GPU 性能与 CPU 性能进行比较，以量化使用 GPU 时的性能提升。
PyCUDA 中的复杂程序
到现在为止，PyCUDA 的语法和术语应该已经熟悉了。我们将利用这些知识来开发高级程序，并学习一些 PyCUDA 的高级概念。在本节中，我们将使用 PyCUDA 开发一个程序，使用三种不同的方法对数组的元素进行平方。我们还将学习在 PyCUDA 中进行矩阵乘法的代码。
PyCUDA 中矩阵的逐元素平方
在本节中，使用三种不同的方法执行矩阵中数字的逐元素平方操作。在这个过程中，详细解释了使用多维线程和块的概念、驱动类中的inout指令以及gpuarray类。
使用多维线程的简单内核调用
本节实现了使用 PyCUDA 对矩阵的每个元素进行平方的简单内核函数。以下是一个 5 x 5 矩阵中每个元素平方的内核函数示例：
import pycuda.driver as drv
import pycuda.autoinit
from pycuda.compiler import SourceModule
import numpy
mod = SourceModule("""
__global__ void square(float *d_a)
{
int idx = threadIdx.x + threadIdx.y*5;
d_a[idx] = d_a[idx]*d_a[idx];
}
""")
核函数square只接受一个设备指针作为输入，该指针指向矩阵，并将每个元素替换为其平方。随着多维线程的启动，x和y方向上的线程索引被用来索引矩阵中的值。你可以假设一个 5 x 5 矩阵被展平为一个 1 x 25 的向量，以理解索引机制。请注意，在这个代码中，矩阵的大小是硬编码为5的，但它也可以像上一节中的数组大小一样由用户定义。使用此内核函数的 Python 代码如下所示：
start = drv.Event()
end=drv.Event()
h_a = numpy.random.randint(1,5,(5, 5))
h_a = h_a.astype(numpy.float32)
h_b=h_a.copy()
start.record()
d_a = drv.mem_alloc(h_a.size * h_a.dtype.itemsize)
drv.memcpy_htod(d_a, h_a)
square = mod.get_function("square")
square(d_a, block=(5, 5, 1), grid=(1, 1), shared=0)
h_result = numpy.empty_like(h_a)
drv.memcpy_dtoh(h_result, d_a)
end.record()
end.synchronize()
secs = start.time_till(end)*1e-3
print("Time of Squaring on GPU without inout")
print("%fs" % (secs))
print("original array:")
print(h_a)
print("Square with kernel:")
print(h_result)
创建了两个事件来测量内核函数的计时。在主机上，一个 5x5 的矩阵通过 numpy.random 模块的 randint 函数初始化为随机数。这需要三个参数。前两个参数定义了用于生成随机数的数字范围。第一个参数是最小值，第二个参数是用于生成数字的最大值。第三个参数是大小，指定为元组 (5,5)。这个生成的矩阵再次转换为单精度数以加快处理速度。矩阵的内存是在设备上分配的，生成的随机数矩阵被复制到其中。
创建了指向内核函数的指针引用，并通过传递设备内存指针作为参数调用内核。内核调用使用多维线程，x 和 y 方向上的值为 5。因此，总共启动了 25 个线程，每个线程计算矩阵中单个元素的平方。内核计算的结果被复制回主机并在控制台上显示。内核所需的时间以及输入和输出矩阵都显示在控制台上。
计算一个 5x5 矩阵中每个元素的平方需要 149 毫秒。使用驱动类的 inout 指令可以简化相同的计算。这将在下一节中解释。
使用 inout 与内核调用结合使用
如上一节程序的核心函数所示，相同的数组既用作输入也用作输出。PyCUDA 的驱动模块为这类情况提供了一个 inout 指令。它消除了为该数组单独分配内存、上传到设备以及将结果下载回主机的需求。所有操作都在内核调用期间同时进行。这使得代码更简单，更容易阅读。使用驱动类 inout 指令的 Python 代码如下所示：
start.record()
start.synchronize()
square(drv.InOut(h_a), block=(5, 5, 1))
end.record()
end.synchronize()
print("Square with InOut:")
print(h_a)
secs = start.time_till(end)*1e-3
print("Time of Squaring on GPU with inout")
print("%fs" % (secs))
使用 inout 指令初始化 CUDA 事件以测量代码的性能。内核调用与上一节相同，因此在此不再重复。可以看出，在调用平方内核时，通过 drv.inout 指令传递了一个变量作为参数。因此，所有与设备相关的操作都在这一步中完成。内核调用使用多维线程，与上一节的情况相同。计算结果和耗时被打印到控制台，如下所示：
所需时间相对于原始内核来说比较少。因此，通过使用驱动类中的 inout 指令，PyCUDA 代码可以变得高效且易于阅读。PyCUDA 还提供了一个用于数组相关操作的 gpuarray 类。它也可以用于平方操作，这将在下一节中解释。
使用 gpuarray 类
Python 提供了一个 numpy 库，用于在 Python 中进行数值计算。PyCUDA 提供了一个与 numpy 类似的 gpuarray 类，该类在 GPU 设备上存储其数据和执行其计算。数组的形状和数据类型与 numpy 中完全相同。gpuarray 类提供了许多用于计算的算术方法。它消除了使用 SourceModule 在 C 或 C++ 中指定内核代码的需要。因此，PyCUDA 代码将只包含 Python 代码。使用 gpuarray 类对矩阵的每个元素进行平方的代码如下所示：
import pycuda.gpuarray as gpuarray
import numpy
import pycuda.driver as drv
start = drv.Event()
end=drv.Event()
start.record()
start.synchronize()
h_b = numpy.random.randint(1,5,(5, 5))
d_b = gpuarray.to_gpu(h_b.astype(numpy.float32))
h_result = (d_b**2).get()
end.record()
end.synchronize()
print("original array:")
print(h_b)
print("doubled with gpuarray:")
print(h_result)
secs = start.time_till(end)*1e-3
print("Time of Squaring on GPU with gpuarray")
print("%fs" % (secs))
gpuarray 类需要在代码中使用，它位于 pycuda.gpuarray 模块中。矩阵使用从 1 到 5 的随机整数进行初始化，以便进行计算。这个矩阵通过使用 gpuarray 类的 to_gpu() 方法上传到设备内存。要上传的矩阵作为参数传递给此方法。矩阵被转换为单精度数字。所有对这个上传矩阵的操作都将在该设备上执行。平方操作以与我们在 Python 代码中执行的方式类似的方式进行，但由于变量是使用 gpuarray 存储在设备上的，因此此操作也将在该设备上执行。结果通过使用 get 方法下载回主机。以下是在控制台上显示的结果，包括使用 gpuarray 进行逐元素平方所需的时间：
计算平方需要大约 58 毫秒。它完全消除了在 C 语言中定义内核函数的需要，其功能与 numpy 库相似，因此 Python 程序员可以轻松地与之一起工作。
总结来说，在本节中，我们使用 PyCUDA 以三种不同的方式开发了一个逐元素平方程序。我们还看到了 PyCUDA 中的多维线程、inout 指令和 gpuarray 类的概念。
使用 GPU 数组进行点积
两个向量之间的点积是各种应用中重要的数学运算。上一节中使用的 gpuarray 类可以用来计算两个向量之间的点积。gpuarray 方法计算点积的性能与 numpy 操作进行了比较。用于使用 numpy 计算点积的代码如下所示：
import pycuda.gpuarray as gpuarray
import pycuda.driver as drv
import numpy
import time
import pycuda.autoinit
n=100
h_a=numpy.float32(numpy.random.randint(1,5,(1,n)))
h_b=numpy.float32(numpy.random.randint(1,5,(1,n)))
start=time.time()
h_result=numpy.sum(h_a*h_b)
#print(numpy.dot(a,b))
end=time.time()-start
print("Answer of Dot Product using numpy")
print(h_result)
print("Time taken for Dot Product using numpy")
print(end,"s")
初始化了两个各有 100 个元素的向量，并使用随机整数来计算点积。Python 的时间模块用于计算计算点积所需的时间。使用*运算符来计算两个向量的逐元素乘积，然后将这些结果相加以计算总的点积。请注意，这里使用的numpy.dot方法用于矩阵乘法，不能用于点积。计算出的点积和时间将在控制台上显示。以下是如何使用gpuarray在 GPU 上执行相同操作的代码：
d_a = gpuarray.to_gpu(h_a)
d_b = gpuarray.to_gpu(h_b)
start1 = drv.Event()
end1=drv.Event()
start1.record()
d_result = gpuarray.dot(d_a,d_b)
end1.record()
end1.synchronize()
secs = start1.time_till(end1)*1e-3
print("Answer of Dot Product on GPU")
print(d_result.get())
print("Time taken for Dot Product on GPU")
print("%fs" % (secs))
if(h_result==d_result.get()):
print("The computed dor product is correct")
使用to_gpu方法将两个向量上传到 GPU 上以计算点积。gpuarray类提供了一个点积方法，可以直接用于计算点积。它需要一个 GPU 数组作为参数。使用get()方法将计算结果下载回主机。计算结果和通过 CUDA 事件测量的时间将在控制台上显示。程序的结果如下所示：
从输出中可以看出，使用numpy和gpuarray计算点积得到相同的结果。numpy库计算点积需要 37 毫秒，而 GPU 只需 0.1 毫秒即可完成相同的操作。这进一步说明了使用 GPU 和 PyCUDA 进行复杂数学运算的优势。
矩阵乘法
经常使用的一个重要数学运算是矩阵乘法。本节将演示如何使用 PyCUDA 在 GPU 上执行矩阵乘法。当矩阵的大小非常大时，这是一个非常复杂的数学运算。应记住，对于矩阵乘法，第一个矩阵的列数应等于第二个矩阵的行数。矩阵乘法不是累积操作。为了避免复杂性，在这个例子中，我们使用相同大小的方阵。如果你熟悉矩阵乘法的数学，你可能还记得，第一个矩阵的每一行将与第二个矩阵的所有列相乘。这将对第一个矩阵的所有行重复进行。以下是一个 3x3 矩阵乘法的示例：
结果矩阵中的每个元素将通过将第一个矩阵的对应行与第二个矩阵的对应列相乘来计算。这个概念被用来开发以下所示的内核函数：
import numpy as np
from pycuda import driver
from pycuda.compiler import SourceModule
import pycuda.autoinit
MATRIX_SIZE = 3
matrix_mul_kernel = """
__global__ void Matrix_Mul_Kernel(float *d_a, float *d_b, float *d_c)
{
int tx = threadIdx.x;
int ty = threadIdx.y;
float value = 0;
for (int i = 0; i < %(MATRIX_SIZE)s; ++i) {
float d_a_element = d_a[ty * %(MATRIX_SIZE)s + i];
float d_b_element = d_b[i * %(MATRIX_SIZE)s + tx];
value += d_a_element * d_b_element;
}
d_c[ty * %(MATRIX_SIZE)s + tx] = value;
} """
matrix_mul = matrix_mul_kernel % {'MATRIX_SIZE': MATRIX_SIZE}
mod = SourceModule(matrix_mul)
内核函数接受两个输入数组和一个输出数组作为参数。矩阵的大小作为常量传递给内核函数。这样就消除了需要将向量的大小作为内核函数参数之一的需求，正如本章前面所解释的那样。两种方法都是正确的，取决于程序员认为哪种更方便。每个线程计算结果矩阵的一个元素。第一矩阵的行和第二矩阵的列的所有元素在for循环内相乘并求和。答案被复制到结果矩阵中的相应位置。内核函数内部计算索引的细节可以在本书的早期章节中找到。以下是如何使用此内核函数的 Python 代码：
h_a = np.random.randint(1,5,(MATRIX_SIZE, MATRIX_SIZE)).astype(np.float32)
h_b = np.random.randint(1,5,(MATRIX_SIZE, MATRIX_SIZE)).astype(np.float32)
d_a = gpuarray.to_gpu(h_a)
d_b = gpuarray.to_gpu(h_b)
d_c_gpu = gpuarray.empty((MATRIX_SIZE, MATRIX_SIZE), np.float32)
matrixmul = mod.get_function("Matrix_Mul_Kernel")
matrixmul(d_a, d_b,d_c_gpu,
block = (MATRIX_SIZE, MATRIX_SIZE, 1),
)
print("*" * 100)
print("Matrix A:")
print(d_a.get())
print("*" * 100)
print("Matrix B:")
print(d_b.get())
print("*" * 100)
print("Matrix C:")
print(d_c_gpu.get())
# compute on the CPU to verify GPU computation
h_c_cpu = np.dot(h_a, h_b)
if h_c_cpu == d_c_gpu.get() :
print("The computed matrix multiplication is correct")
两个大小为 3 x 3 的矩阵被初始化为从1到5的随机整数。这些矩阵使用gpuarray类的to_gpu方法上传到设备内存。创建一个空的 GPU 数组以在设备上存储结果。这三个变量作为参数传递给内核函数。内核函数调用时，矩阵大小作为x和y方向的维度。结果使用get()方法下载回主机。两个输入矩阵和 GPU 计算的结果在控制台上打印。使用numpy库的 dot 方法在 CPU 上计算矩阵乘法。结果与内核计算的结果进行比较，以验证内核计算的结果。程序的结果如下所示：
总结来说，我们已经开发了一个简单的内核函数，使用 PyCUDA 执行矩阵乘法。这个内核函数可以通过使用共享内存进一步优化，正如本书前面所解释的那样。
PyCUDA 中的高级内核函数
到目前为止，我们已经看到了使用SourceModule类在 C 或 C++中定义内核函数的使用。我们还使用了gpuarray类来进行设备计算，而不需要显式定义内核函数。本节描述了 PyCUDA 中可用的高级内核定义功能。这些功能用于开发各种并行通信模式的内核函数，如映射、归约和扫描操作。
PyCUDA 中的元素级内核
这个特性允许程序员定义一个内核函数，该函数作用于数组的每个元素。它允许程序员将一个或多个操作数组成的复杂表达式执行为一个单一的计算步骤。以下是这样定义大型数组元素级加法内核函数的方式：
import pycuda.gpuarray as gpuarray
import pycuda.driver as drv
from pycuda.elementwise import ElementwiseKernel
from pycuda.curandom import rand as curand
add = ElementwiseKernel(
"float *d_a, float *d_b, float *d_c",
"d_c[i] = d_a[i] + d_b[i]",
"add")
使用PyCuda.elementwise.ElementwiseKernel函数来定义元素级核函数。它需要三个参数。第一个参数是核函数的参数列表。第二个参数定义了对每个元素要执行的操作，第三个参数指定了核函数的名称。以下是如何使用此核函数的 Python 代码示例：
n = 1000000
d_a = curand(n)
d_b = curand(n)
d_c = gpuarray.empty_like(d_a)
start = drv.Event()
end=drv.Event()
start.record()
add(d_a, d_b, d_c)
end.record()
end.synchronize()
secs = start.time_till(end)*1e-3
print("Addition of %d element of GPU"%shape)
print("%fs" % (secs))
# check the result
if d_c == (d_a + d_b):
print("The sum computed on GPU is correct")
使用pycuda.curandom类中的curand函数用随机数初始化两个数组。这又是一个有用的功能，因为它消除了在主机上初始化然后上传到设备内存的需要。创建一个空的 GPU 数组来存储结果。通过将这些三个变量作为参数传递来调用add核函数。使用 CUDA 事件计算一百万个元素加法所需的时间，并在控制台上显示。
程序的输出如下所示：
元素级核函数只需要 0.6 毫秒即可完成数组中一百万个元素的加法操作。这种性能优于本章前面看到的程序。因此，当要对向量执行元素级操作时，元素级核定义是一个非常重要的概念需要记住。
归约核
归约操作可以通过使用某些表达式将元素集合归约到单个值来定义。它在各种并行计算应用中非常有用。以计算向量点积的例子来展示 PyCUDA 中的归约概念。以下是如何使用 PyCUDA 中归约核功能计算点积的程序示例：
import pycuda.gpuarray as gpuarray
import pycuda.driver as drv
import numpy
from pycuda.reduction import ReductionKernel
import pycuda.autoinit
n=5
start = drv.Event()
end=drv.Event()
start.record()
d_a = gpuarray.arange(n,dtype= numpy.uint32)
d_b = gpuarray.arange(n,dtype= numpy.uint32)
kernel = ReductionKernel(numpy.uint32,neutral="0",reduce_expr="a+b",map_expr="d_a[i]*d_b[i]",arguments="int *d_a,int *d_b")
d_result = kernel(d_a,d_b).get()
end.record()
end.synchronize()
secs = start.time_till(end)*1e-3
print("Vector A")
print(d_a)
print("Vector B")
print(d_b)
print("The computed dot product using reduction:")
print(d_result)
print("Dot Product on GPU")
print("%fs" % (secs))
PyCUDA 提供了pycuda.reduction.ReductionKernel类来定义归约核。它需要许多参数。第一个参数是输出数据类型。第二个参数是中值，通常定义为0。第三个参数是用于归约元素集合的表达式。在前面代码中定义了加法操作。第四个参数定义为归约前操作数之间映射操作的表达式。在代码中定义了元素级乘法。最后一个参数定义了核函数的参数。
计算点积的归约核函数需要两个向量之间的元素级乘法，然后对所有元素进行加法。使用arange函数定义了两个向量。它在 Python 中的range函数类似，但arange会将数组保存在设备上。通过将这些两个向量作为参数传递来调用核函数，并将结果检索到主机。使用 CUDA 事件计算所需的计算时间，并在控制台上与点积的结果一起显示，如下所示：
减少内核计算点积大约需要 2.5 秒，与上一节中看到的显式内核相比，这是一个相对较长时间。然而，在需要减少操作的并行计算应用中，它非常有用。
扫描内核
扫描操作再次是一个非常重要的并行计算范式。扫描操作符将指定的函数应用于输入的第一个元素。该函数的结果作为输入提供，并带有原始输入的第二个元素。所有中间结果形成输出序列。这个概念可以用于各种应用。以累积加法为例，演示了 PyCUDA 中的扫描内核概念。累积加法不过是将加法应用于向量的每个元素，顺序进行。示例如下：
Input Vector
[7 5 9 2 9]
Scan Operation for cumulative sum
[7,7+5,7+5+9,7+5+9+2,7+2+9+2+7]
如所示，前一次加法的结果被添加到当前元素中，以计算当前位置的输出。这被称为包含扫描操作。如果输入的当前元素不参与，则称为排除扫描。使用包含扫描执行累积求和的程序如下所示：
import pycuda.gpuarray as gpuarray
import pycuda.driver as drv
import numpy
from pycuda.scan import InclusiveScanKernel
import pycuda.autoinit
n=10
start = drv.Event()
end=drv.Event()
start.record()
kernel = InclusiveScanKernel(numpy.uint32,"a+b")
h_a = numpy.random.randint(1,10,n).astype(numpy.int32)
d_a = gpuarray.to_gpu(h_a)
kernel(d_a)
end.record()
end.synchronize()
secs = start.time_till(end)*1e-3
assert(d_a.get() == numpy.cumsum(h_a,axis=0)).all()
print("The input data:")
print(h_a)
print("The computed cumulative sum using Scan:")
print(d_a.get())
print("Cumulative Sum on GPU")
print("%fs" % (secs))
PyCUDA 提供了pycuda.scan.InclusiveScanKernel类来定义一个包含扫描内核。它需要输出数据类型和用于扫描的操作作为参数。对于累积求和，指定了加法操作。随机整数的数组被应用于这个内核函数的输入。内核输出将与输入具有相同的大小。输入和输出向量以及计算累积和所需的时间将在控制台上显示，如下所示：
在一个数组的 10 个元素上运行扫描操作大约需要 2 毫秒。总结来说，在本节中，我们看到了定义映射、减少和扫描操作内核的各种特殊方法。
摘要
本章展示了 PyCUDA 编程的概念。它从使用 PyCUDA 开发一个简单的 Hello, PyCUDA 程序开始。详细讨论了在 C 或 C++ 中定义内核以及在 Python 代码中调用它的概念，以及从 PyCUDA 程序中访问 GPU 设备属性的 API。通过一个简单的程序解释了 PyCUDA 程序中多线程和多块执行机制。使用一个数组加法的简单示例描述了 PyCUDA 程序的基本结构。通过驱动类指令描述了 PyCUDA 代码的简化。详细解释了使用 CUDA 事件来衡量 PyCUDA 程序性能的方法。使用逐元素平方示例解释了驱动类中的 inout 指令和 gpuarray 类的功能。使用 gpuarray 类开发了使用 PyCUDA 计算点积的代码。详细解释了 PyCUDA 中用于矩阵乘法等复杂数学运算的代码。本章的最后部分描述了用于映射、归约和扫描操作的多种内核定义方法。
下一章将在此基础上构建知识，并描述 PyCUDA 中可用的某些高级内核以及使用 PyCUDA 开发计算机视觉应用程序。
问题
-
在 PyCUDA 中使用
SourceModule类定义内核函数时，使用哪种编程语言？将使用哪种编译器来编译这个内核函数？ -
为本章中使用的
myfirst_kernel函数编写一个内核调用函数，块的数量等于 1024 x 1024，每个块中的线程数等于 512 x 512。 -
判断对错：PyCUDA 程序内部的块执行是按顺序进行的。
-
在 PyCUDA 程序中使用
In、Out和inout驱动类原语的优势是什么？ -
编写一个 PyCUDA 程序，使用
gpuarray类将向量中每个元素的值增加 2，向量的大小是任意的。 -
使用 CUDA 事件来测量内核执行时间的优势是什么？
-
判断对错：
gpuarray类是 Python 中numpy库的 GPU 设备版本。
第十二章：使用 PyCUDA 的基本计算机视觉应用
在上一章中，我们学习了与 PyCUDA 相关的重要编程概念。我们还学习了如何使用这些编程概念在 PyCUDA 中开发一些程序。本章将在此基础上构建知识，我们将使用 PyCUDA 来开发基本的图像处理和计算机视觉应用。原子操作和共享内存的并行编程概念也将被详细解释。图像直方图传达了与图像对比度相关的信息，它还可以用作计算机视觉任务的图像特征。本章将详细解释使用 PyCUDA 计算直方图的程序。其他基本的计算机视觉应用，如使用 PyCUDA 进行颜色空间转换、图像加法和图像反转，也将被描述。
本章将涵盖以下主题：
-
使用原子操作和共享内存进行直方图计算
-
使用 PyCUDA 的基本计算机视觉应用
-
从网络摄像头进行图像和视频的颜色空间转换
-
图像加法
-
图像反转
技术要求
本章要求对 Python 编程语言有良好的理解。它还需要任何带有 Nvidia GPU 的计算机或笔记本电脑。本章中使用的所有代码都可以从以下 GitHub 链接下载：github.com/PacktPublishing/Hands-On-GPU-Accelerated-Computer-Vision-with-OpenCV-and-CUDA。查看以下视频以查看代码的实际运行情况：bit.ly/2prC1wI
PyCUDA 中的直方图计算
图像直方图传达了与图像对比度相关的信息，它还可以用作计算机视觉任务的图像特征。直方图表示特定像素值出现的频率。在计算大小为 256 x 256 的 8 位图像的直方图时，65,535 个像素值将在 0-255 的强度值数组上工作。如果为每个像素启动一个线程，那么 65,535 个线程将在 256 个强度值内存位置上工作。
考虑这样一种情况，即大量线程试图修改内存的一小部分。在计算图像直方图时，必须对所有内存位置执行读取-修改-写入操作。这个操作是 d_out[i] ++，其中首先从内存中读取 d_out[i]，然后增加，最后写回内存。然而，当多个线程在相同的内存位置执行此操作时，可能会给出错误的结果。
假设一个内存位置有一个初始值0，线程a和b尝试递增这个内存位置，那么最终答案应该是2。然而，在执行时，可能会发生a和b线程同时读取这个值的情况，那么这两个线程都将得到值0。它们将其递增到1，并且两个线程都将这个1存储在内存中。所以，计算出的答案是1，这是不正确的。
为了理解这可能会多么危险，考虑 ATM 现金取款的例子。假设你在账户中有$50,000 的余额。你有两张同一账户的 ATM 卡。你和你的朋友同时去两个不同的 ATM 取款$40,000。你们同时刷卡；所以，当 ATM 检查余额时，两个都会得到$50,000。如果你俩都取款$40,000，那么两台机器都会查看初始余额，即$50,000。取款金额小于余额，因此两台机器都会给出$40,000。尽管你的余额是$50,000，但你得到了$80,000，这是危险的。为了避免这些情况，在并行编程中使用原子操作，这将在下一节中解释。
使用原子操作
CUDA 提供了一个名为atomicAdd的 API，用于避免并行访问内存位置时的问题。这是一个阻塞操作，这意味着当多个线程尝试访问相同的内存位置时，一次只能有一个线程可以访问该内存位置。其他线程必须等待此线程完成并将答案写入内存。以下是一个使用atomicAdd操作计算直方图的内核函数示例：
import pycuda.autoinit
import pycuda.driver as drv
import numpy
import matplotlib.pyplot as plt
from pycuda.compiler import SourceModule
mod = SourceModule("""
__global__ void atomic_hist(int *d_b, int *d_a, int SIZE)
{
int tid = threadIdx.x + blockDim.x * blockIdx.x;
int item = d_a[tid];
if (tid < SIZE)
{
atomicAdd(&(d_b[item]), 1);
}
}
""")
内核函数有三个参数。第一个参数是计算后存储直方图输出的输出数组。对于 8 位图像，这个数组的大小将是 256。第二个参数是图像强度的展平数组。第三个参数是展平数组的大小。在线程索引处按像素强度索引的直方图数组的内存位置将为每个线程递增。线程的数量等于展平图像数组的大小。
atomicAdd函数用于递增内存位置。它接受两个参数。第一个是我们想要递增的内存位置，第二个是这个位置需要递增的值。atomicadd函数将增加直方图计算在执行时间上的成本。使用原子操作计算直方图的 Python 代码如下：
atomic_hist = mod.get_function("atomic_hist")
import cv2
h_img = cv2.imread("cameraman.tif",0)
h_a=h_img.flatten()
h_a=h_a.astype(numpy.int)
h_result = numpy.zeros(256).astype(numpy.int)
SIZE = h_img.size
NUM_BIN=256
n_threads= int(numpy.ceil((SIZE+NUM_BIN-1) / NUM_BIN))
start = drv.Event()
end=drv.Event()
start.record()
atomic_hist(
drv.Out(h_result), drv.In(h_a), numpy.uint32(SIZE),
block=(n_threads,1,1), grid=(NUM_BIN,1))
end.record()
end.synchronize()
secs = start.time_till(end)*1e-3
print("Time for Calculating Histogram on GPU with shared memory")
print("%fs" % (secs))
plt.stem(h_result)
plt.xlim([0,256])
plt.title("Histogram on GPU")
使用get_function()方法创建指向内核函数的指针引用。图像读取使用 OpenCV 库。如果 Python 中没有安装，您可以从命令提示符执行以下命令：
$pip install opencv-python
然后，可以使用import cv2命令从任何 Python 程序中导入OpenCV库。图像读取函数与本书前面解释过的类似。图像被读取为灰度图像。在 Python 中，图像被存储为一个numpy数组。这个数组被展平为一个向量，以便它可以被一维线程和块操作。也可以在不展平的情况下使用二维线程处理图像。numpy库提供了一个flatten()方法来执行此操作。
块和线程的总数是根据图像的大小和直方图的桶数来计算的。在调用核函数时，将展平的图像数组、空直方图数组和展平数组的大小作为参数传递，同时传递要启动的块和线程数。核函数返回计算出的直方图，可以显示或绘制。
Python 提供了一个包含丰富绘图函数的matplotlib库。该库中的stem函数用于绘制离散的直方图函数。xlim函数用于设置 X 轴的界限。title函数用于给图表添加标题。程序的输出如下所示：
如果直方图没有所有强度的均匀分布，则可能导致对比度差的图像。可以通过直方图均衡化来增强对比度，这会将这种分布转换为均匀分布。直方图还传达了关于图像亮度的信息。如果直方图集中在图表的左侧，则图像会太暗；如果集中在右侧，则图像会太亮。再次强调，可以使用直方图均衡化来纠正这个问题。
计算直方图的核函数也可以使用并行编程中共享内存的概念来开发。这将在以下部分中说明。
使用共享内存
共享内存可以在 GPU 设备上片上使用，因此它比全局内存快得多。共享内存的延迟大约是全球未缓存内存延迟的 100 倍。来自同一块的线程都可以访问共享内存。这在许多需要线程之间共享结果的场景中非常有用。然而，如果不进行同步，这也可能导致混乱或错误的结果。如果一个线程在另一个线程写入之前从内存中读取数据，可能会导致错误的结果。因此，这种内存访问应该得到适当的控制或管理。这通过__syncthreads()指令来完成，它确保在程序继续前进之前，所有对内存的写入操作都已完成。这也被称为屏障。屏障的含义是所有线程都将到达这一行并等待其他线程完成。当所有线程都到达这个屏障后，它们可以继续前进。本节将演示如何从 PyCUDA 程序中使用共享内存。
这种共享内存的概念可以用于计算图像的直方图。内核函数如下所示：
import pycuda.autoinit
import pycuda.driver as drv
import numpy
import matplotlib.pyplot as plt
from pycuda.compiler import SourceModule
mod1 = SourceModule("""
__global__ void atomic_hist(int *d_b, int *d_a, int SIZE)
{
int tid = threadIdx.x + blockDim.x * blockIdx.x;
int offset = blockDim.x * gridDim.x;
__shared__ int cache[256];
cache[threadIdx.x] = 0;
__syncthreads();
while (tid < SIZE)
{
atomicAdd(&(cache[d_a[tid]]), 1);
tid += offset;
}
__syncthreads();
atomicAdd(&(d_b[threadIdx.x]), cache[threadIdx.x]);
}
""")
对于 8 位图像，bins 的数量是 256，所以我们定义的共享内存大小等于块中的线程数，也就是 bins 的数量。我们将为当前块计算一个直方图，因此共享内存被初始化为零，并按前面讨论的方式计算该块的直方图。但这次，结果存储在共享内存中而不是全局内存中。在这种情况下，只有 256 个线程试图访问共享内存中的 256 个内存元素，而不是之前代码中的所有 65,535 个元素。这将有助于减少原子操作中的开销时间。最后一行的最终原子加操作将一个块的直方图添加到整体直方图值中。由于加法是一个累积操作，我们不必担心每个块执行的顺序。以下是如何使用此内核函数计算直方图的 Python 代码示例：
atomic_hist = mod.get_function("atomic_hist")
import cv2
h_img = cv2.imread("cameraman.tif",0)
h_a=h_img.flatten()
h_a=h_a.astype(numpy.int)
h_result = numpy.zeros(256).astype(numpy.int)
SIZE = h_img.size
NUM_BIN=256
n_threads= int(numpy.ceil((SIZE+NUM_BIN-1) / NUM_BIN))
start = drv.Event()
end=drv.Event()
start.record()
atomic_hist(
drv.Out(h_result), drv.In(h_a), numpy.uint32(SIZE),
block=(n_threads,1,1), grid=(NUM_BIN,1),shared= 256*4)
end.record()
end.synchronize()
secs = start.time_till(end)*1e-3
print("Time for Calculating Histogram on GPU with shared memory")
print("%fs" % (secs))
plt.stem(h_result)
plt.xlim([0,256])
plt.title("Histogram on GPU")
代码几乎与上一节相同。唯一的区别在于内核调用。在调用内核时应该定义共享内存的大小。这可以通过内核调用函数中的共享参数来指定。它被指定为256*4，因为共享内存的大小为 256 个整数元素，每个元素需要 4 字节存储。与上一节显示的相同直方图将被显示。
为了检查计算出的直方图的准确性并比较性能，还使用 OpenCV 内置函数calcHist计算直方图，如下面的代码所示：
start = cv2.getTickCount()
hist = cv2.calcHist([h_img],[0],None,[256],[0,256])
end = cv2.getTickCount()
time = (end - start)/ cv2.getTickFrequency()
print("Time for Calculating Histogram on CPU")
print("%fs" % (secs))
calcHist 函数需要五个参数。第一个参数是图像变量的名称。第二个参数在彩色图像的情况下指定通道。对于灰度图像，它为零。第三个参数指定了如果您想计算图像特定部分的直方图，则需要指定掩码。第四个参数指定了箱子的数量，第五个参数指定了强度值的范围。OpenCV 还在 Python 中提供了 getTickCount 和 getTickFrequency 函数来计算 OpenCV 代码的性能。没有共享内存、使用共享内存以及使用 OpenCV 函数的性能如下所示：
没有共享内存的内核函数耗时 1 毫秒，而使用共享内存时为 0.8 毫秒，这进一步证明了使用共享内存可以提高内核函数性能的观点。总结来说，在本节中，我们看到了在 GPU 上计算直方图的两种不同方法。我们还了解了原子操作和共享内存的概念，以及如何在 PyCUDA 中使用它们。
使用 PyCUDA 的基本计算机视觉操作
本节将演示如何使用 PyCUDA 开发简单的计算机视觉应用。在 Python 中，图像不过是二维或三维的 numpy 数组，因此在 PyCUDA 中处理和操作图像与处理多维数组类似。本节将为您提供一个基本的概念，帮助您开发一个简单的应用，您可以用它来利用 PyCUDA 开发复杂的计算机视觉应用。
PyCUDA 中的颜色空间转换
大多数计算机视觉算法都处理灰度图像，因此需要将相机捕获的彩色图像转换为灰度图像。尽管 OpenCV 提供了内置函数来完成此操作，但您也可以通过开发自己的函数来实现。本节将演示如何开发一个 PyCUDA 函数，用于将彩色图像转换为灰度图像。如果已知将图像从一个颜色空间转换为另一个颜色空间的公式，那么本节中显示的函数可以通过替换公式来为任何颜色空间转换编写。
OpenCV 以 BGR 格式捕获和存储图像，其中蓝色是第一个通道，接着是绿色和红色。从 BGR 格式转换为灰度的公式如下：
gray = 0.299*r+0.587*g+0.114*b Where r,g,b are color intensities of red, green and blue channel at a particular location
该函数在图像和视频中的实现将在以下两个子节中展示。
图像上的 BGR 到灰度转换
在本节中，我们将尝试开发将 BGR 图像转换为灰度图像的内核函数。将彩色图像转换为灰度的内核函数如下所示：
import pycuda.driver as drv
from pycuda.compiler import SourceModule
import numpy as np
import cv2
mod = SourceModule \
(
"""
#include<stdio.h>
#define INDEX(a, b) a*256+b
__global__ void bgr2gray(float *d_result,float *b_img, float *g_img, float *r_img)
{
unsigned int idx = threadIdx.x+(blockIdx.x*(blockDim.x*blockDim.y));
unsigned int a = idx/256;
unsigned int b = idx%256;
d_result[INDEX(a, b)] = (0.299*r_img[INDEX(a, b)]+0.587*g_img[INDEX(a, b)]+0.114*b_img[INDEX(a, b)]);
}
"""
)
定义了一个小的 INDEX 函数来计算一个 256 x 256 大小的二维图像的特定索引值。将彩色图像三个通道的展平图像数组作为核函数的输入，其输出是相同大小的灰度图像。INDEX 函数用于将线程索引转换为图像中的特定像素位置。使用所示函数计算该位置的灰度值。以下是将彩色图像转换为灰度图像的 Python 代码示例：
h_img = cv2.imread('lena_color.tif',1)
h_gray=cv2.cvtColor(h_img,cv2.COLOR_BGR2GRAY)
#print a
b_img = h_img[:, :, 0].reshape(65536).astype(np.float32)
g_img = h_img[:, :, 1].reshape(65536).astype(np.float32)
r_img = h_img[:, :, 2].reshape(65536).astype(np.float32)
h_result=r_img
bgr2gray = mod.get_function("bgr2gray")
bgr2gray(drv.Out(h_result), drv.In(b_img), drv.In(g_img),drv.In(r_img),block=(1024, 1, 1), grid=(64, 1, 1))
h_result=np.reshape(h_result,(256,256)).astype(np.uint8)
cv2.imshow("Grayscale Image",h_result)
cv2.waitKey(0)
cv2.destroyAllWindows()
使用 OpenCV 的 imread 函数读取彩色图像。图像的大小应为 256 x 256，因此如果不是这个大小，则应使用 cv2.resize 函数将其转换为该大小。彩色图像以 BGR 格式存储，因此使用 Python 的数组切片将其蓝色、绿色和红色通道分离。这些数组被展平，以便可以将它们传递给核函数。
核函数使用三个颜色通道作为输入和一个用于存储输出灰度图像的数组进行调用。核函数将在每个像素位置计算灰度值，并返回一个灰度图像的展平数组。使用 numpy 库的 reshape 函数将此结果数组转换回原始图像大小。OpenCV 的 imshow 函数需要无符号整数数据类型来显示图像，因此数组也转换为 uint8 数据类型。灰度图像显示在屏幕上，如下面的截图所示：
在摄像头视频中执行 BGR 到灰度的转换
在上一节中开发的将图像转换为灰度的相同核函数可以用来将来自摄像头的视频转换为灰度。以下是其 Python 代码示例：
cap = cv2.VideoCapture(0)
bgr2gray = mod.get_function("bgr2gray")
while(True):
# Capture frame-by-frame
ret, h_img = cap.read()
h_img = cv2.resize(h_img,(256,256),interpolation = cv2.INTER_CUBIC)
b_img = h_img[:, :, 0].reshape(65536).astype(np.float32)
g_img = h_img[:, :, 1].reshape(65536).astype(np.float32)
r_img = h_img[:, :, 2].reshape(65536).astype(np.float32)
h_result=r_img
bgr2gray(drv.Out(h_result), drv.In(b_img), drv.In(g_img),drv.In(r_img),block=(1024, 1, 1), grid=(64, 1, 1))
h_result=np.reshape(h_result,(256,256)).astype(np.uint8)
cv2.imshow("Grayscale Image",h_result)
# Display the resulting frame
cv2.imshow('Original frame',h_img)
if cv2.waitKey(50) & 0xFF == ord('q'):
break
# When everything done, release the capture
cap.release()
cv2.destroyAllWindows()
Python 中的 OpenCV 提供了一个 VideoCapture 类来从摄像头捕获视频。它需要一个相机设备索引作为参数。对于摄像头，它被指定为零。然后，启动一个连续的 while 循环来从摄像头捕获帧。使用捕获对象的 read 方法读取这些帧。这些帧使用 cv2 库的 resize 函数调整大小到 256 x 256。这些帧是彩色图像，因此从它们中分离出三个通道并展平，以便可以将它们传递给核函数。以与上一节相同的方式调用核函数，并将结果重塑以在屏幕上显示。以下是一帧摄像头流的代码输出：
网络摄像头流将一直持续到在键盘上按下 q 键。总结一下，我们已经在 PyCUDA 中开发了一个内核函数，用于将 BGR 格式的彩色图像转换为灰度图像，这个函数既可以处理图像也可以处理视频。这些内核函数可以通过替换相同的方程来修改，以用于其他颜色空间的转换。
PyCUDA 中的图像加法
当两个图像大小相同时，可以执行两个图像的加法。它执行两个图像的逐像素加法。假设在两个图像中，(0,0) 像素的强度值分别为 100 和 150，那么结果图像中的强度值将是 250，这是两个强度值的和，如下所示：
result = img1 + img2
OpenCV 加法是一个饱和操作，这意味着如果加法的结果超过 255，则它将被饱和在 255。因此，相同的函数性作为 PyCUDA 内核函数实现。执行图像加法的代码如下所示：
import pycuda.driver as drv
from pycuda.compiler import SourceModule
import numpy as np
import cv2
mod = SourceModule \
(
"""
__global__ void add_num(float *d_result, float *d_a, float *d_b,int N)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
while (tid < N)
{
d_result[tid] = d_a[tid] + d_b[tid];
if(d_result[tid]>255)
{
d_result[tid]=255;
}
tid = tid + blockDim.x * gridDim.x;
}
}
"""
)
img1 = cv2.imread('cameraman.tif',0)
img2 = cv2.imread('circles.png',0)
h_img1 = img1.reshape(65536).astype(np.float32)
h_img2 = img2.reshape(65536).astype(np.float32)
N = h_img1.size
h_result=h_img1
add_img = mod.get_function("add_num")
add_img(drv.Out(h_result), drv.In(h_img1), drv.In(h_img2),np.uint32(N),block=(1024, 1, 1), grid=(64, 1, 1))
h_result=np.reshape(h_result,(256,256)).astype(np.uint8)
cv2.imshow("Image after addition",h_result)
cv2.waitKey(0)
cv2.destroyAllWindows()
内核函数与上一章中看到的数组加法内核函数类似。内核函数中添加了饱和条件，表示如果像素强度在加法后超过 255，则它将被饱和在 255。读取两个大小相同的图像，展平，并转换为单精度浮点数据类型。这些展平的图像及其大小作为参数传递给内核函数。内核函数计算的结果被重塑为原始图像大小，并使用 imshow 函数转换为无符号整型以显示。结果如下面的截图所示，包括原始图像：
同样的内核函数经过轻微修改后，可以用于其他算术和逻辑运算。
使用 gpuarray 在 PyCUDA 中进行图像反转
除了算术运算外，NOT 运算也广泛用于图像反转，其中黑色转换为白色，白色转换为黑色。它可以表示为以下方程：
result_image = 255 - input_image
在前面的方程中，255 表示 8 位图像的最大强度值。PyCUDA 提供的 gpuarray 类用于开发图像反转程序，如下所示：
import pycuda.driver as drv
import numpy as np
import cv2
import pycuda.gpuarray as gpuarray
import pycuda.autoinit
img = cv2.imread('circles.png',0)
h_img = img.reshape(65536).astype(np.float32)
d_img = gpuarray.to_gpu(h_img)
d_result = 255- d_img
h_result = d_result.get()
h_result=np.reshape(h_result,(256,256)).astype(np.uint8)
cv2.imshow("Image after addition",h_result)
cv2.waitKey(0)
cv2.destroyAllWindows()
图像被读取为灰度图像，展平并转换为单精度浮点数据类型以进行进一步处理。它使用 gpuarray 类的 to_gpu 方法上传到 GPU。使用前面的方程在 GPU 上执行反转，然后使用 get() 方法将结果下载回主机。结果通过重塑为原始图像大小显示在屏幕上，如下面的截图所示：
总结，本节展示了 PyCUDA 在开发基本计算机视觉操作中的应用，如颜色空间转换、图像加法和图像反转。这个概念可以用于使用 PyCUDA 开发复杂的计算机视觉应用。
摘要
本章描述了在开发简单的计算机视觉应用中使用 PyCUDA。它描述了使用 PyCUDA 计算数组直方图的过程。直方图是图像的一个非常重要的统计全局特征，可以用来获取关于图像的重要信息。本章以直方图计算为例，详细解释了原子操作和共享内存的概念。Python 中的图像存储为numpy数组，因此在 PyCUDA 中操作图像类似于修改多维numpy数组。本章描述了 PyCUDA 在多种基本计算机视觉应用中的使用，例如图像加法、图像反转和颜色空间转换。本章中描述的概念可以用于使用 PyCUDA 开发复杂的计算机视觉应用。
本章也标志着本书的结束，本书描述了使用 CUDA 编程和 GPU 硬件加速计算机视觉应用。
问题
-
判断对错：使用
d_out[i]++行而不是atomicadd操作将在直方图计算中产生准确的结果。 -
使用原子操作共享内存的优势是什么？
-
当在内核中使用共享内存时，内核调用函数中的修改是什么？
-
通过计算图像的直方图可以获得哪些信息？
-
判断对错：本章开发的将 BGR 转换为灰度的内核函数也将适用于 RGB 转换为灰度。
-
为什么本章中展示的所有示例都将图像展平？这是一个强制性的步骤吗？
-
为什么在显示之前将图像从
numpy库的uint8数据类型转换？
第十三章：评估
第一章
-
提高性能的三种选项如下：
-
拥有更快的时钟速度
-
单个处理器每个时钟周期完成更多的工作
-
许多可以并行工作的小型处理器。这个选项被 GPU 用来提高性能。
-
-
正确
-
CPU 的设计是为了提高延迟，而 GPU 的设计是为了提高吞吐量。
-
汽车需要 4 小时才能到达目的地，但它只能容纳 5 人，而可以容纳 40 人的公交车需要 6 小时才能到达目的地。公交车每小时可以运输 6.66 人，而汽车每小时可以运输 1.2 人。因此，汽车具有更好的延迟，而公交车具有更好的吞吐量。
-
图像不过是一个二维数组。大多数计算机视觉应用都涉及这些二维数组的处理。这涉及到对大量数据进行类似操作，这些操作可以通过 GPU 高效地执行。因此，GPU 和 CUDA 在计算机视觉应用中非常有用。
-
错误
-
printf语句在主机上执行
第二章
- 通过传递参数作为值来减去两个数字的 CUDA 程序如下：
include <iostream>
#include <cuda.h>
#include <cuda_runtime.h>
__global__ void gpuSub(int d_a, int d_b, int *d_c)
{
*d_c = d_a - d_b;
}
int main(void)
{
int h_c;
int *d_c;
cudaMalloc((void**)&d_c, sizeof(int));
gpuSub << <1, 1 >> > (4, 1, d_c);
cudaMemcpy(&h_c, d_c, sizeof(int), cudaMemcpyDeviceToHost);
printf("4-1 = %d\n", h_c);
cudaFree(d_c);
return 0;
}
- 通过传递参数作为引用来乘以两个数字的 CUDA 程序如下：
#include <iostream>
#include <cuda.h>
#include <cuda_runtime.h>
__global__ void gpuMul(int *d_a, int *d_b, int *d_c)
{
*d_c = *d_a * *d_b;
}
int main(void)
{
int h_a,h_b, h_c;
int *d_a,*d_b,*d_c;
h_a = 1;
h_b = 4;
cudaMalloc((void**)&d_a, sizeof(int));
cudaMalloc((void**)&d_b, sizeof(int));
cudaMalloc((void**)&d_c, sizeof(int));
cudaMemcpy(d_a, &h_a, sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(d_b, &h_b, sizeof(int), cudaMemcpyHostToDevice);
gpuMul << <1, 1 >> > (d_a, d_b, d_c);
cudaMemcpy(&h_c, d_c, sizeof(int), cudaMemcpyDeviceToHost);
printf("Passing Parameter by Reference Output: %d + %d = %d\n", h_a, h_b, h_c);
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
return 0;
}
- 为
gpuMul内核启动 5000 个线程的三种方法如下：
1\. gpuMul << <25, 200 >> > (d_a, d_b, d_c);
2\. gpuMul << <50, 100 >> > (d_a, d_b, d_c);
3\. gpuMul << <10, 500 >> > (d_a, d_b, d_c);
-
错误
-
查找具有 5.0 或更高版本的 GPU 设备的程序如下
int main(void)
{
int device;
cudaDeviceProp device_property;
cudaGetDevice(&device);
printf("ID of device: %d\n", device);
memset(&device_property, 0, sizeof(cudaDeviceProp));
device_property.major = 5;
device_property.minor = 0;
cudaChooseDevice(&device, &device_property);
printf("ID of device which supports double precision is: %d\n", device);
cudaSetDevice(device);
}
- 查找数字立方的 CUDA 程序如下：
#include "stdio.h"
#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>
#define N 50
__global__ void gpuCube(float *d_in, float *d_out)
{
//Getting thread index for current kernel
int tid = threadIdx.x; // handle the data at this index
float temp = d_in[tid];
d_out[tid] = temp*temp*temp;
}
int main(void)
{
float h_in[N], h_out[N];
float *d_in, *d_out;
cudaMalloc((void**)&d_in, N * sizeof(float));
cudaMalloc((void**)&d_out, N * sizeof(float));
for (int i = 0; i < N; i++)
{
h_in[i] = i;
}
cudaMemcpy(d_in, h_in, N * sizeof(float), cudaMemcpyHostToDevice);
gpuSquare << <1, N >> >(d_in, d_out);
cudaMemcpy(h_out, d_out, N * sizeof(float), cudaMemcpyDeviceToHost);
printf("Cube of Number on GPU \n");
for (int i = 0; i < N; i++)
{
printf("The cube of %f is %f\n", h_in[i], h_out[i]);
}
cudaFree(d_in);
cudaFree(d_out);
return 0;
}
-
特定应用的通信模式如下所示：
-
图像处理 - 模板
-
移动平均 - 聚合
-
按升序排序数组 - 散射
-
在数组中查找数字的立方 - 映射
-
第三章
- 选择线程数和块数的最佳方法如下：
gpuAdd << <512, 512 >> >(d_a, d_b, d_c);
每个块可以启动的线程数量有限，对于最新的处理器来说，这个数量是 512 或 1024。同样，每个网格的块数量也有限制。因此，如果有大量线程，那么最好通过少量块和线程来启动内核，如上所述。
- 以下是为查找 50000 个数字的立方而编写的 CUDA 程序：
#include "stdio.h"
#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>
#define N 50000
__global__ void gpuCube(float *d_in, float *d_out)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
while (tid < N)
{
float temp = d_in[tid];
d_out[tid] = temp*temp*temp;
tid += blockDim.x * gridDim.x;
}
}
int main(void)
{
float h_in[N], h_out[N];
float *d_in, *d_out;
cudaMalloc((void**)&d_in, N * sizeof(float));
cudaMalloc((void**)&d_out, N * sizeof(float));
for (int i = 0; i < N; i++)
{
h_in[i] = i;
}
cudaMemcpy(d_in, h_in, N * sizeof(float), cudaMemcpyHostToDevice);
gpuSquare << <512, 512 >> >(d_in, d_out);
cudaMemcpy(h_out, d_out, N * sizeof(float), cudaMemcpyDeviceToHost);
printf("Cube of Number on GPU \n");
for (int i = 0; i < N; i++)
{
printf("The cube of %f is %f\n", h_in[i], h_out[i]);
}
cudaFree(d_in);
cudaFree(d_out);
return 0;
}
-
正确，因为它只需要访问局部内存，这是一种更快的内存。
-
当内核的变量不适合寄存器文件时，它们使用局部内存。这被称为寄存器溢出。因为一些数据不在寄存器中，它将需要更多时间从内存中检索它。这将花费更多时间，因此程序的性能将受到影响。
-
不，因为所有线程都在并行运行。所以数据可能在写入之前就被读取，因此可能不会给出期望的输出。
-
正确。在原子操作中，当一个线程正在访问特定的内存位置时，其他所有线程都必须等待。当许多线程访问相同的内存位置时，这将产生时间开销。因此，原子操作会增加 CUDA 程序的执行时间。
-
Stencil 通信模式非常适合纹理内存。
-
当在
if语句中使用__syncthreads指令时，对于具有此条件的线程，false永远不会到达这个点，__syncthreads将持续等待所有线程到达这个点。因此，程序将永远不会终止。
第四章
-
CPU 计时器将包括操作系统中的线程延迟和调度的时间开销，以及其他许多因素。使用 CPU 测量的时间也将取决于高精度 CPU 计时器的可用性。主机在 GPU 内核运行时经常执行异步计算，因此 CPU 计时器可能无法给出内核执行的准确时间。
-
从
C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0\libnvvp打开 Nvidia Visual profiler。然后，转到 -> 新会话并选择矩阵乘法示例的.exe文件。您可以可视化您代码的性能。 -
除以零、变量类型或大小不正确、不存在变量、下标超出范围等是语义错误的例子。
-
可以给出线程发散的例子如下：
__global__ void gpuCube(float *d_in, float *d_out)
{
int tid = threadIdx.x;
if(tid%2 == 0)
{
float temp = d_in[tid];
d_out[tid] = temp*temp*temp;
}
else
{
float temp = d_in[tid];
d_out[tid] = temp*temp*temp;
}
}
在代码中，奇数和偶数线程执行不同的操作，因此它们完成所需的时间不同。在 if 语句之后，这些线程将再次合并。这将产生时间开销，因为快速线程必须等待慢速线程。这将降低代码的性能。
-
cudaHostAlloc函数应谨慎使用，因为这种内存不会被交换到磁盘上；您的系统可能会耗尽内存。这可能会影响系统上运行的其他应用程序的性能。 -
在 CUDA 流操作中，操作顺序很重要，因为我们希望重叠内存复制操作与内核执行操作。因此，操作队列应设置为这些操作可以相互重叠，否则使用 CUDA 流不会提高程序的性能。
-
对于 1024 x 1024 的图像，线程数应为 32x32（如果您的系统支持每个块 1024 个线程），块数也应为 32 x 32，这可以通过将图像大小除以每个块线程数来确定。
第五章
-
图像处理和计算机视觉领域之间存在差异。图像处理关注通过修改像素值来提高图像的视觉质量，而计算机视觉关注从图像中提取重要信息。因此，在图像处理中，输入和输出都是图像，而在计算机视觉中，输入是图像，但输出是从该图像中提取的信息。
-
OpenCV 库在 C、C++、Java 和 Python 语言中都有接口，并且可以在 Windows、Linux、Mac 和 Android 等所有操作系统上使用，而无需修改单行代码。这个库还可以利用多核处理。它可以利用 OpenGL 和 CUDA 进行并行处理。由于 OpenCV 轻量级，它也可以在树莓派等嵌入式平台上使用。这使得它在实际场景中部署计算机视觉应用成为理想选择。
-
初始化图像为红色的命令如下：
Mat img3(1960,1960, CV_64FC3, Scalar(0,0,255) )
- 从网络摄像头捕获视频并将其存储在磁盘上的程序如下：
#include <opencv2/opencv.hpp>
#include <iostream>
using namespace cv;
using namespace std;
int main(int argc, char* argv[])
{
VideoCapture cap(0);
if (cap.isOpened() == false)
{
cout << "Cannot open Webcam" << endl;
return -1;
}
Size frame_size(640, 640);
int frames_per_second = 30;
VideoWriter v_writer("images/video.avi", VideoWriter::fourcc('M', 'J', 'P', 'G'), frames_per_second, frame_size, true);
cout<<"Press Q to Quit" <<endl;
String win_name = "Webcam Video";
namedWindow(win_name); //create a window
while (true)
{
Mat frame;
bool flag = cap.read(frame); // read a new frame from video
imshow(win_name, frame);
v_writer.write(frame);
if (waitKey(1) == 'q')
{
v_writer.release();
break;
}
}
return 0;
}
-
OpenCV 使用 BGR 颜色格式来读取和显示图像。
-
从网络摄像头捕获视频并将其转换为灰度的程序如下：
#include <opencv2/opencv.hpp>
#include <iostream>
using namespace cv;
using namespace std;
int main(int argc, char* argv[])
{
VideoCapture cap(0);
if (cap.isOpened() == false)
{
cout << "Cannot open Webcam" << endl;
return -1;
}
cout<<"Press Q to Quit" <<endl;
String win_name = "Webcam Video";
namedWindow(win_name); //create a window
while (true)
{
Mat frame;
bool flag = cap.read(frame); // read a new frame from video
cvtColor(frame, frame,cv::COLOR_BGR2GRAY);
imshow(win_name, frame);
if (waitKey(1) == 'q')
{
break;
}
}
return 0;
}
- 测量加法和减法操作性能的 OpenCV 程序如下：
#include <iostream>
#include "opencv2/opencv.hpp"
int main (int argc, char* argv[])
{
//Read Two Images
cv::Mat h_img1 = cv::imread("images/cameraman.tif");
cv::Mat h_img2 = cv::imread("images/circles.png");
//Create Memory for storing Images on device
cv::cuda::GpuMat d_result1,d_result2,d_img1, d_img2;
cv::Mat h_result1,h_result2;
int64 work_begin = getTickCount();
//Upload Images to device
d_img1.upload(h_img1);
d_img2.upload(h_img2);
cv::cuda::add(d_img1,d_img2, d_result1);
cv::cuda::subtract(d_img1, d_img2,d_result2);
//Download Result back to host
d_result1.download(h_result1);
d_result2.download(h_result2);
int64 delta = getTickCount() - work_begin;
//Frequency of timer
double freq = getTickFrequency();
double work_fps = freq / delta;
std::cout<<"Performance of Thresholding on CPU: " <<std::endl;
std::cout <<"Time: " << (1/work_fps) <<std::endl;
cv::waitKey();
return 0;
}
- OpenCV 程序用于执行位运算 AND 和 OR 操作如下：
include <iostream>
#include "opencv2/opencv.hpp"
int main (int argc, char* argv[])
{
cv::Mat h_img1 = cv::imread("images/cameraman.tif");
cv::Mat h_img2 = cv::imread("images/circles.png");
cv::cuda::GpuMat d_result1,d_result2,d_img1, d_img2;
cv::Mat h_result1,h_result2;
d_img1.upload(h_img1);
d_img2.upload(h_img2);
cv::cuda::bitwise_and(d_img1,d_img2, d_result1);
cv::cuda::biwise_or(d_img1, d_img2,d_result2);
d_result1.download(h_result1);
d_result2.download(h_result2);
cv::imshow("Image1 ", h_img1);
cv::imshow("Image2 ", h_img2);
cv::imshow("Result AND operation ", h_result1);
cv::imshow("Result OR operation ", h_result2);
cv::waitKey();
return 0;
}
第六章
- 打印任何颜色图像在
(200,200)位置像素强度的 OpenCV 函数如下：
cv::Mat h_img2 = cv::imread("images/autumn.tif",1);
cv::Vec3b intensity1 = h_img1.at<cv::Vec3b>(cv::Point(200, 200));
std::cout<<"Pixel Intensity of color Image at (200,200) is:" << intensity1 << std::endl;
- 使用双线性插值方法将图像调整大小到
(300,200)像素的 OpenCV 函数如下：
cv::cuda::resize(d_img1,d_result1,cv::Size(300, 200), cv::INTER_LINEAR);
- 使用
AREA插值将图像上采样 2 倍的 OpenCV 函数如下：
int width= d_img1.cols;
int height = d_img1.size().height;
cv::cuda::resize(d_img1,d_result2,cv::Size(2*width, 2*height), cv::INTER_AREA);
-
错误。随着滤波器大小的增加，模糊程度也会增加。
-
错误。中值滤波器不能去除高斯噪声。它可以去除椒盐噪声。
-
在应用拉普拉斯算子以去除噪声敏感性之前，必须使用平均或高斯滤波器对图像进行模糊处理。
-
实现顶帽和黑帽形态学操作的 OpenCV 函数如下：
cv::Mat element = cv::getStructuringElement(cv::MORPH_RECT,cv::Size(5,5));
d_img1.upload(h_img1);
cv::Ptr<cv::cuda::Filter> filtert,filterb;
filtert = cv::cuda::createMorphologyFilter(cv::MORPH_TOPHAT,CV_8UC1,element);
filtert->apply(d_img1, d_resulte);
filterb = cv::cuda::createMorphologyFilter(cv::MORPH_BLACKHAT,CV_8UC1,element);
filterb->apply(d_img1, d_resultd);
第七章
- 从视频中检测黄色物体的 OpenCV 代码如下：请注意，这里没有重复样板代码。
cuda::cvtColor(d_frame, d_frame_hsv, COLOR_BGR2HSV);
//Split HSV 3 channels
cuda::split(d_frame_hsv, d_frame_shsv);
//Threshold HSV channels for Yellow color
cuda::threshold(d_frame_shsv[0], d_thresc[0], 20, 30, THRESH_BINARY);
cuda::threshold(d_frame_shsv[1], d_thresc[1], 100, 255, THRESH_BINARY);
cuda::threshold(d_frame_shsv[2], d_thresc[2], 100, 255, THRESH_BINARY);
//Bitwise AND the channels
cv::cuda::bitwise_and(d_thresc[0], d_thresc[1],d_intermediate);
cv::cuda::bitwise_and(d_intermediate, d_thresc[2], d_result);
d_result.download(h_result);
imshow("Thresholded Image", h_result);
imshow("Original", frame);
-
当物体的颜色与背景颜色相同时，基于颜色的目标检测将失败。即使有光照变化，也可能失败。
-
Canny 边缘检测算法的第一步是高斯模糊，这可以去除图像中存在的噪声。之后，计算梯度。因此，检测到的边缘将比之前看到的其他边缘检测算法受噪声影响更小。
-
当图像受到高斯或椒盐噪声的影响时，霍夫变换的结果非常差。为了改善结果，必须在预处理步骤中通过高斯和中值滤波器对图像进行滤波。
-
当计算 FAST 关键点的强度阈值较低时，则更多的关键点将通过段测试并被分类为关键点。随着这个阈值的增加，检测到的关键点数量将逐渐减少。
-
在 SURF 中，Hessian 阈值的较大值将导致更少但更显著的特征点，而较小值将导致更多但不太显著的特征点。
-
当 Haar 级联的缩放因子从 1.01 增加到 1.05 时，图像大小在每一尺度上都会以更大的因子减小。因此，每帧需要处理的图像更少，这减少了计算时间；然而，这可能导致无法检测到某些对象。
-
MoG相比于GMG算法在背景减法方面更快且噪声更少。可以将开闭等形态学操作应用于 GMG 的输出，以减少存在的噪声。
第八章
-
Jetson TX1 在每秒 Tera 级浮点运算性能方面优于 Raspberry Pi。因此，Jetson TX1 可以用于计算密集型应用，如计算机视觉和深度学习，以实现实时部署。
-
Jetson TX1 开发板支持多达六个 2 通道或三个 4 通道相机。它附带一个 500 万像素的相机。
-
必须使用 USB 集线器来连接 Jetson TX1 与超过两个 USB 外设。
-
True
-
False. Jetson TX1 包含一个 1.73 GHz 运行的 ARM Cortex A57 四核 CPU。
-
尽管 Jetson TX1 预装了 Ubuntu，但它不包含计算机视觉应用所需的任何软件包。Jetpack 包含 Tegra (L4T) 板支持包的 Linux，TensorRT，用于计算机视觉应用中的深度学习推理，最新的 CUDA 工具包，cuDNN，这是 CUDA 深度神经网络库，Visionworks，也用于计算机视觉和深度学习应用，以及 OpenCV。因此，通过安装 Jetpack，我们可以快速安装构建计算机视觉应用所需的全部软件包。
第九章
-
Jetson TX1 上的 GPU 设备全局内存大约为 4 GB，GPU 时钟速度约为 1 GHz。这个时钟速度比本书之前使用的 GeForce 940 GPU 慢。内存时钟速度仅为 13 MHz，而 GeForce 940 为 2.505 GHz，这使得 Jetson TX1 更慢。L2 缓存为 256 KB，而 GeForce 940 为 1 MB。大多数其他属性与 GeForce 940 相似。
-
True
-
在最新的 Jetpack 中，OpenCV 没有编译 CUDA 支持，也没有 GStreamer 支持，这是从代码中访问相机所需的。因此，移除 Jetpack 中包含的 OpenCV 安装，并使用 CUGA 和 GStreamer 支持编译新的 OpenCV 版本是个好主意。
-
False. OpenCV 可以从连接到 Jetson TX1 板的 USB 和 CSI 相机捕获视频。
-
True. CSI 相机更接近硬件，因此读取帧的速度比 USB 相机快，因此在计算密集型应用中最好使用 CSI 相机。
-
Python OpenCV 绑定不支持 CUDA 加速，因此对于计算密集型任务，最好使用 C++ OpenCV 绑定。
-
No. Jetson TX1 预装了 python2 和 python3 解释器，同时 OpenCV 也为 Jetson TX1 编译了；它还安装了 python 二进制文件，因此无需单独安装 python OpenCV 绑定。
第十章
-
Python 是开源的，拥有庞大的用户社区，他们通过模块为语言做出贡献。这些模块可以轻松地用少量代码在短时间内开发应用程序。Python 语言的语法易于阅读和解释，这使得它对新程序员来说更容易学习。它是一种允许逐行执行代码的解释型语言。这些都是 Python 相对于 C/C++ 的几个优点。
-
在编译型语言中，整个代码被检查并转换为机器代码，而在解释型语言中，每次只翻译一条语句。解释型语言分析源代码所需的时间较少，但与编译型语言相比，整体执行时间较慢。解释型语言不会像编译型语言那样生成中间代码。
-
错误。Python 是一种解释型语言，这使得它比 C/C++ 慢。
-
PyOpenCL 可以利用任何图形处理单元，而 PyCUDA 需要 Nvidia GPU 和 CUDA 工具包。
-
正确。Python 允许在 Python 脚本中包含 C/C++ 代码，因此计算密集型任务可以写成 C/C++ 代码以实现更快的处理，并为它创建 Python 包装器。PyCUDA 可以利用这一功能来处理内核代码。
第十一章
-
C/C++ 编程语言用于在
SourceModule类中编写内核函数，并且这个内核函数由nvcc（Nvidia C）编译器编译。 -
内核调用函数如下：
myfirst_kernel(block=(512,512,1),grid=(1024,1014,1))
-
错误。在 PyCUDA 程序中，块执行的顺序是随机的，PyCUDA 程序员无法确定。
-
驱动类中的指令消除了为数组单独分配内存、将其上传到设备以及将结果下载回主机的要求。所有操作都在内核调用期间同时执行。这使得代码更简单，更容易阅读。
-
在数组中每个元素加二的 PyCUDA 代码如下所示：
import pycuda.gpuarray as gpuarray
import numpy
import pycuda.driver as drv
start = drv.Event()
end=drv.Event()
start.record()
start.synchronize()
n=10
h_b = numpy.random.randint(1,5,(1,n))
d_b = gpuarray.to_gpu(h_b.astype(numpy.float32))
h_result = (d_b + 2).get()
end.record()
end.synchronize()
print("original array:")
print(h_b)
print("doubled with gpuarray:")
print(h_result)
secs = start.time_till(end)*1e-3
print("Time of adding 2 on GPU with gpuarray")
print("%fs" % (secs))
-
使用 Python 时间测量选项来测量 PyCUDA 程序的性能不会给出准确的结果。它将包括许多其他因素中的线程延迟在操作系统中的时间开销和调度。使用时间类测量的时间也将取决于高精度 CPU 定时器的可用性。很多时候，主机在进行异步计算的同时 GPU 内核正在运行，因此 Python 的 CPU 计时器可能无法给出内核执行的正确时间。我们可以通过使用 CUDA 事件来克服这些缺点。
-
正确
第十二章
-
错误。这一行代表一个读取-修改-写入操作，当多个线程试图增加相同的内存位置时，如直方图计算的情况，可能会产生错误的结果。
-
在使用共享内存的情况下，较少的线程试图访问共享内存中的 256 个内存元素，而不是没有共享内存时所有线程的情况。这将有助于减少原子操作中的时间开销。
-
在使用共享内存的情况下，内核调用函数如下：
atomic_hist(
drv.Out(h_result), drv.In(h_a), numpy.uint32(SIZE),
block=(n_threads,1,1), grid=(NUM_BIN,1),shared= 256*4)
在调用内核时，应该定义共享内存的大小。这可以通过在内核调用函数中使用共享参数来指定。
-
直方图是一种统计特征，它提供了关于图像对比度和亮度的关键信息。如果它具有均匀分布，那么图像将具有良好的对比度。直方图还传达了关于图像亮度的信息。如果直方图集中在图表的左侧，那么图像将太暗，如果集中在右侧，那么图像将太亮。
-
真的。因为 RGB 和 BGR 颜色格式相同，只是通道的顺序不同。转换的方程式仍然保持不变。
-
与多维线程和块相比，处理单维线程和块更简单。它简化了内核函数内部的索引机制，因此在每个章节中出现的示例中都进行了这种简化。如果我们正在处理多维线程和块，则这不是强制性的。
-
imshow函数，用于在屏幕上显示图像，需要一个无符号整数的图像。因此，在屏幕上显示之前，所有由内核函数计算的结果都转换为numpy库的uint8数据类型。
扫一扫
21
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
