放置常量内存的最佳位置，该内存在内核启动之前已知且从未更改
我有一个整数数组，其大小在内核启动之前已知，但在编译阶段未知。大小的上限约为 10000 个 float3 元素（我猜这意味着 10000 * 3 * 4 = ~120KB）。它在编译时是未知的。
所有线程线性扫描（最多）数组中的所有元素。
I have an array of integers which size is known before the kernel launch but not during the compilation stage. The upper bound on the size is around 10000 float3 elements (I guess that means 10000 * 3 * 4 = ~120KB). It is not known at the compile time.
All threads scan linearly through (at most) all of the elements in the array.
如果你对这篇内容有疑问，欢迎到本站社区发帖提问 参与讨论，获取更多帮助，或者扫码二维码加入 Web 技术交流群。
绑定邮箱获取回复消息
由于您还没有绑定你的真实邮箱，如果其他用户或者作者回复了您的评论，将不能在第一时间通知您！
发布评论
评论（3）
您可以在运行时检查大小，然后是否适合使用 cudaMemcpyToSymbol，或者使用纹理或全局内存。这有点混乱，您必须有一些参数来告诉内核数据在哪里。一如既往，始终测试实际性能。在不同类型的内存中，不同的访问模式可能具有截然不同的速度。
另一个想法是退后一步，再次审视算法。通常有多种方法可以以不同的方式划分问题，以使常量表始终适合常量内存。
You could check the size at runtime, then if it will fit use
cudaMemcpyToSymbol, or otherwise use texture or global memory. This is slightly messy, you will have to have some parameter to tell the kernel where the data is. As always, always test actual performance. Different access patterns can have drastically different speeds in different types of memory.Another thought is to take a step back and look at the algorithm again. There are often ways of dividing the problem differently to get the constant table to always fit into constant memory.
如果扭曲中的所有线程同时访问相同的元素，那么您可能应该考虑使用常量内存，因为这不仅被缓存，而且还具有广播功能，所有线程都可以在单个周期中读取相同的地址。
If all threads in a warp access the same elements at the same time then you should probably consider using constant memory, since this is not only cached, but it also has a broadcast capability whereby all threads can read the same address in a single cycle.
您可以在编译内核后计算可用的常量内存并静态分配它。
然后，使用 cudaMemcpyToSymbol() 将数据复制到常量内存。
我认为这可能会回答您的问题，但您对恒定内存的要求超出了 GPU 的限制。
我会推荐其他方法，即使用共享内存，如果 halfwarp 中的所有线程都从同一位置读取，则可以广播数据。
You could calculate the free constant memory after compile your kernels and allocate it statically.
Then, copy your data to constant memory using cudaMemcpyToSymbol().
I think this might answer your question but your requirement for constant memory exceed the limits of the GPU.
I'll recommend other approaches, i.e. use the share memory which can broadcast data if all threads in a halfwarp read from the same location.