- html - 出于某种原因,IE8 对我的 Sass 文件中继承的 html5 CSS 不友好?
- JMeter 在响应断言中使用 span 标签的问题
- html - 在 :hover and :active? 上具有不同效果的 CSS 动画
- html - 相对于居中的 html 内容固定的 CSS 重复背景?
我对 Cuda 和 Thrust 很陌生,但我的印象是,如果使用得当,Thrust 应该会提供比单纯编写的 Cuda 内核更好的性能。我是否以次优方式使用 Thrust?下面是一个完整的、最小的例子,它接受一个数组 u
。长度N+2
,并且对于每个 i
在 1
之间和 N
计算平均值 0.5*(u[i-1] + u[i+1])
并将结果放入 uNew[i]
. (uNew[0]
设置为 u[0]
,u[N+1]
设置为 u[N+1]
,这样边界项就不会改变)。该代码多次执行此平均操作以获得用于计时测试的合理时间。在我的硬件上,Thrust 计算花费的时间大约是原始代码的两倍。有没有办法改进我的 Thrust 代码?
#include <iostream>
#include <thrust/device_vector.h>
#include <boost/timer.hpp>
#include <thrust/device_malloc.h>
typedef double numtype;
template <typename T> class NeighborAverageFunctor{
int N;
public:
NeighborAverageFunctor(int _N){
N = _N;
}
template <typename Tuple>
__host__ __device__ void operator()(Tuple t){
T uL = thrust::get<0>(t);
T uR = thrust::get<1>(t);
thrust::get<2>(t) = 0.5*(uL + uR);
}
int getN(){
return N;
}
};
template <typename T> void thrust_sweep(thrust::device_ptr<T> u, thrust::device_ptr<T> uNew, NeighborAverageFunctor<T>& op){
int N = op.getN();
thrust::for_each(thrust::make_zip_iterator(thrust::make_tuple(u, u + 2, uNew + 1)), thrust::make_zip_iterator(thrust::make_tuple(u + N, u + N+2, uNew + N+1)), op);
// Propagate boundary values without changing them
uNew[0] = u[0];
uNew[N+1] = u[N+1];
}
template <typename T> __global__ void initialization_kernel(int n, T* u){
const int i = blockIdx.x * blockDim.x + threadIdx.x;
if(i < n+2){
if(i == 0){
u[i] = 1.0;
}
else{
u[i] = 0.0;
}
}
}
template <typename T> __global__ void sweep_kernel(int n, T, T* u, T* uNew){
const int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i >= 1 && i < n-1){
uNew[i] = 0.5*(u[i+1] + u[i-1]);
}
else if(i == 0 || i == n+1){
uNew[i] = u[i];
}
}
int main(void){
int sweeps = 2000;
int N = 4096*2048;
numtype h = 1.0/N;
numtype hSquared = pow(h, 2);
NeighborAverageFunctor<numtype> op(N);
thrust::device_ptr<numtype> u_d = thrust::device_malloc<numtype>(N+2);
thrust::device_ptr<numtype> uNew_d = thrust::device_malloc<numtype>(N+2);
thrust::device_ptr<numtype> uTemp_d;
thrust::fill(u_d, u_d + (N+2), 0.0);
u_d[0] = 1.0;
boost::timer::timer timer1;
for(int k = 0; k < sweeps; k++){
thrust_sweep<numtype>(u_d, uNew_d, op);
uTemp_d = u_d;
u_d = uNew_d;
uNew_d = uTemp_d;
}
double thrust_time = timer1.elapsed();
thrust::host_vector<numtype> u_h(N+2);
thrust::copy(u_d, u_d + N+2, u_h.begin());
for(int i = 0; i < 10; i++){
std::cout << u_h[i] << " ";
}
std::cout << std::endl;
thrust::device_free(u_d);
thrust::device_free(uNew_d);
numtype * u_raw_d, * uNew_raw_d, * uTemp_raw_d;
cudaMalloc(&u_raw_d, (N+2)*sizeof(numtype));
cudaMalloc(&uNew_raw_d, (N+2)*sizeof(numtype));
numtype * u_raw_h = (numtype*)malloc((N+2)*sizeof(numtype));
int block_size = 256;
int grid_size = ((N+2) + block_size - 1) / block_size;
initialization_kernel<numtype><<<grid_size, block_size>>>(N, u_raw_d);
boost::timer::timer timer2;
for(int k = 0; k < sweeps; k++){
sweep_kernel<numtype><<<grid_size, block_size>>>(N+2, hSquared, u_raw_d, uNew_raw_d);
uTemp_raw_d = u_raw_d;
u_raw_d = uNew_raw_d;
uNew_raw_d = uTemp_raw_d;
}
double raw_time = timer2.elapsed();
cudaMemcpy(u_raw_h, u_raw_d, (N+2)*sizeof(numtype), cudaMemcpyDeviceToHost);
for(int i = 0; i < 10; i++){
std::cout << u_raw_h[i] << " ";
}
std::cout << std::endl;
std::cout << "Thrust: " << thrust_time << " s" << std::endl;
std::cout << "Raw: " << raw_time << " s" << std::endl;
free(u_raw_h);
cudaFree(u_raw_d);
cudaFree(uNew_raw_d);
return 0;
}
最佳答案
根据我的测试,这些行:
uNew[0] = u[0];
uNew[N+1] = u[N+1];
相对于内核方法,正在扼杀您的推力性能。当我消除它们时,结果似乎没有任何不同。与内核处理边界情况的方式相比,推力代码使用非常昂贵的方法(cudaMemcpy
操作,在幕后)执行边界处理。
由于推力仿函数实际上从未写入边界位置,因此只写入这些值一次就足够了,而不是在循环中写入。
通过更好地处理边界情况,您可以显着加快推力性能。
关于c++ - 简单的 Thrust 代码执行速度大约是我的原始 cuda 内核的一半。我使用推力错了吗?,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/29421482/
我需要围绕半径大约为 X 米的点 (lon,lat) 创建一个圆。 该点是通过等效于 geomFromEwkt('SRID=1;POINT(lon lat)') 生成的。 我知道 postgis 的缓
代码实现了读取文件(包含大量url)的功能,每个url都通过“evhttp_uri_parse”获取主机和路径。但是有一个错误,evhttp_uri_parse解析失败,返回NULL。可能原因是堆栈溢
所以我有两个进程,一个客户端进程,一个服务器进程。用户可以向客户端发出命令,当用户输入命令时客户端会将命令长度发送给服务器,之后再发送实际的命令。 服务器首先发回响应的长度,然后发送回响应。 我可以执
我从enwiki-latest-pagelinks.sql.gz下载了dumps.wikimedia.org/enwiki/latest/转储。 我开始将表导入到mysql数据库中: mysql -D
我有一个带有 Jw 音频播放器的 php 页面,并且有大约 5500 个链接,在每个链接的 onclick 事件上都附加了一个 javascript 函数。 php 正在生成文件的确切相对路径和名称,
我有一个大约 2GB 的巨大文本文件,我试图用 C# 解析它。该文件具有行和列的自定义分隔符。我想解析文件并提取数据并通过插入列标题并将 RowDelimiter 替换为换行符并将 ColumnDel
我已经建立了几个网站,出于某种原因,当我“喜欢”一篇博客文章或喜欢这个网站时,它们都不会再贴到我的墙上了。 (使用 iframe)示例: http://madhatterulti.com/ http:
关闭。这个问题需要更多focused .它目前不接受答案。 想改进这个问题吗? 更新问题,使其只关注一个问题 editing this post . 关闭 8 年前。 Improve this qu
我有一个与 LongPoll 一起工作的服务,当我收到我的数据时一切正常,但是当我没有收到数据时,我收到的是空结果(长轮询最大时间 == 25 秒)我的服务有时会关闭手动(我没有在服务列表中看到它)。
对于用于 Android 和应用内购买的最佳支付 API 是否达成共识? 在谷歌上搜索“Android 支付 api”,有大量来自 paypal、sms 和信用卡公司等的点击。但其中大部分文章已有多年
我的查询执行时间很长,大约 120 秒。 任何人都可以帮我重写这个查询。 请参阅下面的解释计划和表格结构。 我们经常在慢日志中收到此查询。 查询: select count(*) as col_
我正在尝试找到一种方法来对墓 map 像进行近似分割(在文化科学中的 CBIR 背景下 - 但这不是主题)。到目前为止,我正在使用这个策略: 模糊图像两次(实验结果) 应用 Canny 边缘检测器 寻
当您在 Google 中搜索时(我几乎可以肯定 Altavista 做了同样的事情),它会显示“关于 xxxx 的结果 1-10”... 这一直让我感到惊讶......“关于”是什么意思? 他们怎么能
今天我们的一台 Linux 服务器在打开出站请求时遇到问题。我已经查看了这个答案,Increasing the maximum number of tcp/ip connections in linu
我在 MVC4 站点中使用 NuGet 的最新 SignalR。使用sample hub code (或任何代码),我遇到一些奇怪的连接问题。一切加载正常,SignalR 进行协商调用并记录“Even
有人可以帮帮我吗?我的 GUI 有问题。这会过快地接收过多数据,以至于事件处理程序会完全阻塞 GUI 以供用户输入。 后台程序用于向 GUI 发送解决方案(作为文本),但 GUI 处理数据的速度不够快
假设我有一段代码,例如 for (j = 0; j 180) { c.fillStyle = 'red' } c.fi
我正在创建 2 名玩家的游戏,每个玩家都有自己的区域。我希望每个玩家都可以用一根手指触摸自己的区域。所以首先我将“达阵”功能限制为“2 个指针”。 public boolean touchDown(i
进程运行卡在 32000 (± 5%) 左右 ~# cat/proc/sys/kernel/threads-max127862 ~# ulimit -s堆栈大小(千字节,-s)2048 可用内存:3,
我有一个问题。我的应用程序在某些 Android 设备(Android 版本 5 到 6)的后台运行(如前台服务)。应用程序连接到服务器(TCP 连接)并且它们至少每 45 秒交换一次数据。 如果屏幕
我是一名优秀的程序员,十分优秀!