- iOS/Objective-C 元类和类别
- objective-c - -1001 错误,当 NSURLSession 通过 httpproxy 和/etc/hosts
- java - 使用网络类获取 url 地址
- ios - 推送通知中不播放声音
第一次来这里提问。我在我的 CUDA 程序中遇到了一些问题。
我有数百万个四面体,一个点在 (0,0,0),所以我可以使用公式:
获取四面体的体积。
所以,这是代码:
struct Triangle
{
double x1;
double y1;
double z1;
double x2;
double y2;
double z2;
double x3;
double y3;
double z3;
};
和 CUDA 代码:
__global__ void getResult(double *d_volume ,Triangle *d_triangles, Origin *d_point)
{
extern __shared__ Triangle s_data[];
int tid = threadIdx.x;
int i = blockDim.x * blockIdx.x + threadIdx.x;
s_data[tid] = d_triangles[i];
__syncthreads();
d_volume[i] =s_data[tid].x1 * s_data[tid].y2 * s_data[tid].z3 + \
s_data[tid].y1 * s_data[tid].z2 * s_data[tid].x3 + \
s_data[tid].x2 * s_data[tid].y3 * s_data[tid].z1 - \
s_data[tid].x3 * s_data[tid].y2 * s_data[tid].z1 - \
s_data[tid].x2 * s_data[tid].y1 * s_data[tid].z3 - \
s_data[tid].y3 * s_data[tid].z2 * s_data[tid].x1;
}
我从其他函数中得到了数百万个四面体作为数组。
// Host
Triangle *h_triangles = triangles;
double *h_volume;
// Device
Triangle *d_triangles;
double *d_volume;
// define grid and block size
int numThreadsPerBlock = numTriangles;
int numBlocks = numTrianges / 512;
// Shard memory size
int sharedMemSize = numThreadsPerBlock * sizeof(Triangle);
// allocate host and device memory
size_t memSize_triangles = numBlocks * numThreadsPerBlock * sizeof(Triangle);
size_t memSize_volume = numBlocks * numThreadsPerBlock * sizeof(double);
cudaMalloc( (void **) &d_triangles, memSize_triangles );
cudaMalloc( (void **) &d_volume, memSize_volume );
// Copy host array to device array
cudaMemcpy( d_triangles, h_triangles, memSize_triangles, cudaMemcpyHostToDevice );
cudaMemcpy( d_point, h_point, memSize_point, cudaMemcpyHostToDevice );
// launch kernel
dim3 dimGrid(numBlocks);
dim3 dimBlock(numThreadsPerBlock);
getResult<<< dimGrid, dimBlock, sharedMemSize >>>( d_volume, d_triangles);
// block until the device has completed
cudaThreadSynchronize();
// device to host copy
cudaMemcpy( h_volume, d_volume, memSize_volume, cudaMemcpyDeviceToHost );
// free device memory
cudaFree(d_triangles);
cudaFree(d_volume);
// free host memory
free(h_triangles);
free(h_volume);
到目前为止,一切正常。但是我花了比我想象的更多的时间来获得音量。我的设备是 Tesla C2050(515Gflops),比我的 CPU(单核,20.25Gflops)快 20 倍。但只提速10倍左右(不包括设备和主机之间复制内存的时间。)
我想知道如何才能使它比 CPU 代码快 20 倍左右(for 循环以获取音量。)。
谢谢!
PS:也许 cudaMallocPitch() 会帮助我,但三角形不是矩阵,我不能使用 cudaMemcpy2D() 而不是 cudaMemcpy() 来复制内存。谁能帮我解决这个问题?
最佳答案
与 CPU 相比,GPU 上的峰值性能通常更难获得。原因之一是许多内核受带宽限制而不是计算限制。
因为你的内核的计算复杂度是 O(n)。您可能应该使用带宽指标来计算理论峰值性能,如下所示
1024*1024*64 * sizeof(double) * (9 + 1) / (144e9 * 8/9) = 42 ms
#tetrahedron #input #output peak mem bw ECC cost
另一方面,您的内核可以进一步优化。
__syncthreads()
可能会被消除。获得了一个具有大型 L1 缓存设置和最佳 blockDim/gridDim 选择的新内核。它快了 15%。这是代码和配置文件结果。我的设备是 M2090。
#include <stdlib.h>
#include <thrust/transform.h>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <iterator>
#include <thrust/inner_product.h>
using namespace thrust::placeholders;
struct Triangle
{
double x1;
double y1;
double z1;
double x2;
double y2;
double z2;
double x3;
double y3;
double z3;
};
__global__ void getResultNoSMem(double *d_volume, Triangle *d_triangles)
{
int i = blockDim.x * blockIdx.x + threadIdx.x;
d_volume[i] = d_triangles[i].x1 * d_triangles[i].y2 * d_triangles[i].z3 +
d_triangles[i].y1 * d_triangles[i].z2 * d_triangles[i].x3 +
d_triangles[i].x2 * d_triangles[i].y3 * d_triangles[i].z1 -
d_triangles[i].x3 * d_triangles[i].y2 * d_triangles[i].z1 -
d_triangles[i].x2 * d_triangles[i].y1 * d_triangles[i].z3 -
d_triangles[i].y3 * d_triangles[i].z2 * d_triangles[i].x1;
}
__global__ void getResult(double *d_volume, Triangle *d_triangles)
{
extern __shared__ Triangle s_data[];
int tid = threadIdx.x;
int i = blockDim.x * blockIdx.x + threadIdx.x;
s_data[tid] = d_triangles[i];
// __syncthreads();
d_volume[i] = s_data[tid].x1 * s_data[tid].y2 * s_data[tid].z3 +
s_data[tid].y1 * s_data[tid].z2 * s_data[tid].x3 +
s_data[tid].x2 * s_data[tid].y3 * s_data[tid].z1 -
s_data[tid].x3 * s_data[tid].y2 * s_data[tid].z1 -
s_data[tid].x2 * s_data[tid].y1 * s_data[tid].z3 -
s_data[tid].y3 * s_data[tid].z2 * s_data[tid].x1;
}
__global__ void getResultOpt(double *d_volume, Triangle *d_triangles, int len)
{
const int gridSize = blockDim.x * gridDim.x;
int i = blockDim.x * blockIdx.x + threadIdx.x;
while (i < len)
{
d_volume[i] = d_triangles[i].x1 * d_triangles[i].y2 * d_triangles[i].z3 +
d_triangles[i].y1 * d_triangles[i].z2 * d_triangles[i].x3 +
d_triangles[i].x2 * d_triangles[i].y3 * d_triangles[i].z1 -
d_triangles[i].x3 * d_triangles[i].y2 * d_triangles[i].z1 -
d_triangles[i].x2 * d_triangles[i].y1 * d_triangles[i].z3 -
d_triangles[i].y3 * d_triangles[i].z2 * d_triangles[i].x1;
i += gridSize;
}
}
int main(void)
{
const int m = 1024 * 1024;
thrust::host_vector<Triangle> data(m);
for (int i = 0; i < m; i++)
{
data[i].x1 = (double) rand() / RAND_MAX;
data[i].y1 = (double) rand() / RAND_MAX;
data[i].z1 = (double) rand() / RAND_MAX;
data[i].x2 = (double) rand() / RAND_MAX;
data[i].y2 = (double) rand() / RAND_MAX;
data[i].z2 = (double) rand() / RAND_MAX;
data[i].x3 = (double) rand() / RAND_MAX;
data[i].y3 = (double) rand() / RAND_MAX;
data[i].z3 = (double) rand() / RAND_MAX;
}
thrust::device_vector<Triangle> triangles = data;
thrust::device_vector<double> volume(m);
thrust::device_vector<double> volumeOpt(m);
Triangle* dTriangles = thrust::raw_pointer_cast(&triangles[0]);
double* dVolume = thrust::raw_pointer_cast(&volume[0]);
double* dVolumeOpt = thrust::raw_pointer_cast(&volumeOpt[0]);
int g;
int b;
int threadUpperLimit = 48 * 1024 / sizeof(Triangle);
//for (b = 32; b <= 1024; b += 32)
{
b = 64;
int gridDim = (m + b - 1) / b;
getResultNoSMem<<<gridDim, b, 0, 0>>>(dVolume, dTriangles);
}
// for (b = 32; b <= threadUpperLimit; b += 32)
{
b = 64;
int gridDim = (m + b - 1) / b;
getResult<<<gridDim, b, b * sizeof(Triangle), 0>>>(dVolume, dTriangles);
}
//for (g = 32; g <= 512; g += 32)
// for (b = 32; b <= 1024; b += 32)
{
b = 64;
g = 64;
getResultOpt<<<g, b, 0, 0>>>(dVolumeOpt, dTriangles, m);
}
//for (g = 32; g <= 512; g += 32)
// for (b = 32; b <= 1024; b += 32)
{
b = 64;
g = 512;
cudaFuncSetCacheConfig(getResultOpt, cudaFuncCachePreferL1);
getResultOpt<<<g, b, 0, 0>>>(dVolumeOpt, dTriangles, m);
}
thrust::device_vector<double> X = volume;
thrust::device_vector<double> Y = volumeOpt;
thrust::transform(X.begin(), X.end(), Y.begin(), X.begin(), _1 - _2);
double result = thrust::inner_product(X.begin(), X.end(), X.begin(), 0.0);
std::cout << "difference: " << result << std::endl;
return 0;
}
关于c++ - CUDA并行计算加速体积计算,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/18480975/
有没有办法同时运行 2 个不同的代码块。我一直在研究 R 中的并行包,它们似乎都基于在循环中运行相同的函数。我正在寻找一种同时运行不同函数的方法(循环的 1 次迭代)。例如,我想在某个数据对象上创建一
无论如何增加 Parallel.For 启动后的循环次数?示例如下: var start = 0; var end = 5; Parallel.For(start, end, i => { C
我是 Golang 的新手,正在尝试了解并发和并行。我阅读了下面提到的关于并发和并行的文章。我执行了相同的程序。但没有得到相同的(混合字母和字符)输出。首先获取所有字母,然后获取字符。似乎并发不工作,
我正在寻找同时迭代 R 中两个或多个字符向量/列表的方法,例如。有没有办法做这样的事情: foo <- c('a','c','d') bar <- c('aa','cc','dd') for(i in
我对 Raku 很陌生,我对函数式方法有疑问,尤其是 reduce。 我最初有这样的方法: sub standardab{ my $mittel = mittel(@_); my $foo =
我最近花了很多时间来学习实时音频处理的细节,我发现的大多数库/工具都是c / c++代码或脚本/图形语言的形式,并在其中编译了c / c++代码。引擎盖。 使用基于回调的API,与GUI或App中的其
我正在使用 JMeter 进行图像负载测试。我有一个图像名称数组并遍历该数组,我通过 HTTP 请求获取所有图像。 -> loop_over_image - for loop controller
我整个晚上都在困惑这个问题...... makeflags = ['--prefix=/usr','--libdir=/usr/lib'] rootdir='/tmp/project' ps = se
我正在尝试提高计算图像平均值的方法的性能。 为此,我使用了两个 For 语句来迭代所有图像,因此我尝试使用一个 Parallel For 来改进它,但结果并不相同。 我做错了吗?或者是什么导致了差异?
假设您有一个并行 for 循环实现,例如ConcRT parallel_for,将所有工作放在一个 for 循环体内总是最好的吗? 举个例子: for(size_t i = 0; i < size()
我想并行运行一部分代码。目前我正在使用 Parallel.For 如何让10、20或40个线程同时运行 我当前的代码是: Parallel.For(1, total, (ii) =>
我使用 PAY API 进行了 PayPal 自适应并行支付,其中无论用户(买家)购买什么,都假设用户购买了总计 100 美元的商品。在我的自适应并行支付中,有 2 个接收方:Receiver1 和
我正在考虑让玩家加入游戏的高效算法。由于会有大量玩家,因此算法应该是异步的(即可扩展到集群中任意数量的机器)。有细节:想象有一个无向图(每个节点都是一个玩家)。玩家之间的每条边意味着玩家可以参加同一场
我有一个全局变量 volatile i = 0; 和两个线程。每个都执行以下操作: i++; System.out.print(i); 我收到以下组合。 12、21 和 22。 我理解为什么我没有得到
我有以下称为 pgain 的方法,它调用我试图并行化的方法 dist: /***************************************************************
我有一个 ruby 脚本读取一个巨大的表(约 2000 万行),进行一些处理并将其提供给 Solr 用于索引目的。这一直是我们流程中的一大瓶颈。我打算在这里加快速度,我想实现某种并行性。我对 Ru
我正在研究 Golang 并遇到一个问题,我已经研究了几天,我似乎无法理解 go routines 的概念以及它们的使用方式。 基本上我是在尝试生成数百万条随机记录。我有生成随机数据的函数,并将创建一
我希望 for 循环使用 go 例程并行。我尝试使用 channel ,但没有用。我的主要问题是,我想在继续之前等待所有迭代完成。这就是为什么在它不起作用之前简单地编写 go 的原因。我尝试使用 ch
我正在使用 import Control.Concurrent.ParallelIO.Global main = parallel_ (map processI [1..(sdNumber runPa
我正在尝试通过 makePSOCKcluster 连接到另一台计算机: library(parallel) cl ... doTryCatch -> recvData -> makeSOCKm
我是一名优秀的程序员,十分优秀!