如何在内核中动态分配数组？
Posted
技术标签:
【中文标题】如何在内核中动态分配数组？【英文标题】：How to dynamically allocate arrays inside a kernel? 【发布时间】：2012-11-08 22:24:54 【问题描述】：我需要在内核函数内部动态分配一些数组。我该怎么做？
我的代码是这样的：
__global__ func(float *grid_d,int n, int nn)
int i,j;
float x[n],y[nn];
//Do some really cool and heavy computations here that takes hours.
但这行不通。如果这是在主机代码中，我可以使用 malloc。 cudaMalloc 需要主机上的指针和设备上的其他指针。在内核函数内部我没有主机指针。
那么，我该怎么办？
如果分配所有数组的时间太长（几秒钟）（我需要大约 4 个大小为 n 的数组和 5 个大小为 nn 的数组），这不会有问题。因为内核可能至少会运行 20 分钟。
【问题讨论】：
您可能想阅读CUDA C programmers guide 中设备代码中有关dynamic memory allocation 的部分。此功能需要您的 GPU 具有 2.0 或更高的计算能力。 运行这个内核的配置（块、线程）是什么？n 和 nn 的典型范围是多少（对于小尺寸，您可能会将它们挤入寄存器或共享内存中）。
【参考方案1】：
仅计算能力 2.x 和更新的硬件支持动态内存分配。您可以在内核中使用 C++ new 关键字或 malloc，因此您的示例可以变为：
__global__ func(float *grid_d,int n, int nn)
int i,j;
float *x = new float[n], *y = new float[nn];
这会在具有上下文生命周期的本地内存运行时堆上分配内存，因此如果您不打算再次使用内存，请确保在内核完成运行后释放内存。您还应该注意，不能直接从主机 API 访问运行时堆内存，因此您不能将内核中分配的指针作为参数传递给 cudaMemcpy，例如。
【讨论】：
我有类似的情况，我需要动态分配数组。每个线程都必须访问这些数组以进行写入。我很困惑，如果我在内核中声明这个动态分配过程，那么如果内核的维度是（1,4），即 nThreads = 4 和 nBlocks = 1，它会创建 4 次这样的数组吗？free 在这里合适吗，还是有另一个函数可以从内核中的本地堆中释放出来？
@landau 不，你只是免费使用或删除【参考方案2】：
@talonmies 回答了您关于如何在内核中动态分配内存的问题。这旨在作为补充答案，解决 __device__ malloc() 的性能问题以及您可能要考虑的替代方案。
在内核中动态分配内存可能很诱人，因为它允许 GPU 代码看起来更像 CPU 代码。但它会严重影响性能。我写了一个独立的测试并将其包含在下面。该测试启动了大约 260 万个线程。每个线程使用从线程索引派生的一些值填充 16 个整数的全局内存，然后将这些值相加并返回总和。
测试实现了两种方法。第一种方法使用__device__ malloc()，第二种方法使用在内核运行之前分配的内存。
在我的 2.0 设备上，使用 __device__ malloc() 时内核运行时间为 1500 毫秒，使用预分配内存时运行时间为 27 毫秒。换句话说，当内存在内核中动态分配时，测试需要 56 倍 才能运行。时间包括外部循环cudaMalloc()/cudaFree()，它不是内核的一部分。如果使用相同数量的线程多次启动同一个内核，通常情况下，cudaMalloc() / cudaFree() 的成本将分摊到所有内核启动中。这使差异更大，达到 60 倍左右。
推测，我认为性能下降部分是由隐式序列化引起的。 GPU 可能必须序列化所有同时对__device__ malloc() 的调用，以便为每个调用者提供单独的内存块。
不使用__device__ malloc()的版本会在运行内核之前分配所有GPU内存。指向内存的指针被传递给内核。每个线程计算一个索引到先前分配的内存中，而不是使用__device__ malloc()。
预先分配内存的潜在问题是，如果只有一些线程需要分配内存，并且不知道是哪些线程，则有必要为所有线程分配内存。如果没有足够的内存，那么减少每个内核调用的线程数可能比使用__device__ malloc() 更有效。其他解决方法可能最终会重新实现 __device__ malloc() 在后台执行的操作，并且会看到类似的性能损失。
测试__device__ malloc()的性能：
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
const int N_ITEMS(16);
#define USE_DYNAMIC_MALLOC
__global__ void test_malloc(int* totals)
int tx(blockIdx.x * blockDim.x + threadIdx.x);
int* s(new int[N_ITEMS]);
for (int i(0); i < N_ITEMS; ++i)
s[i] = tx * i;
int total(0);
for (int i(0); i < N_ITEMS; ++i)
total += s[i];
totals[tx] = total;
delete[] s;
__global__ void test_malloc_2(int* items, int* totals)
int tx(blockIdx.x * blockDim.x + threadIdx.x);
int* s(items + tx * N_ITEMS);
for (int i(0); i < N_ITEMS; ++i)
s[i] = tx * i;
int total(0);
for (int i(0); i < N_ITEMS; ++i)
total += s[i];
totals[tx] = total;
int main()
cudaError_t cuda_status;
cudaSetDevice(0);
int blocks_per_launch(1024 * 10);
int threads_per_block(256);
int threads_per_launch(blocks_per_launch * threads_per_block);
int* totals_d;
cudaMalloc((void**)&totals_d, threads_per_launch * sizeof(int));
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaDeviceSynchronize();
cudaEventRecord(start, 0);
#ifdef USE_DYNAMIC_MALLOC
cudaDeviceSetLimit(cudaLimitMallocHeapSize, threads_per_launch * N_ITEMS * sizeof(int));
test_malloc<<<blocks_per_launch, threads_per_block>>>(totals_d);
#else
int* items_d;
cudaMalloc((void**)&items_d, threads_per_launch * sizeof(int) * N_ITEMS);
test_malloc_2<<<blocks_per_launch, threads_per_block>>>(items_d, totals_d);
cudaFree(items_d);
#endif
cuda_status = cudaDeviceSynchronize();
if (cuda_status != cudaSuccess)
printf("Error: %d\n", cuda_status);
exit(1);
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
float elapsedTime;
cudaEventElapsedTime(&elapsedTime, start, stop);
printf("Elapsed: %f\n", elapsedTime);
int* totals_h(new int[threads_per_launch]);
cuda_status = cudaMemcpy(totals_h, totals_d, threads_per_launch * sizeof(int), cudaMemcpyDeviceToHost);
if (cuda_status != cudaSuccess)
printf("Error: %d\n", cuda_status);
exit(1);
for (int i(0); i < 10; ++i)
printf("%d ", totals_h[i]);
printf("\n");
cudaFree(totals_d);
delete[] totals_h;
return cuda_status;
输出：
C:\rd\projects\test_cuda_malloc\Release>test_cuda_malloc.exe
Elapsed: 27.311169
0 120 240 360 480 600 720 840 960 1080
C:\rd\projects\test_cuda_malloc\Release>test_cuda_malloc.exe
Elapsed: 1516.711914
0 120 240 360 480 600 720 840 960 1080
【讨论】：
您应该在第二个测试中计时 cudaMalloc。否则，您将准备运行的汽车（第二次测试）与车库中停止的汽车（第一次测试）进行比较。两个内核需要相同的存储要求。 除了 pQB 反对意见：您的cudaMalloc 分配了一个大数组，这与分配 250 万个小矩阵（每个线程一个）相比。这样的过程当然会更慢，并且对 CPU 的测试表明，您报告的 60 倍减速实际上是一项不错的工作（我得到 1000 倍的减速，前提是代码没有段错误——分配器需要处理这么多的矩阵）。公平测试是：分配相同的（一个）数组，（1）每个cudaMalloc，（2）每个kernel<<<1,1>>>。我看到kernel 的分配速度变慢了~3 倍。所以这是真正的性能冲击。
@pQB：谢谢。我已经把 cudaMalloc() 排除在时间之外，假设它是不可测量的。令我惊讶的是，添加它确实引起了变化，从 60 倍变为 56 倍。我已经更新了答案并添加了关于在时间中包含 cudaMalloc() / cudaFree() 的含义的简介。
@PMarecki：测试的目的是展示使用__device__ malloc() 对性能的影响，并展示另一种完成任务的方法，许多人会考虑使用__device__ malloc()。目的不是比较单个 cudaMalloc() 和单个 __device__ malloc() 的性能。
@RogerDahl 整洁的测试！我认为重点是展示分配许多小矩阵的区别，无论是在设备上还是在主机上。但是，无论如何，使用相同数量的 malloc 调用。我认为“当然”单个 malloc 调用会比许多单个 malloc 调用更快。【参考方案3】：
如果在调用内核之前n和nn的值是已知的，那为什么不 cudaMalloc 把主机端的内存传给内核呢？
【讨论】：
因为每个内核必须拥有一个数组。 您是否同时启动多个 kenel？难道你不能分配足够的空间，每个内核只共享一部分吗？ 如果我启动，例如，1000 个内核，如果我需要 10 个大小为 n 的数组。我应该制作 10 个大小为 n*1000 的数组？并使用 threadid 和 blockid 在内核之间共享？【参考方案4】：根据@rogerdahl 帖子中的概念进行实验。假设：
4MB 内存以 64B 块分配。 该块中有 1 个 GPU 块和 32 个扭曲线程 在 P100 上运行GPU 本地的 malloc+free 调用似乎比 cudaMalloc + cudaFree 调用快得多。程序的输出：
Starting timer for cuda malloc timer
Stopping timer for cuda malloc timer
timer for cuda malloc timer took 1.169631s
Starting timer for device malloc timer
Stopping timer for device malloc timer
timer for device malloc timer took 0.029794s
我将省略 timer.h 和 timer.cpp 的代码，但这里是测试本身的代码：
#include "cuda_runtime.h"
#include <stdio.h>
#include <thrust/system/cuda/error.h>
#include "timer.h"
static void CheckCudaErrorAux (const char *, unsigned, const char *, cudaError_t);
#define CUDA_CHECK_RETURN(value) CheckCudaErrorAux(__FILE__,__LINE__, #value, value)
const int BLOCK_COUNT = 1;
const int THREADS_PER_BLOCK = 32;
const int ITERATIONS = 1 << 12;
const int ITERATIONS_PER_BLOCKTHREAD = ITERATIONS / (BLOCK_COUNT * THREADS_PER_BLOCK);
const int ARRAY_SIZE = 64;
void CheckCudaErrorAux (const char *file, unsigned line, const char *statement, cudaError_t err)
if (err == cudaSuccess)
return;
std::cerr << statement<<" returned " << cudaGetErrorString(err) << "("<<err<< ") at "<<file<<":"<<line << std::endl;
exit (1);
__global__ void mallocai()
for (int i = 0; i < ITERATIONS_PER_BLOCKTHREAD; ++i)
int * foo;
foo = (int *) malloc(sizeof(int) * ARRAY_SIZE);
free(foo);
int main()
Timer cuda_malloc_timer("cuda malloc timer");
for (int i = 0; i < ITERATIONS; ++ i)
if (i == 1) cuda_malloc_timer.start(); // let it warm up one cycle
int * foo;
cudaMalloc(&foo, sizeof(int) * ARRAY_SIZE);
cudaFree(foo);
cuda_malloc_timer.stop_and_report();
CUDA_CHECK_RETURN(cudaDeviceSynchronize());
Timer device_malloc_timer("device malloc timer");
device_malloc_timer.start();
mallocai<<<BLOCK_COUNT, THREADS_PER_BLOCK>>>();
CUDA_CHECK_RETURN(cudaDeviceSynchronize());
device_malloc_timer.stop_and_report();
如果您发现错误，请在 cmets 中 lmk，我会尽力修复它们。
然后我用更大的东西再次运行它们：
const int BLOCK_COUNT = 56;
const int THREADS_PER_BLOCK = 1024;
const int ITERATIONS = 1 << 18;
const int ITERATIONS_PER_BLOCKTHREAD = ITERATIONS / (BLOCK_COUNT * THREADS_PER_BLOCK);
const int ARRAY_SIZE = 1024;
而且 cudaMalloc 仍然慢了很多：
Starting timer for cuda malloc timer
Stopping timer for cuda malloc timer
timer for cuda malloc timer took 74.878016s
Starting timer for device malloc timer
Stopping timer for device malloc timer
timer for device malloc timer took 0.167331s
【讨论】：
另外值得注意的是，malloc + free 与new 和delete 花费的时间基本相同。【参考方案5】：
也许你应该测试一下
cudaMalloc(&foo,sizeof(int) * ARRAY_SIZE * ITERATIONS);
cudaFree(foo);
改为
for (int i = 0; i < ITERATIONS; ++ i)
if (i == 1) cuda_malloc_timer.start(); // let it warm up one cycle
int * foo;
cudaMalloc(&foo, sizeof(int) * ARRAY_SIZE);
cudaFree(foo);
【讨论】：
以上是关于如何在内核中动态分配数组？的主要内容，如果未能解决你的问题，请参考以下文章