使用 Nvidia 的推力库规范化一堆矢量
Normalize a bunch of vectors using Nvidia's Thrust library
我刚刚了解了Nvidia的推力库。只是为了尝试写一个小的例子,它应该是标准化一堆向量。
#include <cstdio>
#include <thrust/transform.h>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
struct normalize_functor: public thrust::unary_function<double4, double4>
{
__device__ __host__ double4 operator()(double4 v)
{
double len = sqrt(v.x*v.x + v.y*v.y + v.z*v.z);
v.x /= len;
v.y /= len;
v.z /= len;
printf("%f %f %fn", v.x, v.y, v.z);
}
};
int main()
{
thrust::host_vector<double4> v(2);
v[0].x = 1; v[0].y = 2; v[0].z = 3;
v[1].x = 4; v[1].y = 5; v[1].z = 6;
thrust::device_vector<double4> v_d = v;
thrust::for_each(v_d.begin(), v_d.end(), normalize_functor());
// This doesn't seem to copy back
v = v_d;
// Neither this does..
thrust::host_vector<double4> result = v_d;
for(int i=0; i<v.size(); i++)
printf("[ %f %f %f ]n", result[i].x, result[i].y, result[i].z);
return 0;
}
上面的例子似乎工作,但是我无法复制数据回来。我认为一个简单的赋值会调用cudaMemcpy。它的工作原理是将数据从主机复制到设备,但不返回?
其次,我不确定我这样做是否正确。for_each的文档说:
for_each将函数对象f应用于[first, last)范围内的每个元素;F的返回值(如果有的话)将被忽略。
但是unary_function结构模板需要两个模板参数(一个用于返回值)并强制操作符()也返回值,这在编译时导致警告。我不明白为什么我应该写一个没有返回值的一元函子。
下一步是数据排列。我只是选择了double4,因为这将导致两个读取指令ld.v2。f64和ld.f64 IIRC。然而,我想知道如何在引擎盖下获取数据(以及创建多少cuda线程/块)。如果我选择一个由4个向量组成的结构体,它是否能够以合并的方式获取数据?
最后,thrust提供了元组。那么元组数组呢?在这种情况下,数据如何排列?
我查看了示例,但我还没有找到一个示例来解释为一堆向量(dot_products_with_zip)选择哪种数据结构。这个例子说的是"数组的结构"而不是"结构的数组",但是我在这个例子中没有看到任何结构。
我修复了上面的代码,并尝试运行一个更大的例子,这次规范化了10k个向量。
#include <cstdio>
#include <thrust/transform.h>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
struct normalize_functor
{
__device__ __host__ void operator()(double4& v)
{
double len = sqrt(v.x*v.x + v.y*v.y + v.z*v.z);
v.x /= len;
v.y /= len;
v.z /= len;
}
};
int main()
{
int n = 10000;
thrust::host_vector<double4> v(n);
for(int i=0; i<n; i++) {
v[i].x = rand();
v[i].y = rand();
v[i].z = rand();
}
thrust::device_vector<double4> v_d = v;
thrust::for_each(v_d.begin(), v_d.end(), normalize_functor());
v = v_d;
return 0;
}
用computeprof分析显示了一个低占用和非合并的内存访问:
Kernel Occupancy Analysis
Kernel details : Grid size: 23 x 1 x 1, Block size: 448 x 1 x 1
Register Ratio = 0.984375 ( 32256 / 32768 ) [24 registers per thread]
Shared Memory Ratio = 0 ( 0 / 49152 ) [0 bytes per Block]
Active Blocks per SM = 3 / 8
Active threads per SM = 1344 / 1536
Potential Occupancy = 0.875 ( 42 / 48 )
Max achieved occupancy = 0.583333 (on 9 SMs)
Min achieved occupancy = 0.291667 (on 5 SMs)
Occupancy limiting factor = Block-Size
Memory Throughput Analysis for kernel launch_closure_by_value on device GeForce GTX 470
Kernel requested global memory read throughput(GB/s): 29.21
Kernel requested global memory write throughput(GB/s): 17.52
Kernel requested global memory throughput(GB/s): 46.73
L1 cache read throughput(GB/s): 100.40
L1 cache global hit ratio (%): 48.15
Texture cache memory throughput(GB/s): 0.00
Texture cache hit rate(%): 0.00
L2 cache texture memory read throughput(GB/s): 0.00
L2 cache global memory read throughput(GB/s): 42.44
L2 cache global memory write throughput(GB/s): 46.73
L2 cache global memory throughput(GB/s): 89.17
L2 cache read hit ratio(%): 88.86
L2 cache write hit ratio(%): 3.09
Local memory bus traffic(%): 0.00
Global memory excess load(%): 31.18
Global memory excess store(%): 62.50
Achieved global memory read throughput(GB/s): 4.73
Achieved global memory write throughput(GB/s): 45.29
Achieved global memory throughput(GB/s): 50.01
Peak global memory throughput(GB/s): 133.92
我想知道如何优化这个?
如果你想用for_each
就地修改序列,那么你需要在函子中通过引用获取实参:
struct normalize_functor
{
__device__ __host__ void operator()(double4& ref)
{
double v = ref;
double len = sqrt(v.x*v.x + v.y*v.y + v.z*v.z);
v.x /= len;
v.y /= len;
v.z /= len;
printf("%f %f %fn", v.x, v.y, v.z);
ref = v;
}
};
或者,您可以将normalize_functor
的定义与transform
算法一起使用,指定v_d
作为源和目标范围:
thrust::transform(v_d.begin(), v_d.end(), v_d.begin(), normalize_functor());
我个人倾向于在这种情况下使用transform
,但是在任何一种情况下性能应该是相同的。
关于优化的问题,没有太多可以用Thrust做的事情——这并不是库的真正意图。Nathan Bell是Thrust的作者之一,他已经在这个帖子中发帖了,我不想为他说话,我们的目标是以一种简单、直观的方式为GPU提供一系列数据并行算法,而无需编写太多CUDA代码。在我看来,它在这方面非常成功。许多thrust内核的内核性能接近最先进的水平,但是总有一些优化可以在特定情况下完成,这在通用模板代码中是不容易做到的。这是您为Thrust提供的易用性和灵活性所付出的部分代价。
话虽如此,我怀疑在你的操作符函数中有一些调整可以尝试,这可能会改善事情。我通常会这样写:
struct normalize_functor
{
__device__ __host__ void operator()(double4& v)
{
double4 nv = v;
double len = sqrt(nv.x*nv.x + nv.y*nv.y + nv.z*nv.z);
nv.x /= len;
nv.y /= len;
nv.z /= len;
(void)nv.h;
v = nv;
};
};
现在,虽然它不像原来那样漂亮,但它应该确保编译器发出矢量化的加载和存储指令。我在过去看到过这样的情况,编译器会优化掉vector类型中未使用成员的加载和存储,这会导致PTX生成器发出标量的加载和存储,从而导致合并中断。通过使用一个明确的float4加载和存储,并确保使用了结构的每个元素,它可以避免至少在2中出现的不必要的"优化"。X和3。X NVCC版本。我不确定4.0编译器是否仍然是这种情况。
- 如何使用结构和指针推动和弹出一堆双打
- 如何改进一堆在已知值范围内评估变量的 else-if 条件?
- 当我在结构中包含多个数组时,我的程序会跳过一堆代码
- 在共享指针的值中调用 std::swap 调用一堆构造函数和析构函数
- 使用带有.lib和一堆.dll和.h文件的Python CFFI
- 用 c++ 编写一堆类似的 if 语句的漂亮方法
- 当我尝试编译程序时,我遇到了一堆错误,例如:'std::max':找不到匹配的重载函数
- Swig C++ to python:编译一堆.cpp和.h文件
- template元编程:将一堆模板参数相乘
- 在.NET中存储一堆恒定值的最佳方法
- 我的一堆函数出现"undefined reference"错误,我不知道为什么
- 制作一堆 int 数组
- 为什么glut.h会在CodeBlocks中弹出一堆未定义的引用
- 更换一堆std ::矢量条目
- 在Windows下编译程序会给出一堆"error: template with C linkage"报告
- C++将字符从静态数组复制到动态数组会添加一堆随机元素
- 如何创建一堆多维数组
- 如何在Linux中使用Makefile一次编译一堆测试C++文件
- 如何在一堆数字中找到一个不同的值
- 使用 Nvidia 的推力库规范化一堆矢量