我有以下代码：
__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动态并行编程指南，我在全局内存中分配了一个 169 _的 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 ，几乎所有时间都花在了 cudaMalloc API上 .
因为在动态编程中，人们经常需要分配内存空间来与子内核交换数据，唯一的解决方案似乎是全局内存需要花费很多时间，在我看来，动态编程的有用性的一个严重瓶颈就是数据交换，除非有办法规避全局内存分配问题 .
那么问题是：上述问题是否有任何解决方法，即从内核中分配全局内存需要花费大量时间？谢谢
SOLUTION PROPOSED IN THE COMMENTS
从父内核外部分配所需的全局内存 . 我已经验证从父内核之外分配所需的全局内存要快得多 .
1 回答
你从每个线程调用cudaMalloc，其中i <M表示你正在进行M cudaMalloc调用 .
M越大，它就越糟糕 .
相反，你可以从块的第一个线程调用一个cudaMalloc，分配M次你之前使用的大小（实际上在你的情况下你应该分配更多，所以每个块都正确对齐） . 在同步线程之后，您可以为每个子内核使用正确计算的phi_cap地址启动子内核 .
或者（如果您的特定情况允许您分配足够的内存，您可以在内核调用之间保留），您可以在内核之外分配内存并重用它 . 那会更快 . 如果M在内核调用之间变化，则可以根据最大M的需要进行分配 .