C语言如何在内核中动态分配数组？

klh5stk1 于 11个月前发布在其他

关注(0)|答案(5)|浏览(114)

我需要在内核函数中动态地分配一些数组。我该怎么做呢？
我的代码是这样的：

__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分钟。

来源：https://stackoverflow.com/questions/13480213/how-to-dynamically-allocate-arrays-inside-a-kernel

5条答案

按热度按时间

wecizke31#

动态内存分配仅在compute capability 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。

赞(0）回复(0）举报 11个月前

bwleehnv2#

@talonmies回答了你关于如何在内核中动态分配内存的问题。这是一个补充答案，解决了__device__ malloc()的性能问题，也是你可能想考虑的替代方案。
在内核中动态分配内存可能很诱人，因为它允许GPU代码看起来更像CPU代码。但它可能会严重影响性能。我写了一个自包含的测试，并在下面包含它。测试启动了大约260万个线程。每个线程使用从线程索引派生的一些值填充16个整数的全局内存，然后将这些值相加并返回总和。
测试实现了两种方法，第一种方法使用__device__ malloc()，第二种方法使用在内核运行之前分配的内存。
在我的2.0设备上，使用__device__ malloc()时，内核运行时间为1500 ms，使用预分配内存时为27 ms。换句话说，在内核内动态分配内存时，测试运行时间延长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;
}

字符串
输出量：

(test_malloc_2)
C:\rd\projects\test_cuda_malloc\Release>test_cuda_malloc.exe
Elapsed: 27.311169
0 120 240 360 480 600 720 840 960 1080

(test_malloc)
C:\rd\projects\test_cuda_malloc\Release>test_cuda_malloc.exe
Elapsed: 1516.711914
0 120 240 360 480 600 720 840 960 1080

型

赞(0）回复(0）举报 11个月前

piok6c0g3#

如果n和nn的值在内核被调用之前就已经知道了，那么为什么不在主机端cudaMalloc内存并将设备内存指针传递给内核呢？

赞(0）回复(0）举报 11个月前

u0njafvf4#

根据@rogerdahl的帖子中的概念运行了一个实验。假设：

4 MB的内存分配在64 B块中。
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();
}

型
如果你发现了错误，请在评论中注明，我会尽力修复它们。
我又用更大的东西做了一遍：

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

型

赞(0）回复(0）举报 11个月前

jjjwad0x5#

也许你应该测试一下

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);
}

型

赞(0）回复(0）举报 11个月前

我来回答

C语言如何在内核中动态分配数组？

5条答案

相关问题

热门标签

最新问答

C语言 如何在内核中动态分配数组？

5条答案

相关问题

热门标签

最新问答

C语言如何在内核中动态分配数组？