1

我有以下代码：

__global__ void interpolation(const double2* __restrict__ data, double2* __restrict__ result, const double* __restrict__ x, const double* __restrict__ y, const int N1, const int N2, int M)
{
    int i = threadIdx.x + blockDim.x * blockIdx.x;

    [...]        

    double phi_cap1, phi_cap2;

    if(i<M) {   

         for(int m=0; m<(2*K+1); m++) {

              [calculate phi_cap1];

              for(int n=0; n<(2*K+1); n++) {

                 [calculate phi_cap2];

                 [calculate phi_cap=phi_cap1*phi_cap2];

                 [use phi_cap];

             }
    }

}

}

我想在 Kepler K20 卡上使用动态编程来调度处理phi_cap1并phi_cap2并行处理一堆线程，以减少计算时间。K=6在我的代码中，所以我启动了一个13x13线程块。

按照 CUDA 动态并行编程指南，我在全局内存中分配与子内核交换数据所需phi_cap的169元素矩阵（由 和 的乘积形成phi_cap1）。确实，引用指南，phi_cap2

作为一般规则，传递给子内核的所有存储都应该从全局内存堆中显式分配。

然后我结束了以下代码

__global__ void interpolation(const double2* __restrict__ data, double2* __restrict__ result, const double* __restrict__ x, const double* __restrict__ y, const int N1, const int N2, int M)
{
    int i = threadIdx.x + blockDim.x * blockIdx.x;

    [...]   

    dim3 dimBlock(2*K+1,2*K+1); dim3 dimGrid(1,1);

    if(i<M) {   

    double* phi_cap; cudaMalloc((void**)&phi_cap,sizeof(double)*(2*K+1)*(2*K+1));

    child_kernel<<<dimGrid,dimBlock>>>(cc_diff1,cc_diff2,phi_cap);

    for(int m=0; m<(2*K+1); m++) {

        for(int n=0; n<(2*K+1); n++) {

                        [use phi_cap];

        }
    }

}

}

问题是第一个例程需要5ms运行，而第二个例程，即使通过评论child_kernel启动，也需要23ms，几乎所有时间都花在cudaMallocAPI 上。

由于在动态编程中经常需要分配内存空间来与子内核交换数据，而唯一的解决方案似乎是全局内存占用了这么多时间，在我看来，动态编程有用性的一个严重瓶颈是数据交换，除非有办法规避全局内存分配问题。

那么问题是：对于上述问题是否有任何解决方法，即从内核中分配全局内存时需要花费这么多时间？. 谢谢

评论中提出的解决方案

从父内核外部分配所需的全局内存。我已经验证从父内核外部分配所需的全局内存要快得多。

4

1 回答 1

4

您正在从 i < M 的每个线程调用 cudaMalloc，这意味着您正在进行 M cudaMalloc 调用。

M越大，情况就越糟。

相反，您可以从分配 M 倍于您之前使用的大小的块的第一个线程进行单个 cudaMalloc 调用（实际上在您的情况下您应该分配更多，因此每个块都正确对齐）。之后同步线程，您可以使用为每个子内核正确计算的 phi_cap 地址启动子内核。

或者（如果您的特定情况允许您分配足够的内存以在内核调用之间保留），您可以在内核外部分配内存并重用它。那会快很多。如果 M 在内核调用之间变化，您可以为最大的 M 分配尽可能多的内存。

于 2013-02-15T12:31:39.743 回答