如何以最高性能规范化 CUDA 中的矩阵列?

2022-08-01 21:00:30 other 阅读：814
本文介绍了如何以最高性能规范化 CUDA 中的矩阵列?的处理方法，对大家解决问题具有一定的参考价值

问题描述

如何在 CUDA 中有效地规范化矩阵列?

How to effectively normalize matrix columns in CUDA?

我的矩阵是列优先存储的，典型的大小是2000x200.

My matrix is stored in column-major, and the typical size is 2000x200.

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

The operation can be represented in the following matlab code.

A = rand(2000,200);

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

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

Can this be done effectively by Thrust, cuBLAS and/or cuNPP?

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

A rapid implementation including 4 kernels is shown as follows.

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

Wondering if these can be done in 1 or 2 kernels to improve the performance, especially for the column summation step implemented by 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;
}

推荐答案

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

I compared the performance of 3 approaches on M2090 with CUDA 5.0.

  1. [173.179 us] cublas 实现如问题所示
  2. [733.734 us] 使用来自@talonmies 的 thrust::reduce_by_key 实现纯 Thrust 实现
  3. [1.508 ms] 使用 thrust::inclusive_scan_by_key
    4. 的纯 Thrust 实现
  1. [173.179 us] cublas implementation as shown in the question
  2. [733.734 us] pure Thrust implementation with thrust::reduce_by_key from @talonmies
  3. [1.508 ms] pure Thrust implementation with thrust::inclusive_scan_by_key

可以看出，

  1. cublas 在这种情况下性能最高；
  2. thrust::reduce_by_key 和thrust::inclusive_scan_by_key 启动多个内核，导致额外开销；
  3. thrust::inclusive_scan_by_key 与 thrust::reduce_by_key 相比，向 DRAM 写入的数据要多得多，这可能是内核时间较长的原因之一；
  4. cublas 和推力方法之间的主要性能差异是矩阵列求和.推力较慢可能是因为 thrust::reduce_by_key 旨在对具有可变长度的段进行缩减，但 cublas_gemv() 只能应用于固定长度的段(行/列).
  1. cublas has highest performance in this case;
  2. both thrust::reduce_by_key & thrust::inclusive_scan_by_key launch multiple kernels, which lead to extra overhead;
  3. thrust::inclusive_scan_by_key writes much more data to DRAM compared to thrust::reduce_by_key, which can be one of the reasons for longer kernel time;
  4. the main performance difference between cublas and thrust approach is the matrix column summation. thrust is slower possibly because thrust::reduce_by_key is designed to do reduction on segments with variant length, but cublas_gemv() can only apply to fixed length segments (row/col).

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

When the matrix A is large enough to ignore the kernel launching overhead, the cublas appoach still performs best. The profiling result on A_{20,000 x 2,000} is shown as follows.

将第一个 for_each 操作与 @talonmies 指示的 cublasSgemv 调用融合可能会进一步提高性能，但我认为应该使用手工编写的内核而不是 <代码>推力::reduce_by_key.

Fusing the first for_each operation with the cublasSgemv call as indicated by @talonmies may further improve the performance, but I think kernel written by hand should be used instead of thrust::reduce_by_key.

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

The code for the 3 approaches is shown as follows.

#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;
}

这篇关于如何以最高性能规范化 CUDA 中的矩阵列?的文章就介绍到这了，希望我们推荐的答案对大家有所帮助，WP2

admin_action_{$_REQUEST[‘action’]}

do_action( "admin_action_{$_REQUEST[‘action’]}" )动作钩子：：在发送“Action”请求变量时激发。Action Hook: Fires when an ‘action’ request variable is sent.目录锚点：#说明#源码说明（Description）钩子名称的动态部分$_REQUEST['action']引用从GET或POST请求派生的操作。源码（Source）更新版本源码位置使用被使用2.6.0 wp-admin/admin.php:...

日期：2020-09-02 17:44:16 浏览：1159

admin_footer-{$GLOBALS[‘hook_suffix’]}

do_action( "admin_footer-{$GLOBALS[‘hook_suffix’]}", string $hook_suffix )操作挂钩：在默认页脚脚本之后打印脚本或数据。Action Hook: Print scripts or data after the default footer scripts.目录锚点：#说明#参数#源码说明（Description）钩子名的动态部分，$GLOBALS['hook_suffix']引用当前页的全局钩子后缀。参数（Parameters）参数类...

日期：2020-09-02 17:44:20 浏览：1060

customize_save_{$this->id_data[‘base’]}

do_action( "customize_save_{$this-&gt;id_data[‘base’]}", WP_Customize_Setting $this )动作钩子：：在调用WP_Customize_Setting：：save（）方法时激发。Action Hook: Fires when the WP_Customize_Setting::save() method is called.目录锚点：#说明#参数#源码说明（Description）钩子名称的动态部分，$this->id_data...

日期：2020-08-15 15:47:24 浏览：800

customize_value_{$this->id_data[‘base’]}

apply_filters( "customize_value_{$this-&gt;id_data[‘base’]}", mixed $default )过滤器：：过滤未作为主题模式或选项处理的自定义设置值。Filter Hook: Filter a Customize setting value not handled as a theme_mod or option.目录锚点：#说明#参数#源码说明（Description）钩子名称的动态部分，$this->id_date['base']，指的是设置...

日期：2020-08-15 15:47:24 浏览：888

get_comment_author_url

过滤钩子：过滤评论作者的URL。Filter Hook: Filters the comment author’s URL.目录锚点：#源码源码（Source）更新版本源码位置使用被使用 wp-includes/comment-template.php:32610...

日期：2020-08-10 23:06:14 浏览：925

network_admin_edit_{$_GET[‘action’]}

do_action( "network_admin_edit_{$_GET[‘action’]}" )操作挂钩：启动请求的处理程序操作。Action Hook: Fires the requested handler action.目录锚点：#说明#源码说明（Description）钩子名称的动态部分$u GET['action']引用请求的操作的名称。源码（Source）更新版本源码位置使用被使用3.1.0 wp-admin/network/edit.php:3600...

日期：2020-08-02 09:56:09 浏览：873

network_sites_updated_message_{$_GET[‘updated’]}

apply_filters( "network_sites_updated_message_{$_GET[‘updated’]}", string $msg )筛选器挂钩：在网络管理中筛选特定的非默认站点更新消息。Filter Hook: Filters a specific, non-default site-updated message in the Network admin.目录锚点：#说明#参数#源码说明（Description）钩子名称的动态部分$_GET['updated']引用了非默认的...

日期：2020-08-02 09:56:03 浏览：855

pre_wp_is_site_initialized

过滤器：：过滤在访问数据库之前是否初始化站点的检查。Filter Hook: Filters the check for whether a site is initialized before the database is accessed.目录锚点：#源码源码（Source）更新版本源码位置使用被使用 wp-includes/ms-site.php:93910...

日期：2020-07-29 10:15:38 浏览：825

WordPress 的SEO 教学：如何在网站中加入关键字（Meta Keywords）与Meta 描述（Meta Description）？

你想在WordPress 中添加关键字和meta 描述吗？关键字和meta 描述使你能够提高网站的SEO。在本文中，我们将向你展示如何在WordPress 中正确添加关键字和meta 描述。为什么要在WordPress 中添加关键字和Meta 描述？关键字和说明让搜寻引擎更了解您的帖子和页面的内容。关键词是人们寻找您发布的内容时，可能会搜索的重要词语或片语。而Meta Description则是对你的页面和文章的简要描述。如果你想要了解更多关于中继标签的资讯，可以参考Google的说明。Meta 关键字和描...

日期：2020-10-03 21:18:25 浏览：1695

谷歌的SEO是什么

SEO (Search Engine Optimization)中文是搜寻引擎最佳化，意思近于「关键字自然排序」、「网站排名优化」。简言之，SEO是以搜索引擎（如Google、Bing）为曝光媒体的行销手法。例如搜寻「wordpress教学」，会看到本站的「WordPress教学：12个课程…」排行Google第一：关键字：wordpress教学、wordpress课程…若搜寻「网站架设」，则会看到另一个网页排名第1：关键字：网站架设、架站…以上两个网页，每月从搜寻引擎导入自然流量，达2万4千：每月「有机搜...

日期：2020-10-30 17:23:57 浏览：1298