cuda - 随机访问时共享内存中的预期存储体冲突数

标签 cuda gpu-shared-memory bank-conflict

A 为共享内存中正确对齐的 32 位整数数组。

如果单个 warp 尝试随机获取 A 的元素,则银行冲突的预期数量是多少?

换句话说:

__shared__ int A[N];          //N is some big constant integer
...
int v = A[ random(0..N-1) ];  // <-- expected number of bank conflicts here?

请假设特斯拉或费米架构。我不想详细讨论 Kepler 的 32 位与 64 位组配置。另外,为了简单起见,我们假设所有随机数都不同(因此没有广播机制)。

我的直觉建议是 4 到 6 之间的数字,但我想找到它的一些数学评估。

<小时/>

我相信这个问题可以从 CUDA 中抽象出来并呈现为一个数学问题。我将其作为生日悖论的扩展进行搜索,但我在那里发现了非常可怕的公式,并且没有找到最终的公式。我希望有一个更简单的方法......

最佳答案

在数学中,这被认为是“箱子里的球”问题 - 32 个球被随机放入 32 个箱子中。您可以枚举可能的模式并计算它们的概率以确定分布。简单的方法是行不通的,因为模式的数量很大:(63!)/(32!)(31!)“几乎”是五亿。

如果您递归地构建解决方案并使用条件概率,则可以解决这个问题。

查找 Charles J. Corrado 撰写的题为“多项式/狄利克雷和多元超几何频率的最大值、最小值和范围的精确分布”的论文。

接下来,我们从最左边的桶开始,计算可能落入其中的每个球的概率。然后,我们向右移动一个球,并在给定已使用的球和桶的数量的情况下确定该桶中可能存在的每个球数量的条件概率。

对 VBA 代码表示歉意,但当我有动力回答时,VBA 是我唯一可用的:)。

Function nCr#(ByVal n#, ByVal r#)
    Static combin#()
    Static size#
    Dim i#, j#

    If n = r Then
        nCr = 1
        Exit Function
    End If

    If n > size Then
        ReDim combin(0 To n, 0 To n)
        combin(0, 0) = 1
        For i = 1 To n
            combin(i, 0) = 1
            For j = 1 To i
                combin(i, j) = combin(i - 1, j - 1) + combin(i - 1, j)
            Next
        Next
        size = n
    End If

    nCr = combin(n, r)
End Function

Function p_binom#(n#, r#, p#)
    p_binom = nCr(n, r) * p ^ r * (1 - p) ^ (n - r)
End Function

Function p_next_bucket_balls#(balls#, balls_used#, total_balls#, _
  bucket#, total_buckets#, bucket_capacity#)

    If balls > bucket_capacity Then
        p_next_bucket_balls = 0
    Else
        p_next_bucket_balls = p_binom(total_balls - balls_used, balls, 1 / (total_buckets - bucket + 1))
    End If
End Function

Function p_capped_buckets#(n#, cap#)
    Dim p_prior, p_update
    Dim bucket#, balls#, prior_balls#

    ReDim p_prior(0 To n)
    ReDim p_update(0 To n)

    p_prior(0) = 1

    For bucket = 1 To n
        For balls = 0 To n
            p_update(balls) = 0
            For prior_balls = 0 To balls
                p_update(balls) = p_update(balls) + p_prior(prior_balls) * _
                   p_next_bucket_balls(balls - prior_balls, prior_balls, n, bucket, n, cap)
            Next
        Next
        p_prior = p_update
    Next

    p_capped_buckets = p_update(n)
End Function

Function expected_max_buckets#(n#)
    Dim cap#

    For cap = 0 To n
        expected_max_buckets = expected_max_buckets + (1 - p_capped_buckets(n, cap))
    Next
End Function

Sub test32()
    Dim p_cumm#(0 To 32)
    Dim cap#

    For cap# = 0 To 32
        p_cumm(cap) = p_capped_buckets(32, cap)
    Next

    For cap = 1 To 32
        Debug.Print "    ", cap, Format(p_cumm(cap) - p_cumm(cap - 1), "0.000000")
    Next
End Sub

对于 32 个球和桶,我预计桶中的最大球数约为 3.532941。

与艾哈迈德的输出进行比较:

           1            0.000000
           2            0.029273
           3            0.516311
           4            0.361736
           5            0.079307
           6            0.011800
           7            0.001417
           8            0.000143
           9            0.000012
           10           0.000001
           11           0.000000
           12           0.000000
           13           0.000000
           14           0.000000
           15           0.000000
           16           0.000000
           17           0.000000
           18           0.000000
           19           0.000000
           20           0.000000
           21           0.000000
           22           0.000000
           23           0.000000
           24           0.000000
           25           0.000000
           26           0.000000
           27           0.000000
           28           0.000000
           29           0.000000
           30           0.000000
           31           0.000000
           32           0.000000

关于cuda - 随机访问时共享内存中的预期存储体冲突数,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/12823597/

相关文章:

c++ - 将 OpenCV 图像数据类型转换为 Devil 图像格式,反之亦然

c++ - 获取CUDA错误 “declaration is incompatible with previous ” variable_name“

CUDA共享内存

CUDA 内联 PTX ld.shared 遇到 cudaErrorIllegalAddress 错误

c - 非 GPU 硬件上是否会发生 Bank 冲突?

cuda - 共享内存库与 char 数组冲突

c - 将 openMp 程序移植到 cuda c : correct grid_size/block_size and reduction

c++ - 在内核上工作的 CUDA 上的 vector

c++ - CMakeLists 问题:在 thrust::device_vector 上调用 resize() 时出现 bad_alloc 错误

c++ - CUDA - 确定共享内存中的银行数量