cuda内核for循环中的Break语句给出了问题

Posted 2023-03-23

【中文标题】cuda内核for循环中的Break语句给出了问题

【英文标题】：Break statement inside a cuda kernel for loop is giving issues

【发布时间】：2018-03-25 03:13:33

【问题描述】：

我最近在玩 cuda/numba 代码。我有一个 MxN 矩阵（例如 cumul_A），其中每一行都是累积概率分布。我想通过从均匀随机分布中映射样本来从这些累积分布中抽取样本。简单来说，假设从均匀随机分布中抽取的样本是 0.3。 cuda 内核应该选择一行“cumul_A”并将该行的每个元素（从该行的第一个元素开始）与 0.3 进行比较。一旦它得到一个大于 0.3 的值，内核应该将元素的索引存储在输出参数中并打破 for 循环。我无法让这个看似简单的内核工作。 break 语句是否会在内核内部造成任何问题？ 下面提供了最小的工作示例。

    from __future__ import division
    from __future__ import print_function

    import numpy as np

    from numba import vectorize, cuda, jit
    np.set_printoptions(precision=4, suppress=True)

    # Number of rows
    M = 10
    # Number of columns
    N = 20

    # ======= 1-D GRIDS =======
    # Set the number of threads in a block
    threadsperblock_1d = 4
    # Calculate the number of thread blocks in the grid
    blockspergrid_1d = np.int(np.ceil(M / threadsperblock_1d))
    # ======= 1-D GRIDS =======

    @cuda.jit('void(float32[:, :], float32[:], int32[:])')
    def get_randomchoice(cumul_a, random_nos, output):
      x = cuda.grid(1) 
      if x < cumul_a.shape[0]:
        for y in range(cumul_a.shape[1]):
          if random_nos[x] > cumul_a[x, y]:
            output[x] = y
            break # return

    if __name__ == '__main__':
      # Prepare the matrix whise each row is a cumulative probability distribution
      A = np.random.rand(M, N).astype(np.float32)
      A = np.divide(A,np.sum(A,axis=1,keepdims=True))
      cumul_A = np.cumsum(A, axis=1)

      # Put an assertion that cumul_A is indeed cumulative
      assert np.allclose(cumul_A[:,-1],np.ones(M))

      # Draw values from uniform distribution
      RandValues = np.random.rand(M).astype(np.float32)

      # Output array in numpy
      Y = np.zeros(M, dtype=np.int32)
      for iStep in range(M):
        Y[iStep] = np.argwhere(RandValues[iStep] <= cumul_A[iStep])[0]

      print('From numpy:\n'.format(Y))

      # Transfer to GPU
      cumul_A_gpu = cuda.to_device(cumul_A)
      RandValues_gpu = cuda.to_device(RandValues)
      # Return array from GPU
      random_idx_gpu = cuda.device_array(M, dtype=np.int32)
      get_randomchoice[blockspergrid_1d, threadsperblock_1d](cumul_A_gpu, RandValues_gpu, random_idx_gpu)
      random_idx = random_idx_gpu.copy_to_host()

      print('From cuda:\n'.format(random_idx))

我们将不胜感激。

【参考方案1】：

这是一场虚惊！代码有问题。线if random_nos[x] > cumul_a[x, y]: 将是if random_nos[x] < cumul_a[x, y]:。 for 循环中的“break”没有问题。