前言：為了處理大批量資料，專案需要高並行處理，快速完成計算。開發平臺為兩臺伺服器，每臺伺服器擁有8個V100 GPU。本人目前使用工作站只有兩個GPU（RTX 3080）。直接上程式碼，註釋也比較清楚，有CUDA程式碼開發經驗的能直接讀懂，多GPU開發使用還不熟練，程式碼效率感覺不高，有哪裡可以改進的，希望有心之人可以指出，每一次交流都是一次進步！

1.平臺介紹

win10 + VS2015 + CUDA11.0 + RXT3080*2

2.程式碼

//多GPU程式設計，現在還不熟悉使用,有些函式不熟悉。
#include<iostream>
#include<cuda.h>
#include<cuda_runtime.h>
#include<device_launch_parameters.h>
//#include<omp.h>
using namespace std;
//__global__ void add(double *input, const double val);
void showGPU(int gpuNum)
{
	cudaDeviceProp prop;
	cudaGetDeviceProperties(&prop, gpuNum);
	cout << "正在呼叫第  " << gpuNum << " 個GPU，型號為：" << prop.name << endl;
}
__global__ void add(double *input, const double val);
int main() {
//1)讀取裝置數量*********************************************************
	cout << "1)裝置資訊讀取******************************************************************" << endl;
	int gpus;
	cudaGetDeviceCount(&gpus);
	cout << "該主機擁有GPU的數量為 ： " << gpus << endl;
	cudaDeviceProp prop;
	size_t freeSize;
	size_t TotalSize;
	for (int i = 0; i < gpus; i++)
	{
		cudaSetDevice(i);
		cudaGetDeviceProperties(&prop, i);
		cout << "第 "<<i<<" 個GPU的型號為："<<prop.name << endl;
		cudaMemGetInfo(&freeSize, &TotalSize);
		cout << "可利用記憶體大小 : " << freeSize / 1024 / 1024 / 1024 << " G " << endl;
		cout << "總記憶體大小 : " << TotalSize / 1024 / 1024 / 1024 << " G" << endl;
		cout << "*****************************************" << endl;
	}
	int flag_p2p = 0;
	for (int i = 0; i < gpus; i++)
	{
		for (int j = i+1; j < gpus; j++)
		{
			cudaDeviceCanAccessPeer(&flag_p2p, i, j);
			if (flag_p2p)
				cout << "gpu[" << i << "] 可以與gpu[" << j << "] 進行點對點傳輸" << endl;
			else
				cout << "gpu[" << i << "] 不可以與gpu[" << j << "] 進行點對點傳輸" << endl;
		}
	}
//2）設定GPU上的工作流 *********************************************************
	cout << "2)多GPU工作流設定******************************************************************" << endl;
	int gpu_st_num = 2;
	cudaStream_t *st1 = new cudaStream_t[gpu_st_num];
	cudaStream_t *st2 = new cudaStream_t[gpu_st_num];
	cudaSetDevice(0);
	showGPU(0);
	for (int i = 0; i < gpu_st_num; i++)
	{
		cudaStreamCreate(&st1[i]);
	}
	cudaSetDevice(1);
	showGPU(0);
	for (int i = 0; i < gpu_st_num; i++)
	{
		cudaStreamCreate(&st2[i]);
	}
// openMP 主機端分配執行緒，控制GPU。。。2020-12-17 22:22:19 失敗
//#pragma omp parallel  num_threads(gpus)
//	for (int i = 0; i < 2; i++)
//	{ 
//		cout << "This is threads : " << omp_get_thread_num() << "  Total threads is : " << omp_get_num_threads() << endl;
//	}

//3)獲取處理的資料*********************************************************
	//假設一個資料
	cout << "3)資料獲取******************************************************************" << endl;
	long long data_num = 1 <<  30;
	long long size = data_num * sizeof(double);
	cout << "處理資料量大小（位元組）為 : " << size / 1024 / 1024 << "MB" << endl;
	cout << "每個GPU處理的資料大小（位元組）為" << size / 1024 / 1024  / gpus << "MB" << endl;
	double *data = new double[data_num];
	memset(data, 0, size);

	//加入事件--測時
	cudaSetDevice(0);
	cudaEvent_t start, stop;
	cudaEventCreate(&start);
	cudaEventCreate(&stop);
	cudaEventRecord(start);

//4)在不同的GPU進行資源配置********************************************************
	cout << "4)多GPU任務分配******************************************************************" << endl;
	//4.1）gpu--0  ***********************************
	cudaSetDevice(0);
	showGPU(0);
	double *d_a_gpu0;
	cudaMalloc(&d_a_gpu0, size / 2);
	cudaMemcpyAsync(d_a_gpu0, data, size / 4, cudaMemcpyHostToDevice,st1[0]); //非同步傳輸，非同步流控制，傳輸資料前半部分
	cudaMemcpyAsync(d_a_gpu0 + data_num / 4, data + data_num / 4, size / 4, cudaMemcpyHostToDevice, st1[1]); //非同步傳輸，非同步流控制，傳輸資料前半部分
	//4.2）gpu--1  ***********************************
	cudaSetDevice(1);
	showGPU(1);
	double *d_a_gpu1;
	cudaMalloc(&d_a_gpu1, size / 2);
	cudaMemcpyAsync(d_a_gpu1, data + data_num / 2, size / 4, cudaMemcpyHostToDevice, st2[0]); //非同步傳輸，非同步流控制，傳輸資料後半部分
	cudaMemcpyAsync(d_a_gpu1 + data_num / 4, data + data_num / 2 + data_num / 4, size / 4, cudaMemcpyHostToDevice, st2[1]); //非同步傳輸，非同步流控制，傳輸資料後半部分
//5)核函式配置與呼叫************************************************************
	cout << "5)多GPU配置與呼叫******************************************************************" << endl;
	
	dim3 block(32, 8);
	dim3 grid(4, 1, 1);
	//5.1)gpu--0  ***********************************
	cudaSetDevice(0);
	showGPU(0);
	cudaStreamSynchronize(st1[0]);
	add << <grid, block, 0, st1[0] >> >(d_a_gpu0, -1);
	cudaMemcpyAsync(data, d_a_gpu0, size / 4, cudaMemcpyDeviceToHost, st1[0]);
	cudaStreamSynchronize(st1[1]);
	add << <grid, block, 0, st1[1] >> >(d_a_gpu0 + data_num / 4, -2);
	cudaMemcpyAsync(data + data_num / 4, d_a_gpu0 + data_num / 4, size / 4, cudaMemcpyDeviceToHost, st1[1]);
	//5.2)gpu--1  ***********************************
	cudaSetDevice(1);
	showGPU(1);
	add << <grid, block, 0, st2[0] >> >(d_a_gpu1, 1);
	cudaMemcpyAsync(data + data_num / 2, d_a_gpu1, size / 4, cudaMemcpyDeviceToHost,st2[0]);
	add << <grid, block, 0, st2[1] >> >(d_a_gpu1 + data_num / 4, 2);
	cudaMemcpyAsync(data + data_num / 2 + data_num / 4, d_a_gpu1 + data_num / 4, size / 4, cudaMemcpyDeviceToHost, st2[1]);
	
	
	//複製下面這一段到你想要停止計時的地方。
	cudaSetDevice(0);
	cudaDeviceSynchronize();
	cudaEventRecord(stop);
	cudaEventSynchronize(stop);
	float tm;
	cudaEventElapsedTime(&tm, start, stop);
	printf("cost time %.4f ms\n", tm);
	cudaDeviceSynchronize();
//6)抽檢資料************************************************************
	cout << "6)主機端校驗資料******************************************************************" << endl;
	for (int i = 0; i < 5; i++)
	{
		//cout << "經GPU0處理後：data[" << i << "] = " << data[i] << "\t";
		cout << "經GPU0處理，stream0 --> data[" << i << "] = " << data[i] << "\t";
		cout << "stream1 --> data[" << i << "] = " << data[i+data_num / 4] << "\t";
		cout << "經GPU1處理，stream0 --> data[" << i << "] = " << data[i + data_num / 2] << "\t";
		cout << "stream1 --> 為data[" << i << "] = " << data[i + data_num / 4 + data_num / 2] << endl;
		//cout << "經GPU1處理後：data[" << i + data_num / 2 << "] = " << data[i + data_num / 2] << endl;
	}
	
	cudaFree(d_a_gpu0);
	cudaFree(d_a_gpu1);
	free(data);
	return 0;
}

__global__ void add(double *input, const double val)
{
	int idx = threadIdx.x + blockIdx.x*blockDim.x;
	int idy = threadIdx.y + blockIdx.y*blockDim.y;
	const int row = gridDim.x * blockDim.x;
	const int col = gridDim.y * blockDim.y;
	long id = idx + idy*row;
	input[id] = input[id] + val;
}

程式碼只是簡單地應用了一個給向量中每個元素加1的十分十分簡單的操作，但是該核函式可被替換，應用與其它程式開發中，最重要的搭建了一個開發框架，用於後續參考借鑑，後續有更合理高效的會分享出來。

上述程式碼中，沒有使用openMP在主機端開啟多執行緒來控制多個GPU，只是使用了單執行緒。

我在自己工作站上試用過openMP，但是出現了問題：CUDA需要執行在x64平臺下，但是此時執行openMP（環境已經配置好）會出現錯誤，此時切換到x86平臺後，程式碼執行不發生錯誤。此處問題還未解決！！！如果有知道如何解決的，還希望能分享我一下，萬分感謝！