用于处理大数据的 Pycuda 块和网格答案

【问题标题】：Pycuda Blocks and Grids to work with big datas用于处理大数据的 Pycuda 块和网格
【发布时间】：2013-01-24 15:12:14
【问题描述】：

我需要帮助来了解我的块和网格的大小。我正在构建一个 python 应用程序来执行基于 scipy 的度量计算：欧几里德距离、曼哈顿、皮尔森、余弦、加入其他。

项目是PycudaDistances。

它似乎适用于小型数组。当我进行更详尽的测试时，不幸的是它没有用。我下载了movielens set (http://www.grouplens.org/node/73)。

使用Movielens 100k，我声明了一个形状为 (943, 1682) 的数组。即用户对943部和1682部影片进行评价。不是分类器用户的电影我将值配置为 0。

使用更大的数组算法不再有效。我面临以下错误：

pycuda._driver.LogicError: cuFuncSetBlockShape 失败：值无效。

研究这个错误，我找到了一个解释，告诉 Andrew 支持 512 个线程加入和处理更大的块，必须使用块和网格。

我需要帮助来调整欧几里得距离数组算法，使其适用于从小型数组到巨型数组。

def euclidean_distances(X, Y=None, inverse=True):
    X, Y = check_pairwise_arrays(X,Y)
    rows = X.shape[0]
    cols = Y.shape[0]
    solution = numpy.zeros((rows, cols))
    solution = solution.astype(numpy.float32)

    kernel_code_template = """
    #include <math.h>
    
    __global__ void euclidean(float *x, float *y, float *solution) {

        int idx = threadIdx.x + blockDim.x * blockIdx.x;
        int idy = threadIdx.y + blockDim.y * blockIdx.y;
        
        float result = 0.0;
        
        for(int iter = 0; iter < %(NDIM)s; iter++) {
            
            float x_e = x[%(NDIM)s * idy + iter];
            float y_e = y[%(NDIM)s * idx + iter];
            result += pow((x_e - y_e), 2);
        }
        int pos = idx + %(NCOLS)s * idy;
        solution[pos] = sqrt(result);
    }
    """
    kernel_code = kernel_code_template % {
        'NCOLS': cols,
        'NDIM': X.shape[1]
    }

    mod = SourceModule(kernel_code)

    func = mod.get_function("euclidean")
    func(drv.In(X), drv.In(Y), drv.Out(solution), block=(cols, rows, 1))

    return numpy.divide(1.0, (1.0 + solution)) if inverse else solution

【问题讨论】：

标签： python cuda gpu pycuda euclidean-distance

【解决方案1】：

要调整内核的执行参数，您需要做两件事（按此顺序）：

1。确定块大小

您的块大小主要取决于硬件限制和性能。我建议阅读this answer 以获取更多详细信息，但简短的总结是，您的 GPU 对其可以运行的每个块的线程总数有限制，并且它具有有限的寄存器文件、共享和本地内存大小。您选择的块尺寸必须在这些限制范围内，否则内核将无法运行。块大小也会影响内核的性能，您会发现一个块大小可以提供最佳性能。块大小应该始终是 warp 大小的整数倍，在迄今为止发布的所有 CUDA 兼容硬件上都是 32。

2。确定网格大小

对于您展示的那种内核，您需要的块数与输入数据量和每个块的尺寸直接相关。

例如，如果您的输入数组大小为 943x1682，而您的块大小为 16x16，则您需要一个 59 x 106 的网格，这将在内核启动时产生 944x1696 个线程。在这种情况下，输入数据大小不是块大小的整数倍，您需要修改内核以确保它不会越界读取。一种方法可能是这样的：

__global__ void euclidean(float *x, float *y, float *solution) {
    int idx = threadIdx.x + blockDim.x * blockIdx.x;
    int idy = threadIdx.y + blockDim.y * blockIdx.y;

     if ( ( idx < %(NCOLS)s ) && ( idy < %(NDIM)s ) ) {

        .....
     }
}

启动内核的 Python 代码可能类似于：

bdim = (16, 16, 1)
dx, mx = divmod(cols, bdim[0])
dy, my = divmod(rows, bdim[1])

gdim = ( (dx + (mx>0)) * bdim[0], (dy + (my>0)) * bdim[1]) )
func(drv.In(X), drv.In(Y), drv.Out(solution), block=bdim, grid=gdim)

This question and answer 也可能有助于了解此过程的工作原理。

请注意，以上所有代码都是在浏览器中编写的，从未经过测试。需要您自担风险使用它。

另请注意，它是基于对您的代码的非常简短的阅读，可能不正确，因为您没有真正描述任何关于如何在您的问题中调用代码的内容。

【讨论】：

【解决方案2】：

原则上接受的答案是正确的，但是 talonmies 列出的代码并不完全正确。该行： gdim = ( (dx + (mx>0)) * bdim[0], (dy + (my>0)) * bdim[1]) ) 应该： gdim = ( (dx + (mx>0)), (dy + (my>0)) ) 除了一个明显的额外括号外， gdim 会产生比你想要的太多的线程。 talonmies 在他的文章中已经解释过，线程是 blocksize * gridSize。然而，他列出的 gdim 会为您提供总线程数，而不是所需的正确网格大小。

【讨论】：