- html - 出于某种原因,IE8 对我的 Sass 文件中继承的 html5 CSS 不友好?
- JMeter 在响应断言中使用 span 标签的问题
- html - 在 :hover and :active? 上具有不同效果的 CSS 动画
- html - 相对于居中的 html 内容固定的 CSS 重复背景?
如何在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;
}
最佳答案
我比较了 M2090 和 CUDA 5.0 上 3 种方法的性能。
thrust::reduce_by_key
实现纯 Thrustthrust::inclusive_scan_by_key
实现纯 Thrust
可以看出,
thrust::reduce_by_key
和 thrust::inclusive_scan_by_key
都会启动多个内核,这会导致额外的开销;thrust::reduce_by_key
相比,thrust::inclusive_scan_by_key
向 DRAM 写入了更多数据,这可能是内核时间较长的原因之一;thrust::reduce_by_key
旨在对不同长度的段进行缩减,但 cublas_gemv()
只能应用于固定长度段(行/列) .当矩阵 A 足够大以忽略内核启动开销时,cublas appoach 仍然表现最佳。 A_{20,000 x 2,000} 上的分析结果如下所示。
将第一个 for_each
操作与 @talonmies 所示的 cublasSgemv
调用融合可能会进一步提高性能,但我认为应该使用手工编写的内核而不是 thrust::reduce_by_key
。
这3种方法的代码如下所示。
#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;
}
关于performance - 如何以最大性能标准化 CUDA 中的矩阵列?,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/14211093/
在这段令人惊叹的视频 ( https://www.youtube.com/watch?v=udix3GZouik ) 中,Alex Blom 谈到了 Ember 在移动世界中的“黑客攻击”。 在 22
我们希望通过我们的应用收集使用情况统计信息。因此,我们希望在服务器端的某个地方跟踪用户操作。 就性能而言,哪个选项更合适: 在 App Engine 请求日志中跟踪用户操作。即为每个用户操作写入一个日
在针对对象集合的 LINQ 查询的幕后究竟发生了什么?它只是语法糖还是发生了其他事情使其更有效的查询? 最佳答案 您是指查询表达式,还是查询在幕后的作用? 查询表达式首先扩展为“普通”C#。例如: v
我正在构建一个简单的照片库应用程序,它在列表框中显示图像。 xaml 是:
对于基于 Web 的企业应用程序,使用“静态 Hashmap 存储对象” 和 apache java 缓存系统有何优缺点?哪一个最有利于性能并减少堆内存问题 例如: Map store=Applica
我想知道在性能方面存储类变量的最佳方式是什么。我的意思是,由于 Children() 函数,存储一个 div id 比查找所有其他类名更好。还是把类名写在变量里比较好? 例如这样: var $inne
我已经阅读了所有这些关于 cassandra 有多快的文章,例如单行读取可能需要大约 5 毫秒。 到目前为止,我不太关心我的网站速度,但是随着网站变得越来越大,一些页面开始需要相当多的查询,例如一个页
最近,我在缓存到内存缓存之前的查询一直需要很长时间才能处理!在这个例子中,它花费了 10 秒。在这种情况下,我要做的就是获得 10 个最近的点击。 我感觉它加载了所有 125,592 行然后只返回 1
我找了几篇文章(包括SA中的一些问题),试图找到基本操作的成本。 但是,我尝试制作自己的小程序,以便自己进行测试。在尝试测试加法和减法时,我遇到了一些问题,我用简单的代码向您展示了这一点
这个问题在这里已经有了答案: Will Java app slow down by presence of -Xdebug or only when stepping through code? (
我记得很久以前读过 with() 对 JavaScript 有一些严重的性能影响,因为它可能对范围堆栈进行非确定性更改。我很难找到最近对此的讨论。这仍然是真的吗? 最佳答案 与其说 with 对性能有
我们有一个数据仓库,其中包含非规范化表,行数从 50 万行到 6 多万行不等。我正在开发一个报告解决方案,因此出于性能原因我们正在使用数据库分页。我们的报告有搜索条件,并且我们已经创建了必要的索引,但
我有一条有效的 SQL 语句,但需要很长时间才能处理 我有一个 a_log 表和一个 people 表。我需要在 people 表中找到给定人员的每个 ID 的最后一个事件和关联的用户。 SELECT
很难说出这里问的是什么。这个问题是含糊的、模糊的、不完整的、过于宽泛的或修辞性的,无法以目前的形式得到合理的回答。如需帮助澄清此问题以便重新打开它,visit the help center 。 已关
通常当我建立一个站点时,我将所有的 CSS 放在一个文件中,并且一次性定义与一组元素相关的所有属性。像这样: #myElement { color: #fff; background-
两者之间是否存在任何性能差异: p { margin:0px; padding:0px; } 并省略最后的分号: p { margin:0px; padding:0px } 提前致谢!
我的应用程序 (PHP) 需要执行大量高精度数学运算(甚至可能出现一共100个数字) 通过这个论坛的最后几篇帖子,我发现我必须使用任何高精度库,如 BC Math 或 GMP,因为 float 类型不
我一直在使用 javamail 从 IMAP 服务器(目前是 GMail)检索邮件。 Javamail 非常快速地从服务器检索特定文件夹中的消息列表(仅 id),但是当我实际获取消息(仅包含甚至不包含
我非常渴望开发我的第一个 Ruby 应用程序,因为我的公司终于在内部批准了它的使用。 在我读到的关于 Ruby v1.8 之前的所有内容中,从来没有任何关于性能的正面评价,但我没有发现关于 1.9 版
我是 Redis 的新手,我有一个包含数百万个成员(member) ID、电子邮件和用户名的数据集,并且正在考虑将它们存储在例如列表结构中。我认为 list 和 sorted set 可能最适合我的情
我是一名优秀的程序员,十分优秀!