在PyCUDA的SourceModule中创建数组触发非法内存访问该如何解决？

是的，你遇到的非法内存访问错误核心就是设备端malloc的使用问题，叠加了线程并发分配的内存超额问题：

• CUDA kernel中调用的malloc从设备堆分配内存，默认设备堆大小仅为8MB。你单线程单次分配的内存总大小为 50 * 30000 * 4Byte = 6MB，已经接近默认上限。
• 你启动了2个线程执行kernel，两个线程同时申请内存总需求超过12MB，超出默认设备堆容量，导致malloc返回NULL指针，后续对NULL指针的地址写入操作触发非法内存访问。
• 额外注意：你当前的代码没有显式初始化PyCUDA上下文，部分环境下也会导致未知运行错误。

方案1：临时调大设备堆容量（适合必须在设备端动态分配的场景） #

在初始化CUDA上下文后显式设置更大的设备堆大小，同时增加malloc返回值校验，避免空指针访问，修正后的代码如下：

import pycuda.driver as drv
from pycuda.compiler import SourceModule

# 显式初始化PyCUDA
drv.init()
device = drv.Device(0)
ctx = device.make_context()
# 调大设备堆到32MB
ctx.set_limit(drv.limit.MALLOC_HEAP_SIZE, 32 * 1024 * 1024)

target = SourceModule(r"""
__device__ void Array_test()
{
    float ** PH = (float **)malloc(50 * sizeof(float *));
    // 校验一级指针分配结果
    if (!PH) return;
    for(int i = 0; i < 50; i++) {
        PH[i]=(float*)malloc(30000*sizeof(float));
        // 校验二级指针分配结果
        if (!PH[i]) {
            // 分配失败释放已申请的内存
            for(int j = 0; j < i; j++) free(PH[j]);
            free(PH);
            return;
        }
    }
    PH[49][29999] = 1;

    for(int i = 0; i < 50; i++)
        free(PH[i]);
    free(PH);  
}

void __global__ f()
{
    Array_test();
}
""")

if __name__ == "__main__":
    start = drv.Event()
    end = drv.Event()
    start.record()
    test = target.get_function("f")
    test(grid=(1,), block=(2, 1, 1))
    end.record()
    end.synchronize()
    secs = start.time_till(end) * 1e-3
    print(secs)
    # 销毁上下文
    ctx.pop()

方案2：Host端预分配设备内存（更推荐，性能更高、稳定性更好） #

设备端malloc本身开销极高，绝大多数场景下推荐在Host端完成所有设备内存分配，将指针传入kernel使用：

• 二维数组可以拍平为一维数组，Host端分配50 * 30000大小的float设备内存，kernel内通过PH[i * 30000 + j]的方式访问，无需嵌套指针。
• 完全规避设备堆不足的问题，内存分配开销比设备端malloc低两个数量级以上。

内容的提问来源于stack exchange，提问作者akio_hibi