在__device__函数中分配CUDA内存

Question

在CUDA中，是否有一种方式可以在设备端函数中动态分配内存？我找不到任何这样的示例。

根据CUDA C编程手册:

B.15 动态全局内存分配

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

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

CUDA内核中的malloc()函数至少从设备堆中分配size字节，并返回指向已分配内存的指针，如果没有足够的内存来满足请求，则返回NULL。返回的指针保证对齐到16字节边界。

CUDA内核中的free()函数释放ptr指向的内存，其中ptr必须由先前调用malloc()返回。如果ptr为NULL，则忽略对free()的调用。重复使用相同的指针调用free()具有未定义的行为。

通过malloc()分配的内存在CUDA上下文的生命周期内始终保持分配状态，或直到显式地通过调用free()释放。它可以被任何其他CUDA线程使用，即使是从后续的内核启动。任何CUDA线程都可以释放另一个线程分配的内存，但应注意确保不要多次释放同一指针。

- SparcU

---

是的，我明白这是一个有些奇特的要求，但我正在移植现有的代码库。 - SparcU

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 size); void free(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，并且拥有一张支持大于2倍架构的显卡。

- Nate

---

嗨@Nate，当我在__global__函数中使用malloc和free时，它会给出编译错误，说无法从设备调用主机函数malloc和free。我是否缺少一些头文件？您知道如何检查GPU支持的体系结构吗？谢谢！ - Charles Chow