如何以最大性能标准化 CUDA 中的矩阵列？答案

【问题标题】：How to normalize matrix columns in CUDA with max performance?如何以最大性能标准化 CUDA 中的矩阵列？
【发布时间】：2012-12-22 01:49:42
【问题描述】：

如何有效地对 CUDA 中的矩阵列进行归一化？

我的矩阵以列为主存储，典型大小为 2000x200。

该操作可以用下面的matlab代码表示。

A = rand(2000,200);

A = exp(A);
A = A./repmat(sum(A,1), [size(A,1) 1]);

这可以通过 Thrust、cuBLAS 和/或 cuNPP 有效地完成吗？

一个包含 4 个内核的快速实现如下所示。

想知道这些是否可以在 1 个或 2 个内核中完成以提高性能，特别是对于 cublasDgemv() 实现的列求和步骤。

#include <cuda.h>
#include <curand.h>
#include <cublas_v2.h>
#include <thrust/device_vector.h>
#include <thrust/device_ptr.h>
#include <thrust/transform.h>
#include <thrust/iterator/constant_iterator.h>
#include <math.h>

struct Exp
{
    __host__ __device__ void operator()(double& x)
    {
        x = exp(x);
    }
};

struct Inv
{
    __host__ __device__ void operator()(double& x)
    {
        x = (double) 1.0 / x;
    }
};

int main()
{
    cudaDeviceSetCacheConfig(cudaFuncCachePreferShared);
    cublasHandle_t hd;
    curandGenerator_t rng;
    cublasCreate(&hd);
    curandCreateGenerator(&rng, CURAND_RNG_PSEUDO_DEFAULT);

    const size_t m = 2000, n = 200;
    const double c1 = 1.0;
    const double c0 = 0.0;

    thrust::device_vector<double> A(m * n);
    thrust::device_vector<double> sum(1 * n);
    thrust::device_vector<double> one(m * n, 1.0);

    double* pA = thrust::raw_pointer_cast(&A[0]);
    double* pSum = thrust::raw_pointer_cast(&sum[0]);
    double* pOne = thrust::raw_pointer_cast(&one[0]);

    for (int i = 0; i < 100; i++)
    {
        curandGenerateUniformDouble(rng, pA, A.size());


        thrust::for_each(A.begin(), A.end(), Exp());

        cublasDgemv(hd, CUBLAS_OP_T, m, n,
                &c1, pA, m, pOne, 1, &c0, pSum, 1);

        thrust::for_each(sum.begin(), sum.end(), Inv());

        cublasDdgmm(hd, CUBLAS_SIDE_RIGHT, m, n, pA, m, pSum, 1, pA, m);
    }

    curandDestroyGenerator(rng);
    cublasDestroy(hd);

    return 0;
}

【问题讨论】：

是的，它可以通过 CUDA 有效地完成。展示你为实现你想要的而编写的一些 CUDA 代码。
添加了代码。寻求绩效改进

标签： performance matrix cuda thrust cublas

【解决方案1】：

我比较了 M2090 和 CUDA 5.0 上 3 种方法的性能。

[173.179 us] cublas 实现如问题所示
[733.734 us] 纯推力实现与来自@talonmies 的thrust::reduce_by_key
[1.508 ms] 纯推力实现与thrust::inclusive_scan_by_key

可以看出，

cublas 在这种情况下性能最高；
thrust::reduce_by_key 和thrust::inclusive_scan_by_key 都启动多个内核，这会导致额外的开销；
与thrust::reduce_by_key 相比，thrust::inclusive_scan_by_key 向 DRAM 写入的数据要多得多，这可能是内核时间更长的原因之一；
cublas 和推力方法之间的主要性能差异是矩阵列求和。推力较慢可能是因为thrust::reduce_by_key 旨在对具有不同长度的段进行缩减，但cublas_gemv() 只能适用于固定长度的段（行/列）。

当矩阵 A 大到足以忽略内核启动开销时，cublas 方法仍然表现最佳。 A_{20,000 x 2,000} 上的分析结果如下所示。

如@talonmies 所示，将第一个for_each 操作与cublasSgemv 调用融合可能会进一步提高性能，但我认为应该使用手工编写的内核而不是thrust::reduce_by_key。

三种方法的代码如下所示。

#include <cuda.h>
#include <curand.h>
#include <cublas_v2.h>
#include <thrust/device_vector.h>
#include <thrust/device_ptr.h>
#include <thrust/transform.h>
#include <thrust/reduce.h>
#include <thrust/scan.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/iterator/discard_iterator.h>
#include <thrust/iterator/permutation_iterator.h>
#include <math.h>

struct Exp: public thrust::unary_function<double, double>
{
    __host__ __device__ double operator()(double x)
    {
        return exp(x);
    }
};

struct Inv: public thrust::unary_function<double, double>
{
    __host__ __device__ double operator()(double x)
    {
        return (double) 1.0 / x;
    }
};

template<typename T>
struct MulC: public thrust::unary_function<T, T>
{
    T C;
    __host__ __device__ MulC(T c) :
        C(c)
    {
    }
    __host__ __device__ T operator()(T x)
    {
        return x * C;
    }
};

template<typename T>
struct line2col: public thrust::unary_function<T, T>
{
    T C;
    __host__ __device__ line2col(T C) :
            C(C)
    {
    }

    __host__ __device__ T operator()(T i)
    {
        return i / C;
    }
};

int main()
{
    cudaDeviceSetCacheConfig(cudaFuncCachePreferShared);
    cublasHandle_t hd;
    curandGenerator_t rng;
    cublasCreate(&hd);
    curandCreateGenerator(&rng, CURAND_RNG_PSEUDO_DEFAULT);

    const size_t m = 2000, n = 200;
    const double c1 = 1.0;
    const double c0 = 0.0;

    thrust::device_vector<double> A(m * n);
    thrust::device_vector<double> B(m * n);
    thrust::device_vector<double> C(m * n);
    thrust::device_vector<double> sum1(1 * n);
    thrust::device_vector<double> sum2(1 * n);
    thrust::device_vector<double> one(m * n, 1);

    double* pA = thrust::raw_pointer_cast(&A[0]);
    double* pB = thrust::raw_pointer_cast(&B[0]);
    double* pSum1 = thrust::raw_pointer_cast(&sum1[0]);
    double* pSum2 = thrust::raw_pointer_cast(&sum2[0]);
    double* pOne = thrust::raw_pointer_cast(&one[0]);

    curandGenerateUniformDouble(rng, pA, A.size());

    const int count = 2;

    for (int i = 0; i < count; i++)
    {
        thrust::transform(A.begin(), A.end(), B.begin(), Exp());
        cublasDgemv(hd, CUBLAS_OP_T, m, n, &c1, pB, m, pOne, 1, &c0, pSum1, 1);
        thrust::transform(sum1.begin(), sum1.end(), sum1.begin(), Inv());
        cublasDdgmm(hd, CUBLAS_SIDE_RIGHT, m, n, pB, m, pSum2, 1, pB, m);
    }

    for (int i = 0; i < count; i++)
    {
        thrust::reduce_by_key(
                thrust::make_transform_iterator(thrust::make_counting_iterator(0), line2col<int>(m)),
                thrust::make_transform_iterator(thrust::make_counting_iterator(0), line2col<int>(m)) + A.size(),
                thrust::make_transform_iterator(A.begin(), Exp()),
                thrust::make_discard_iterator(),
                sum2.begin());
        thrust::transform(
                A.begin(), A.end(),
                thrust::make_permutation_iterator(
                        sum2.begin(),
                        thrust::make_transform_iterator(thrust::make_counting_iterator(0), line2col<int>(m))),
                C.begin(),
                thrust::divides<double>());
    }

    for (int i = 0; i < count; i++)
    {
        thrust::inclusive_scan_by_key(
                thrust::make_transform_iterator(thrust::make_counting_iterator(0), line2col<int>(m)),
                thrust::make_transform_iterator(thrust::make_counting_iterator(0), line2col<int>(m)) + A.size(),
                thrust::make_transform_iterator(A.begin(), Exp()),
                C.begin());
        thrust::copy(
                thrust::make_permutation_iterator(
                        C.begin() + m - 1,
                        thrust::make_transform_iterator(thrust::make_counting_iterator(0), MulC<int>(m))),
                thrust::make_permutation_iterator(
                        C.begin() + m - 1,
                        thrust::make_transform_iterator(thrust::make_counting_iterator(0), MulC<int>(m))) + n,
                sum2.begin());
        thrust::transform(
                A.begin(), A.end(),
                thrust::make_permutation_iterator(
                        sum2.begin(),
                        thrust::make_transform_iterator(thrust::make_counting_iterator(0), line2col<int>(m))),
                C.begin(),
                thrust::divides<double>());
    }

    curandDestroyGenerator(rng);
    cublasDestroy(hd);

    return 0;
}

【讨论】：

【解决方案2】：

您应该能够将第一个for_each 操作与cublasSgemv 调用融合为一个reduce_by_key 调用。如果您将函子定义/重新定义为：

struct Accessor : public thrust::unary_function<int,int>
{
    int lda;
    __host__ __device__ Accessor(int _lda) : lda(_lda) {};
    __host__ __device__ int operator()(const int& idx)
    {
        return idx/lda;
    }
};

struct Exp : public thrust::unary_function<double,double>
{
    __host__ __device__ double operator()(const double& x)
    {
        return exp(x);
    }
};

struct Inv : public thrust::unary_function<double,double>
{
    __host__ __device__ double operator()(const double& x)
    {
        return double(1.0) / x;
    }
};

然后您可以将标准化输出计算为

Accessor columns(m);
thrust::reduce_by_key(
        thrust::make_transform_iterator(thrust::make_counting_iterator(int(0)), columns),
        thrust::make_transform_iterator(thrust::make_counting_iterator(int(m*n)), columns),
        thrust::make_transform_iterator(A.begin(), Exp()),
        thrust::make_discard_iterator(),
        sum.begin());

thrust::for_each(sum.begin(), sum.end(), Inv());

cublasDdgmm(hd, CUBLAS_SIDE_RIGHT, m, n, pA, m, pSum, 1, pA, m);

[免责声明：所有代码均由浏览器编写且未经测试，使用风险自负]

除了减少内核调用的数量之外，使用花哨的迭代器还消除了对大型单位矩阵的需求，这应该减少内存占用和内存事务总数来进行求和和求幂运算。

【讨论】：

迭代器真的很花哨。我比较了 cublas 和推力方法。尽管thrust::reduce_by_key 可能需要较低的内存带宽，但与cublasDgemv 相比它仍然较慢。有什么想法吗？
我怀疑相对性能很大程度上取决于您使用的 GPU 和类型。在使用 32 位类型的不同 GPU 上，您可能会发现缩减方法在性能上比纯 CUBLAS 实现更接近。推力开发人员已经承认，自从他们在推力中进行当前实现以来，最先进的缩减技术已经有了一些进步，但一般来说，树状缩减模式总是比以 FMAD 流表示的最优方式效率低，因为在这种情况下。
我还建议尝试使用thrust::transform 而不是thrust_for_each。在某些情况下（诚然是前一段时间），我发现它比for_each 快一点。但它可能不会对性能产生太大影响。

【解决方案3】：

您可以通过以下方式使用ArrayFire

array A = randu(2000, 2000);
A = exp(A);
A /= tile(sum(A, 0), A.dims(0), 1);

你也可以这样做。但是，如果您要使用矩阵（而不是普通向量），则必须在 for 循环中进行，效率不高。

免责声明我是 Accelereyes 的一名开发人员，致力于 arrayfire。

编辑我正在按要求生成新的基准。

编辑由于这个基准，我们在代码中发现了exp 的性能错误。我们正在审查和修复它。

【讨论】：

谢谢！令人印象深刻的是，代码可以像 Matlab 一样简单。您能否将您的代码的性能与我的进行比较？因为我手头没有 ArrayFire 库。
@EricShiyinKang 已更新结果。
我认为您的基准代码中存在问题，导致 cublas/thrust 方法的池计时结果。这是修改后的bench.cu
@EricShiyinKang 您在循环内外生成随机数的任何原因？我还意识到我没有在 timer::stop 之前使用设备同步，导致它对推力和阵列火的结果都产生了偏差。我正在重新修改代码。
在 curandCreateGenerator() 之后首次调用 curandGenerateUniformDouble() 需要额外的时间，如 CURAND 参考手册的Performance Notes 中所述。