VisualStudio2019配置CUDA v10.0 解决CUDA模板无法生成的问题 实现奇偶排序算法

---

前言

初次尝试CUDA编程，在Windows上配置CUDA环境属实麻烦。而且Visual Studio 2019在我做实验时还不支持CUDA，只能自己一步步踩坑配置一个模板…

---

一、 VS2019配置CUDA v10.0

1. 前置条件：已经安装了Visual Studio。建议安装2017及以下版本，会省很多事情。
2. 下载CUDA，可以选择旧版本，但是一定要注意CUDA版本和自己的显卡是否适配，是否支持cuDNN加速。如果选择旧版本安装，就选择Legacy Releases，版本号为v10.0，进行下载即可。地址为

https://developer.nvidia.com/cuda-downloads

3. 下载完成后，运行文件。注意，此时出现的地址为临时文件的存放地址，并不是最终安装的地址，可以随意选择。
   

4. 建议初学者选择精简安装即可，即使选择自定义安装也需要安装所有的组件。截至我当初做实验时，也就是2020年上半年，CUDA10.0乃至最新的10.2都只支持Visual Studio 2017及以下版本，而我安装的是Visual Studio 2019，其中并没有CUDA组件，需要进行一系列操作，相当繁琐，强烈建议安装Visual Studio 2017。如果已经安装了2019的话，之后会提供一个解决办法。

5. 如果正确安装完成后，再打开Visual Studio，新建项目中即可看到对应的CUDA模板工程文件。

---

二、解决模板无法生成的问题

（1）思路

可以得知，目前VS2019并不支持CUDA，2017是支持的。同时VS版本往往是向下兼容的，因此，通过安装2017的编译器，采取一些小措施修改一下，即可解决模板无法生成的问题。

（2）安装VS2017编译器

首先，请打开Visual Studio Installer，安装MSVC v141 – VS 2017 C++ x64/x86生成工具。这个工具对应了2017版本的编译器。

（3）原因分析

在默认安装CUDA的情况下，可以在此处找到sample文件：

C:\ProgramData\NVIDIA Corporation\CUDA Samples\v10.0

选择其中一个，比如：

C:\ProgramData\NVIDIA Corporation\CUDA Samples\v10.0\1_Utilities\deviceQuery

进入后加载deviceQuery_vs2017.sln，会提示：intelliSense不可用。为了追踪项目加载失败的原因，我们需要设置环境变量追踪一下。

打开Developer Command Prompt for VS 2019，输入set TraceDesignTime=true，再输入devenv使环境变量生效。

之后在跳出来的Visual Studio窗口打开之前的sln文件，下方就会显示日志地址了。

找到该日志，在C:\Users\你的用户名\AppData\Local\Temp\，在最下方有以下提示。

现在，就可以解决问题了。

（4）问题解决

打开NVIDIA GPU Computing Toolkit目录下的

CUDA\v10.0\extras\visual_studio_integration\ MSBuildExtension

可以看到其中有4个文件，有一个就是CUDA v10.0.props，就是我们想要的配置文件。将其复制到图片中所说的目录下：

D:\Software\Microsoft Visual Studio 2019\MSBuild\Microsoft\VC\v160\BuildCustomizations

注意，每个人安装Visual Studio的位置不同。之后，再打开deviceQuery_vs2017.sln文件，右上角在“解决方案资源管理器”的项目上右键选择重新生成，就不会有任何报错了。这时，编译一下，可以看到结果。

当出现Result=PASS时，代表没有问题了。之后，需要更改一下设置属性，如下图所示，这样是为了保证可以搜寻到生成的文件。然后，更改一下.cpp文件的名称作为初始名称，并且添加一个文件，将其命名为kernel.cu作为核文件，以后内核函数将会在这里编写。

另外，在编译过程中会出现某些函数比如__syncthreads()，或者内核函数调用中的<<< >>>等，可能会无法识别。无需在意，这是Visual Studio无法识别，CUDA专用编译器nvcc可以正常编译。

最后的最后，选择项目-导出模板，设置一下，比如这样。

之后你就会发现，在新建项目中，已经可以找到CUDA模板了。

---

三、奇偶排序算法实现及结果分析

（1）串行部分

首先，先写串行的算法作为基准。代码如下：

//串行serial.c
#include<stdio.h>
#include<stdlib.h>
#include<string.h>
#include <sys/time.h>

#define GET_TIME(now) { \
struct timeval t; \
gettimeofday(&t, NULL); \
now = t.tv_sec + t.tv_usec/1000000.0; \
}
#define MAXN 130000


void swap(int *a, int *b)
{
    int temp;
    temp = *a;
    *a = *b;
    *b = temp;
}

void Odd_even_sort(int num[], int size)
{
	int sorted = 0;
  	while(!sorted)
  	{
		sorted = 1;
		for(int i = 1; i < size - 1; i += 2)
		{
			if(num[i] > num[i + 1])
			{
				swap(&num[i], &num[i + 1]);
				sorted = 0;
			}
		}
		for(int i = 0; i < size - 1; i += 2)
		{
			if(num[i] > num[i + 1])
			{
				swap(&num[i], &num[i+1]);
				sorted = 0;
			}
		}
	}
}

int main(void)
{
	double start, stop;
	int size = MAXN;
	int num[size];
	
	for(int i = 0; i < size; i++)
	{
		num[i] = rand();
	}
	
	GET_TIME(start);
	
	Odd_even_sort(num, size);
	
	GET_TIME(stop);
	
	printf("Serial run time: %e\n", stop-start);
	return 0;
}

我测试了1E3到1E6数量级的数字排序，时间消耗如下表，均测量多次取平均值。

然后，做了一个函数图像，可以看到，基本上呈指数函数走势。其中，横轴为需要排序的数字数。纵轴为时间。当数字再大时，就会堆栈溢出了。

（2）并行部分

接下来是并行函数的分析。CUDA的并行程序分成两个部分，一个部分是被称为核的.cu文件，一个是调用核进行运算的.cpp文件。理论上其实可以将这两个代码合并，但是一般来说像这样分开比较清楚，也能体现CUDA运算的特性。具体代码如下：

//kernel.cu
#include "cuda_runtime.h"  
#include "device_launch_parameters.h"  

__global__ void swap(int* num, int size, int flag)
{
	int idx = threadIdx.x;
	int temp;
	
	if (idx < size-1 && idx % 2 == flag && num[idx] > num[idx + 1])
	{
		temp = num[idx];
		num[idx] = num[idx + 1];
		num[idx + 1] = temp;
	}

	__syncthreads();
}

extern "C" void swapKernel(int* num, int size)
{
	int* dev_num = 0;

	cudaSetDevice(0);
	cudaMalloc((void**)& dev_num, sizeof(int) * size);
	cudaMemcpy(dev_num, num, sizeof(int) * size, cudaMemcpyHostToDevice);

	for (int i = 0; i < size; i++)
	{
		swap <<<1, size>>> (dev_num, size, i%2);
	}
	

	cudaMemcpy(num, dev_num, size * sizeof(int), cudaMemcpyDeviceToHost);

	cudaFree(dev_num);
}

//test.cpp
#include <stdio.h>  
#include <stdlib.h>

#define MAXN 1000000

extern "C" void swapKernel(int* num, int size);

int main(int argc, char** argv)
{
	int num[MAXN];

	for (int i = 0; i < MAXN; i++)
	{
		num[i] = rand();
		//printf("%d ", num[i]);
	}

	swapKernel(num, MAXN);

	/*
	printf("\n");
	for (int i = 0; i < MAXN; i++)
	{
		printf("%d ", num[i]);
	}
	*/
	
	return 0;
}

结果出乎我的意料，性能差距极其之大。下图是我用Nsight Moniter对程序进行测试得出的时间。这是一个Nvidia自带的进行性能分析和监视的软件。它会自动整合进VS中。首先，验证一下程序的正确性。

可以很明显的看到，结果正确，成功实现了奇偶排序算法。接下来，需要研究一下性能如何。

这是MAXN=1E3时的情况，可以看到，只花了0.47秒，GPU的利用率也极其低，基本上大多数都是CPU运算花费的时间。

当MAXN=1E4时，情况和1E3时大同小异。到MAXN=1E5时，可以看到，CPU占比开始下降了，GPU开始真正发挥作用了。测不出使用率的原因是没能测出函数调用的次数。

之后，为了测出更高的数字，我不得不调大系统分配的堆栈，并设置MAXN=1E6时，也就是一百万个数字排序时，可以看到时间开始上升到达3秒左右，且CPU占比只到了27.2%，而GPU利用率达到了26.9%。此时，已经没必要再开大数组了。可以下定论，两者性能差距相当大。