如何在多个 GPU 上同时执行 cufftXt 和 CUDA 内核

以下是我按照工具包中的simpleCUFFT_2d_MGPU代码示例进行操作的方法。我不确定它是否完全正确。在 2 个 GPU 上比仅使用 1 个 GPU 慢 50%。我在Tesla K40 GPU上测试了这段代码(而不是使用R2C和C2R FFT的另一个代码(。

#include <iostream>
#define _USE_MATH_DEFINES
#include <math.h>
#include <ctime>
#include <fstream>
#include <sstream>
#include <cstdlib>
#include <string>
#include <stdlib.h>
#include <stdio.h>
#include <cuda_runtime.h>
#include <cufft.h>
#include <cufftXt.h>
using namespace std;
__global__ void Cube (cufftComplex *data, cufftComplex *data3, int N, int n, int nGPUs) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;   
    if (i<n){
        data3[i].x = pow(data[i].x, 3.0f);  
    data3[i].y = 0;
    }
    __syncthreads();
}
__global__ void Normalize (cufftComplex *data, int N, int n, int nGPUs){
    int i = blockIdx.x * blockDim.x + threadIdx.x;   
    if (i<n){
    data[i].x /= n;
    }
    __syncthreads();
}
int main (int argc, char **argv) {
    int x, y;
    int N = 8192;
    int n = N*N;
        //int cplx_size = N * (N/2 + 1);
        //int real_size = 2 * cplx_size;
    int mem_size = sizeof(cufftComplex)*n;
    int maxThreads=(N>1024)?1024:N;
    int threadsPerBlock = maxThreads;
    int numBlocks = (n)/threadsPerBlock;
    cout <<"numBlocks " <<numBlocks <<endl;
    cufftComplex *h_data; 
    h_data = (cufftComplex*)malloc(mem_size);
    cufftComplex *h_data3 = (cufftComplex*)malloc(mem_size);
cout <<"host data allocated" <<endl;
    int index;
    float lambda = N*.1;
    for(y=0; y<N; y++){
    for(x=0; x<N; x++){
        //cout <<x <<" " <<y <<endl;
        index = x + y*N;
        h_data[index].x = cos(2*M_PI*(x+y)/lambda);
        h_data[index].y = 0;
    }
    }
cout <<"host data values set" <<endl;
    cufftResult res;
    int  device;
    int nGPUs;
    cudaGetDeviceCount(&nGPUs);
    cout <<nGPUs <<" CUDA devices" <<endl;
    size_t total_mem, free_mem;
    for(int i=0; i<nGPUs; i++){
        cudaMemGetInfo(&free_mem, &total_mem);
        cout <<"GPU" <<i <<" used memory " <<(total_mem-free_mem)/pow(10,9);
    }
    int whichGPUs[nGPUs];
    for(int i=0; i<nGPUs; i++){
        whichGPUs[i]=i;
    }
cout <<"whichgpus set" <<endl;
        size_t* worksize;
        worksize =(size_t*)malloc(sizeof(size_t) * nGPUs);
cout <<"worksize set" <<endl;
    cufftHandle plan_complex;
    res = cufftCreate(&plan_complex);
        if (res != CUFFT_SUCCESS){cout <<"create plan failed" <<endl;}
    res = cufftXtSetGPUs(plan_complex, nGPUs, whichGPUs);
        if (res != CUFFT_SUCCESS){cout <<"setgpus forward failed" <<endl;}
cout <<"set gpus" <<endl;
    res = cufftMakePlan2d(plan_complex, N, N, CUFFT_C2C, worksize);
        if (res != CUFFT_SUCCESS){cout <<"make plan forward failed" <<endl;}
cout <<"plan created" <<endl;
    cudaLibXtDesc *d_data; 
    cudaLibXtDesc *d_data3;
    res = cufftXtMalloc(plan_complex, (cudaLibXtDesc **)&d_data, CUFFT_XT_FORMAT_INPLACE);
        if (res != CUFFT_SUCCESS){cout <<"data malloc failed" <<endl;}
    res = cufftXtMalloc(plan_complex, (cudaLibXtDesc **)&d_data3, CUFFT_XT_FORMAT_INPLACE);
        if (res != CUFFT_SUCCESS){cout <<"data3 malloc failed" <<endl;}
cout <<"xtmalloc done" <<endl;
    res = cufftXtMemcpy (plan_complex, d_data, h_data, CUFFT_COPY_HOST_TO_DEVICE);
        if (res != CUFFT_SUCCESS){cout <<"memcpy to device failed" <<endl;}
cout <<"memcpy h to d" <<endl;
int tmax = 10000;
int start = time(0);
for(int tau=0; tau<tmax; tau++){
    res = cufftXtExecDescriptorC2C(plan_complex, d_data, d_data, CUFFT_FORWARD);
        if (res != CUFFT_SUCCESS){cout <<"cufftXtExec failed" <<endl; return 0;}
    res = cufftXtExecDescriptorC2C(plan_complex, d_data, d_data, CUFFT_INVERSE);
        if (res != CUFFT_SUCCESS){cout <<"cufftXtExec failed" <<endl; return 0;}
    for(int i=0; i<nGPUs; i++){
        device = d_data->descriptor->GPUs[i];
        cudaSetDevice(device);
        Normalize <<<numBlocks, threadsPerBlock>>> ((cufftComplex*) d_data->descriptor->data[i], N, n, nGPUs);
    }
    cudaDeviceSynchronize();
}
int stop = time(0);
cout <<tmax <<" timesteps" <<endl <<(stop-start) <<" seconds"<<endl;
/*
    for(int i=0; i<nGPUs; i++){
        device = d_data->descriptor->GPUs[i];
        cudaSetDevice(device);
        Cube <<<numBlocks, threadsPerBlock>>> ((cufftComplex*) d_data->descriptor->data[i], (cufftComplex*) d_data3->descriptor->data[i], N, real_size);
    }
*/
/*
    cudaDeviceSynchronize();
    res = cufftXtMemcpy (plan_complex, h_data, d_data, CUFFT_COPY_DEVICE_TO_HOST);
        if (res != CUFFT_SUCCESS){cout <<"memcpy to host failed" <<endl;}
cout <<"memcpy d to h" <<endl;
    ofstream fout;
    ostringstream outstr;
    outstr.precision(4);
    outstr <<time(0) <<".dat";
    string filename=outstr.str();
    fout.open(filename.c_str());
    fout.precision(4);
    for (int i = 0; i < n; i++) {
        x = (i % (N));
            y = (i /(N))%N; 
            fout <<x <<" " <<y <<" " <<h_data[i].x <<endl;
    }
    fout.close();  
*/
    //clean up
    res = cufftXtFree(d_data);
        if (res != CUFFT_SUCCESS){cout <<"free data failed" <<endl;}
    res = cufftXtFree(d_data3);
        if (res != CUFFT_SUCCESS){cout <<"free data3 failed" <<endl;}
    cufftDestroy(plan_complex);
    return 0;
}