CUDA C++:CUDA内核不使用__shared__如何定义未知大小数组
CUDA线程私有可变长度数组解决方案
针对你的场景,有以下几种可行的处理方式,优先推荐性能最好的预分配方案:
方案1:预计算+全局内存预分配(生产环境首选)
这是图计算场景下的标准实现方案,完全避开内核内动态内存申请的开销:
- 先运行一次轻量统计内核,遍历所有节点算出每个节点的1跳邻接点数量(度),结果存在长度为
n_IR的degree_arr数组中 - 在主机端对
degree_arr做前缀和计算,得到每个节点存储邻接点的全局内存起始偏移,同时统计所有节点的总度数,按总大小开辟一块全局内存distOne_buf传入计算内核 - 计算内核中直接通过当前节点
seed对应的偏移量访问distOne_buf,读写当前节点的邻接点即可,不需要在内核中动态申请空间
如果你的邻接表本身就是提前构建好的,甚至可以直接在主机端就把distOne_buf填充完,内核直接读取使用,性能最优。
方案2:临时计算省掉存储(适合不需要留存邻接点的场景)
如果你存储邻接点只是为了计算当前节点的指标,不需要留存给后续步骤使用,完全可以边遍历1跳邻接点边做指标计算,不需要开辟任何数组存储邻接点,直接省掉存储开销。
方案3:内核内动态申请(仅适合调试验证)
CUDA内核支持线程私有堆内存申请,你可以先算出当前节点的度,再申请对应长度的数组:
__global__ void expectedForce(int* IR_vec, int* IC_vec, int n_IR) { double ExF = 0; int seed = blockDim.x * blockIdx.x + threadIdx.x+1; if(seed < n_IR) { int valRiga = IR_vec[seed]; // 先计算当前节点的1跳邻接点数量k int k = calc_degree(valRiga); // 替换为你自己的度计算逻辑 int* distOne = new int[k]; // 填充数组、计算指标逻辑 ... // 用完释放,避免内存泄漏 delete[] distOne; } }
注意这个方案性能极低,内核中的new/malloc是全局锁实现,高并发申请时会有严重的竞争开销,仅适合小数据量调试使用。
方案4:可变长度数组(兼容性差,不推荐)
部分NVIDIA GPU架构支持C语言可变长度数组(VLA)扩展,你可以先算出度再声明栈上数组:
int k = calc_degree(valRiga); int distOne[k];
该方案限制较多:线程默认栈大小通常仅1KB左右,邻接点数量稍多就会栈溢出,同时属于非标准扩展,兼容性差,不建议生产环境使用。
内容的提问来源于stack exchange,提问作者Federica De Martin
相关产品推荐
相关产品推荐

