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

Question

我需要帮助知道我的块和网格的大小。 我正在构建一个python应用程序来执行基于scipy as的度量计算：Euclidean distance、Manhattan、Pearson、Cosine、joined other。

该项目是PycudaDistances（https://github.com/vinigracindo/PycudaDistances）。

它在小数组中似乎工作得很好。当我执行更详尽的测试时，不幸的是它没有工作。我下载了movielens集（http://www.grouplens.org/node/73）。

使用Movielens 100k，我声明了一个具有shape（9431682）的数组。也就是说，用户共有943部和1682部电影被评估。电影不是由分类器用户我配置的值为0。

用大得多的数组算法不再有效。我面临以下错误：pycuda.\u driver.logicalerror: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

有关详细信息，请参见：https://github.com/vinigracindo/pycudaDistances/blob/master/distances.py

我很感激你的帮助。 已经非常感谢你了。

Answer 1

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

一。确定块大小

块大小主要由硬件限制和性能决定。我建议您阅读this answer以获得更详细的信息，但简短的总结是，您的GPU对它可以运行的每个块的线程总数有一个限制，并且它有一个有限的寄存器文件、共享内存和本地内存大小。您选择的块维度必须在这些限制内，否则内核将不会运行。块大小也会影响内核的性能，您会发现块大小可以提供最佳性能。块大小应该始终是warp大小的整数倍，到目前为止，所有CUDA兼容硬件上的warp大小都是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也可能有助于理解这个过程是如何工作的。

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

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

Answer 2

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