CUDA在__device__函数中分配内存

Question

在CUDA中有没有办法在__device__函数中分配内存？我找不到这样做的任何例子.

来自手册:B.15动态全局内存分配void*malloc(size_t size); void free(void*ptr); 从全局内存中的固定大小的堆动态分配和释放内存.CUDA内核中malloc()函数至少从设备堆分配大小字节,并返回指向已分配内存的指针,如果存在的内存不足以满足请求,则返回NULL.保证返回的指针与16字节边界对齐.CUDA in-kernel 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 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和支持> 2x体系结构的卡.