【问题标题】:Lookup table using Cuda-C使用 Cuda-C 查找表
【发布时间】:2013-07-15 02:46:54
【问题描述】:

我已经使用算法方法为this post 找到了解决方案。我也很想尝试帖子中的一个 cmets 中建议的查找表方法。我对 CUDA C 相当陌生,并试图寻找有关如何做到这一点的示例/信息。我将值存储在下表中。我知道我需要关联每个线程来提取 4 个值中的每一个。这些值分别对应于每个线程的索引 SubBlkIdxA、SubBlkIdxB、BlkIdxA 和 BlkIdxB。一旦从表中读取它们,它们就会被传递给一个函数来计算一些东西。

我知道如果我说 m_aIdx[3][0],它将进入 {3,0,0,1,},进入表格并读取第一个条目'3'。为了读取这个位置的每个条目到上面提到的索引,我这样想:

我的桌子是这样的:

static __constant__ int16 m_aIdx[64][4] =
{
    {0,1,0,0,},
    {2,3,0,0,},
    {1,0,0,1,},
    {3,0,0,1,},
    {1,2,0,1,},
    {3,2,0,1,},
    and so on ... upto 64 entries
}

函数如下:

static __device__ void func()
{
    SubBlkIdxA = m_aIdx[3][0];
    SubBlkIdxB = m_aIdx[3][1];
    BlkIdxA = m_aIdx[3][2];
    BlkIdxB = m_aIdx[3][3];

    func1(SubBlkIdxA, SubBlkIdxB, BlkIdxA, BlkIdxB);
}

内核执行速度也是我关心的问题。那么,很想知道这种方法是否是一种好的做法(生成索引的有效方法)?

【问题讨论】:

  • 从另一篇文章看来,您有一个根据“算法”方法生成索引的代码。上面你有一个“非算法”方法的代码草案(似乎你只需要通过线程/块索引来处理m_aIdx 矩阵)。为什么不比较两种解决方案的相对速度?

标签: c cuda gpgpu nvidia lookup-tables


【解决方案1】:

两者都应该很快。在您的“算法”方法中,您根据存储在寄存器中的数据计算索引,这将非常快。在这种方法中,您正在对 512 字节的常量内存进行相对良好的合并内存访问,这也非常快。 (即使合并得不好,它也会很快被缓存)。

我关心的是如何在 func1 中使用这些索引。如果关于这些索引的语句可能会导致一些糟糕的分歧,并且使用这些索引进行内存访问可能会导致一些不良合并传输。

要考虑的一件事是在相同的子块中保持连续的 tid。如果它们是基于每个子块的,那么这样做将导致更清晰的内存传输。

附:我不太确定您的子块是如何构造的,因为我没有掌握您的索引模式,也不明白为什么您要在块内创建子块而不是使用更小的块。

【讨论】:

  • 也许对于 Kepler 架构,他也可以使用只读纹理缓存而不是常量内存?
  • 谢谢大家。我确实比较了速度,两种方法似乎给出了几乎相同的速度数字,但没有显着不同。
  • 在索引模式上,我需要区分子块边缘之间的边界以及块边缘之间的边界,因为我正在计算沿这些边缘的边界强度值。这就是处理数据的方式。
猜你喜欢
  • 2013-06-09
  • 2021-08-29
  • 2015-09-26
  • 2019-05-31
  • 2014-06-18
  • 2021-05-25
  • 1970-01-01
  • 1970-01-01
  • 1970-01-01
相关资源
最近更新 更多