我正在写一个n体模拟,基本上整个操作是:
-Prepare CUDA memory
loop {
-Copy data to CUDA
-Launch kernel
-Copy data to host
-Operations using data (drawing etc.)
}
我注意到几乎 90% 的时间都花在了将数据写入内核中的全局设备内存上。这是内核:
__global__ void calculateForcesCuda(float *deviceXpos, float *deviceYpos, float *deviceZpos,
float *deviceXforces, float *deviceYforces, float *deviceZforces,
float *deviceMasses, int particlesNumber) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid <= particlesNumber) {
float particleXpos = deviceXpos[tid];
float particleYpos = deviceYpos[tid];
float particleZpos = deviceZpos[tid];
float xForce = 0.0f;
float yForce = 0.0f;
float zForce = 0.0f;
for (int index=0; index<particlesNumber; index++) {
if (tid != index) {
float otherXpos = deviceXpos[index];
float otherYpos = deviceYpos[index];
float otherZpos = deviceZpos[index];
float mass = deviceMasses[index];
float distx = particleXpos - otherXpos;
float disty = particleYpos - otherYpos;
float distz = particleZpos - otherZpos;
float distance = sqrt((distx*distx + disty*disty + distz*distz) + 0.01f);
xForce += 10.0f * mass / distance * (otherXpos - particleXpos);
yForce += 10.0f * mass / distance * (otherYpos - particleYpos);
zForce += 10.0f * mass / distance * (otherZpos - particleZpos);
}
}
deviceXforces[tid] += xForce;
deviceYforces[tid] += yForce;
deviceZforces[tid] += zForce;
}
}
运行它的设备是 GTX 970。执行时间约为 8.0 秒,但是在添加这些标志后:-gencode arch=compute_52,code=sm_52,性能提升到6.7秒左右。注释掉写入全局设备内存的代码后:
deviceXforces[tid] += xForce;
deviceYforces[tid] += yForce;
deviceZforces[tid] += zForce;
... 总执行时间减少到 0.92 秒左右,这意味着写入全局设备内存大约占执行时间的 86%。有什么方法可以提高这些写入的性能?
最佳答案
内存通常是这种计算的瓶颈,即使它没有像您测量的那样占用 90% 的时间。我会建议两件事。
将设备...[索引]
加载到共享内存中
就目前而言,所有线程都读取相同的 deviceXpos[index]
、deviceYpos[index]
、deviceZpos[index]
和 deviceMasses[索引]
。相反,您可以将它们加载到共享内存中:
static const int blockSize = ....;
__shared__ float shXpos[blockSize];
__shared__ float shYpos[blockSize];
__shared__ float shZpos[blockSize];
__shared__ float shMasses[blockSize];
for (int mainIndex=0; mainIndex<particlesNumber; index+=blockSize) {
__syncthreads(); //ensure computation from previous iteration has completed
shXpos[threadIdx.x] = deviceXpos[mainIndex + threadIdx.x];
shYpos[threadIdx.x] = deviceYpos[mainIndex + threadIdx.x];
shZpos[threadIdx.x] = deviceZpos[mainIndex + threadIdx.x];
shMasses[threadIdx.x] = deviceMasses[mainIndex + threadIdx.x];
__syncthreads(); //ensure all data is read before computation starts
for (int index=0; index<blockSize; ++index) {
.... //your computation, using sh....[index] values
}
}
这应该减少全局内存读取量,因为每个线程读取不同的数据,而不是所有线程都读取相同的东西。
但是请注意,如果驱动程序正确管理 L1 缓存,则此建议可能不会那么有效。试试吧!
每个线程处理超过 1 个(接收)粒子
您可能希望一次对多个粒子执行计算。您可以使用一组 {particleX/Y/Zpos
, x/y/zForce
} 来代表单个粒子接受力,而不是只使用一组那些同时。
这样,通过在循环中加载一次源,您可以处理多个接收器。
这可能会显着降低您的内存压力,但同时会增加您的寄存器数量。寄存器太多 - 您将无法启动那么多线程。
检查您的线程已有多少个寄存器,并查阅 CUDA 占用计算器以了解您还可以使用多少。也许将占用率从 1 降低到 0.5 或 0.75,但同时处理更多的粒子会有好处吗?您将需要进行试验,因为这可能因 GPU 而异。
关于c++ - CUDA:有没有更快的写入全局内存的方法?,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/38877670/