VS中将cuda项目编译为DLL并调用

本篇来介绍一下如何将cuda工程编译为DLL， 并在其他工程中调用。
配置环境 ： VS2022+CUDA12.2， 工程目录如下

普通C++工程

首先在VS中创建一个空项目或者控制台程序MainTest 作为主程序。

其次，再创建一个空项目或者控制台程序作为生成DLL的工程（CPlusPlusProject）。
在工程中写一个简单的打印helloworld函数
Demo1.h

#pragma once
#include <iostream>
__declspec(dllexport) void PrintfHello();

Demo1.cpp

#include "Demo1.h"


void PrintfHello() {
	for (int i = 0; i < 10; i++) {
		printf("hello world Cpu\n");
	}
}

注： __declspec(dllexport) 含义是我要把这个函数、类或变量导出到 DLL 中，让其他程序可以通过这个 DLL 使用它。

step1:

修改项目属性->改为动态库.DLL。 其中，输出目录可以更改生成的DLL、Lib文件的存储路径。

对CPlusPlusProject工程进行编译， 会得到如下文件（路径大概率在MainTest工程下\x64\Release文件夹中。 ）

step2:

在MainTest中进行配置， 和配置其他第三方库一样。
包含目录：
库目录：
附加依赖项：


在MainTest添加主函数进行调用
MainTest.cpp

#include <iostream>
#include "Demo1.h"
//#include "ParalleReduction.h"

int main() {

	PrintfHello();
	//paralleReduction();
}

结果如下。

编译CUDA程序

首先， 右键解决方案->添加->新建项目->CUDA项目（CudaRuntime）， 添加代码文件
ParalleReduction.h

#pragma once
#include <iostream>

extern "C"
__declspec(dllexport) void paralleReduction();

可以看出， 相比与普通c++程序， 多了extern “C”， 这是为了告诉编译器请用 C 语言的方式（不带名字修饰）来导出这个函数， 防止因编译器的不同导致的名称更改。
为了偷懒使用前面cuda并行规约的代码进行测试， 代码如下。
ParallelReduction.cu

//#define Time

//  并行规约
#include <iostream>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <chrono>
#include "ParalleReduction.h"


// 每个内核只计算一个block中的和
__global__ void reduceNeighbored(int *g_idata, int *g_odata, unsigned int n) {

	int idx = threadIdx.x;
	unsigned int index = blockIdx.x * blockDim.x + threadIdx.x;
	int* idata = g_idata + blockIdx.x * blockDim.x;
	if (index > n)
	{
		return; 
	}

	for (int stride = 1; stride < blockDim.x; stride *= 2) {
		if (idx % (2 * stride) == 0) {
			idata[idx] += idata[idx + stride];
		}
		// 等待线程同步
		__syncthreads();
	}

	if (idx == 0){
		g_odata[blockIdx.x] = idata[0];
	}
}

// 相邻配对方法
__global__ void reduceNeighboredLess(int* g_idata, int* g_odata, unsigned int n) {   // 避免线程分化版本

	unsigned int tid = threadIdx.x;
	unsigned int index = blockIdx.x * blockDim.x + threadIdx.x;
	// get data
	int* idata = g_idata + blockIdx.x * blockDim.x;   // 对应的数据

	if (index > n) {
		return;
	}

	for (int stride = 1; stride < blockDim.x; stride *= 2) {
		int sumIndex = tid * 2 * stride;
		if (sumIndex < blockDim.x) {
			idata[sumIndex] += idata[sumIndex + stride];
		}
		// 等待线程同步
		__syncthreads();

	}
	if (tid == 0) {
		g_odata[blockIdx.x] = idata[0];
	}

}


// 交错规约
__global__ void reduceInterleaved(int* g_idata, int* g_odata, unsigned int n) {
	unsigned int tid = threadIdx.x;

	unsigned int index = blockIdx.x * blockDim.x + threadIdx.x;

	int* idata = g_idata + blockIdx.x * blockDim.x;

	if (index >= n) {
		return;
	}
	int dataSize = blockDim.x;
	for (int stride = blockDim.x / 2; stride > 0; stride /= 2) {
		if (tid < stride) {
			idata[tid] += idata[tid + stride];
		}
		// 等待线程同步
		__syncthreads();
	}

	if (tid == 0) {
		g_odata[blockIdx.x] = idata[0];
	}
}

// 循环展开
__global__ void reduceUnrolling2(int* g_idata, int* g_odata, unsigned int n) {   // 将相邻两个block数据合并成一组数据， 对合并后的数据进行求和计算

	unsigned int tid = threadIdx.x;
	unsigned int idx = blockIdx.x * blockDim.x * 2 + threadIdx.x;     // 全局索引

	int* idata = g_idata + blockIdx.x * blockDim.x * 2;   // 每次归并两个block   blockId.x 取值范围从0-n

	if (idx + blockDim.x < n)                // block0 与 block1 相同的thread 对应的数据   全局数据
		g_idata[idx] += g_idata[idx + blockDim.x];

	__syncthreads();

	//
	for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
		if (tid < stride) {
			idata[tid] += idata[tid + stride];
		}
		__syncthreads();
	}
	
	if (tid == 0) {
		g_odata[blockIdx.x] = idata[0];
	}
}

// reduceUnrolling8
__global__ void reduceUnrolling8(int* g_idata, int* g_odata, unsigned int n) {
	unsigned int tid = threadIdx.x;

	unsigned int index = blockIdx.x * blockDim.x * 8 + threadIdx.x;

	int* idata = g_idata + blockIdx.x * blockDim.x * 8;

	if (index + 7*blockDim.x < n) {   // 越界条件检查
		int a1 = g_idata[index];
		int a2 = g_idata[index + blockDim.x];
		int a3 = g_idata[index + blockDim.x * 2];
		int a4 = g_idata[index + blockDim.x * 3];
		int a5 = g_idata[index + blockDim.x * 4];
		int a6 = g_idata[index + blockDim.x * 5];
		int a7 = g_idata[index + blockDim.x * 6];
		int a8 = g_idata[index + blockDim.x * 7];
		g_idata[index] = a1 + a2 + a3 + a4 + a5 + a6 + a7 + a8;
	}
	__syncthreads();
	
	for (int stride = blockDim.x / 2; stride > 32; stride >>= 1) {
		if (tid < stride) {
			idata[tid] += idata[tid + stride];
		}
		__syncthreads();
	}
	if (tid < 32) {
		volatile int* vmem = idata;
		vmem[tid] += vmem[tid + 32];
		vmem[tid] += vmem[tid + 16];
		vmem[tid] += vmem[tid + 8];
		vmem[tid] += vmem[tid + 4];
		vmem[tid] += vmem[tid + 2];
		vmem[tid] += vmem[tid + 1];
	}
	if (tid == 0) {
		g_odata[blockIdx.x] = idata[0];
	}
}



__host__ void GetSumCpu(int* idata, int& odata, const int n) {

	for (int i = 0; i < n; i++) {
		odata += idata[i];
	}
}


void paralleReduction() {
	
	// 选择设备， 打印设备信息
	int dev = 0;
	cudaDeviceProp deviceProp;
	cudaGetDeviceProperties(&deviceProp, dev);
	printf("device %d: %s \n", dev, deviceProp.name);

	// 申请host 端的memory
	int size = 1 << 24;
	int nByte = size * sizeof(int);
	dim3 block(512, 1);
	dim3 grid((size + block.x -1) / block.x, 1);

	int* idata_h = (int*)malloc(nByte);
	int* odata_h = (int*)malloc(grid.x * sizeof(int));
	int* temp = (int*)malloc(nByte);

	// 初始化数组
	for (int i = 0; i < size; i++) {
		idata_h[i] = (int)(rand() & 0xFF);
	}
	memcpy(temp, idata_h, nByte);


#ifdef Time
	auto start = std::chrono::high_resolution_clock::now();
#endif // TIme
	int result_cpu = 0;
	GetSumCpu(temp, result_cpu, size);
	printf("CPU result: %d \n", result_cpu);
#ifdef Time
	auto end = std::chrono::high_resolution_clock::now();
	std::chrono::duration<double, std::milli> elapsed = end - start;
	printf("CPU use Time: %f \n", elapsed);
#endif // Time

	// cuda memory
	int* idata_d = nullptr;
	int* odata_d = nullptr;
	cudaMalloc((int **)&idata_d, nByte);
	cudaMalloc((int **)&odata_d, grid.x *sizeof(int));

	// 内存拷贝 -- 耗时
	cudaMemcpy(idata_d, idata_h, nByte, cudaMemcpyHostToDevice);
#ifdef Time
	float milliseconds = 0;
	cudaEvent_t k_start, k_stop;
	cudaEventCreate(&k_start);
	cudaEventCreate(&k_stop);
	cudaEventRecord(k_start);
#endif // Time
	
	reduceNeighbored<< <grid, block>> >(idata_d, odata_d, size);
	cudaDeviceSynchronize();
	
#ifdef Time
	cudaEventRecord(k_stop);
	cudaEventSynchronize(k_stop);  // 等待事件完成
	cudaEventElapsedTime(&milliseconds, k_start, k_stop);  // 计算耗时（毫秒）
	printf("GPU use Time %.5f ms\n", milliseconds);

	cudaEventDestroy(k_start);
	cudaEventDestroy(k_stop);
#endif // Time
	cudaMemcpy(odata_h, odata_d, grid.x * sizeof(int), cudaMemcpyDeviceToHost);
	int gpu_sum = 0;
	for (int i = 0; i < grid.x; i++) {
		gpu_sum += odata_h[i];
	}
	printf("GPU result %d \n", gpu_sum);
	


	//  改进后的并行规约算法

	cudaMemcpy(idata_d, idata_h, nByte, cudaMemcpyHostToDevice);
#ifdef Time
	float milliseconds_L = 0;
	cudaEvent_t kl_start, kl_stop;
	cudaEventCreate(&kl_start);
	cudaEventCreate(&kl_stop);
	cudaEventRecord(kl_start);
#endif // Time
	reduceNeighboredLess << <grid, block >> > (idata_d, odata_d, size);
	cudaDeviceSynchronize();

#ifdef Time
	cudaEventRecord(kl_stop);
	cudaEventSynchronize(kl_stop);  // 等待事件完成
	cudaEventElapsedTime(&milliseconds_L, kl_start, kl_stop);  // 计算耗时（毫秒）
	printf("GPU use Time reduceNeighboredLess %.5f ms\n", milliseconds_L);

	cudaEventDestroy(kl_start);  // 销毁事件
	cudaEventDestroy(kl_stop);
#endif // Time
	cudaMemcpy(odata_h, odata_d, grid.x * sizeof(int), cudaMemcpyDeviceToHost);
	gpu_sum = 0;
	for (int i = 0; i < grid.x; i++) {
		gpu_sum += odata_h[i];
	}
	printf("GPU result reduceNeighboredLess %d \n", gpu_sum);


	// 交错规约算法

	cudaMemcpy(idata_d, idata_h, nByte, cudaMemcpyHostToDevice);
#ifdef Time
	float milliseconds_2 = 0;
	cudaEvent_t k2_start, k2_stop;
	cudaEventCreate(&k2_start);
	cudaEventCreate(&k2_stop);
	cudaEventRecord(k2_start);
#endif // Time

	reduceInterleaved<<<grid, block>>>(idata_d, odata_d, size);
	cudaDeviceSynchronize();

#ifdef Time
	cudaEventRecord(k2_stop);
	cudaEventSynchronize(k2_stop);  // 等待事件完成
	cudaEventElapsedTime(&milliseconds_2, k2_start, k2_stop);  // 计算耗时（毫秒）
	printf("GPU use Time reduceInterleaved %.5f ms\n", milliseconds_2);

	cudaEventDestroy(k2_start);  // 销毁事件
	cudaEventDestroy(k2_stop);
#endif // !Time

	cudaMemcpy(odata_h, odata_d, grid.x * sizeof(int), cudaMemcpyDeviceToHost);
	gpu_sum = 0;
	for (int i = 0; i < grid.x; i++) {
		gpu_sum += odata_h[i];
	}
	printf("GPU result reduceInterleaved %d \n", gpu_sum);


	// 展开规约算法 1 每次规约两个block
	// 可以降低耗时的原因：一个线程有更多的独立内存加载、存储操作会产生更好的性能。GPU会将多个内存操作请求合并成一个进行下发
	// 更高效的利用带宽

	cudaMemcpy(idata_d, idata_h, nByte, cudaMemcpyHostToDevice);
	reduceUnrolling2 << <grid.x / 2, block >> > (idata_d, odata_d, size);
	cudaMemcpy(odata_h, odata_d, grid.x / 2 * sizeof(int), cudaMemcpyDeviceToHost);
	gpu_sum = 0;
	for (int i = 0; i < grid.x / 2; i++) {
		gpu_sum += odata_h[i];
	}
	printf("GPU result reduceUnrolling2 %d \n", gpu_sum);

	free(idata_h);
	free(odata_h);
	free(temp);

	cudaFree(idata_d);
	cudaFree(odata_d);

}

配置CUDA工程属性

step1:

1. 右键CudaRuntime工程->生成依赖项->自定义生成->选中已安装的cuda版本 , 点击应用.
   

step2:

2.将项目类型更改为DLL

step3:

参考普通c++程序DLL调用配置的库目录、包含目录、附加依赖项的方式， 将cuda工程的配置 添加进去。

编译成功后，运行主函数

#include <iostream>
#include "Demo1.h"
#include "ParalleReduction.h"

int main() {

	PrintfHello();
	paralleReduction();
}

结果如下：