CUDA在__device__函数中分配内存

Question

CUDA中是否有一种方法可以在设备端功能中动态分配内存？我找不到执行此操作的任何示例。

摘自CUDA C编程手册：

B.15动态全局内存分配

void* malloc(size_t size); 
void free(void* ptr); 

从全局内存中的固定大小的堆中动态分配和释放内存。

CUDA内核malloc()函数至少从设备堆中分配大小字节，并返回指向分配的内存的指针；如果内存不足，无法满足请求，则返回NULL。保证返回的指针与16字节边界对齐。

CUDA内核free()函数将分配由ptr指向的内存，该内存必须由先前对malloc()的调用返回。如果ptr为NULL，则对free（）的调用将被忽略。重复调用具有相同ptr的free（）具有未定义的行为。

由给定CUDA线程通过malloc()分配的内存在CUDA上下文的生存期内一直保持分配，或者直到调用free()显式释放它为止。它可以被其他任何CUDA线程使用，即使随后的内核启动也可以使用。任何CUDA线程都可以释放另一个线程分配的内存，但应注意确保同一指针不会被释放超过一次。

Answer 1

根据http://developer.download.nvidia.com/compute/cuda/3_2_prod/toolkit/docs/CUDA_C_Programming_Guide.pdf，您应该能够在设备函数中使用malloc（）和free（）。

第122页

B.15动态全局内存分配void * malloc（size_t大小）;无空（void * ptr）;从全局内存中的固定大小的堆中动态分配和释放内存。

手册中给出的示例。

__global__ void mallocTest()
{
    char* ptr = (char*)malloc(123);
    printf(“Thread %d got pointer: %p\n”, threadIdx.x, ptr);
    free(ptr);
}

void main()
{
    // Set a heap size of 128 megabytes. Note that this must
    // be done before any kernel is launched.
    cudaThreadSetLimit(cudaLimitMallocHeapSize, 128*1024*1024);
    mallocTest<<<1, 5>>>();
    cudaThreadSynchronize();
}

您需要编译器参数-arch = sm_20和支持> 2x体系结构的卡。