python - 处理大数据的 Pycuda block 和网格

Question

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

项目是PycudaDistances .

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

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

对于更大的数组算法不再有效。我遇到以下错误:

pycuda._driver.LogicError: cuFuncSetBlockShape failed: invalid value.

研究这个错误，我找到了一个解释，告诉安德鲁支持 512 个线程加入和处理更大的 block ，有必要使用 block 和网格。

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

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

要调整内核的执行参数大小，您需要做两件事(按此顺序):

1。确定 block 大小

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

2。确定网格大小

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

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

__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也可能有助于了解此过程的工作原理。

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

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