我最近使用 Cuda 实现了 Successive Over Relaxation 作为我的课程项目的一部分，我很想知道如何使代码更高效。
我正在使用红/黑 SOR 方案，这是并行 SOR 的流行版本。要更新网格点的值，我需要访问 4 个相邻的值，这些值在内存中通常不连续，因此 GPU 无法使用 cuda 的内存合并功能，这使得算法内存受限。为了克服这个问题，我将奇数和偶数网格点的存储重新排序在单独的向量中，如本文所述：http: //link.springer.com/chapter/10.1007%2F978-3-642-31464-3_60这需要内存合并的优势。
实施论文中提到的技术，我发现与 1 个 CPU 相比，速度提高了约 15 倍（与传统的红/黑 SOR 相比约 10 倍）。我应该期待多少性能加速？我觉得〜15的加速非常低？
我正在调用内核函数以以下方式从主机 CPU 依次更新奇数值和偶数值：
for(size_t it = 0; it < itmax; it++){
cuda_odd_update <<<dimGrid_odd,dimBlock>>> (odd,even);
cuda_even_update <<<dimGrid_even,dimBlock>>> (odd,even);
}
其中odd和even是指向设备内存中重新排序的向量的指针。我正在使用这种方法，因为我不知道有什么方法可以同步设备中的所有线程。有没有更好的方法来实现执行循环？
下面的代码实现了论文中提到的技术，而不是传统的红/黑 SOR：
__global__ void cuda_odd_update(double* odd,double* even){
size_t tx = blockIdx.x*blockDim.x + threadIdx.x;
size_t ty = blockIdx.y*blockDim.y + threadIdx.y;
size_t odd_index = tx*height_odd+ty;
size_t even_index = tx*height_even+ty;
if (( (ty == 0 && tx%2 == 0) || (ty == height_odd-1 && tx%2 == 1) || (ty > 0 && ty < height_odd-1) ) && (tx > 0 && tx < width-1 && ty < height_odd) ){
odd[odd_index] = (1.0 - omega)*odd[odd_index] + omega/(2*(1+beta))
* ( (tx%2)*even[even_index]+(1-tx%2)*even[even_index+1]
+ (tx%2)*even[even_index-1]+(1-tx%2)*even[even_index]
+ beta * ( even[even_index-height_even]
+ even[even_index+height_even] )) ;
}
}
__global__ void cuda_even_update(double* odd,double* even){
size_t tx = blockIdx.x*blockDim.x + threadIdx.x;
size_t ty = blockIdx.y*blockDim.y + threadIdx.y;
size_t odd_index = tx*height_odd+ty;
size_t even_index = tx*height_even+ty;
if(( (ty == 0 & tx%2 == 1) || (ty == height_even-1 && tx%2 == 0) || (ty > 0 && ty < height_even-1) ) && ( tx > 0 && tx < width-1 && ty < height_even)){
even[even_index] = (1.0 - omega)*even[even_index] + omega/(2*(1+beta))
* ( (1-tx%2)*odd[odd_index] + (tx%2)*odd[odd_index+1]
+ (1-tx%2)*odd[odd_index-1] + (tx%2)*odd[odd_index]
+ beta * ( odd[odd_index-height_odd]
+ odd[odd_index+height_odd] )) ;
}
}
我目前正在使用 Intel Xeon CPU E5-2660-2.20GHz（8 核）和 NVIDIA TESLA M2090。给出以下nvprof结果（1000 次迭代）：
CPU time : 62.69
==27665== NVPROF is profiling process 27665, command: ./sor
GPU time : 3.73
End!
==27665== Profiling application: ./sor
==27665== Profiling result:
Time(%) Time Calls Avg Min Max Name
49.81% 1.83752s 1000 1.8375ms 1.8333ms 1.8422ms cuda_even_update(double*, double*)
49.67% 1.83226s 1000 1.8323ms 1.8290ms 1.8364ms cuda_odd_update(double*, double*)
0.26% 9.6281ms 2 4.8141ms 4.6306ms 4.9975ms [CUDA memcpy DtoH]
0.26% 9.4569ms 2 4.7284ms 4.7065ms 4.7504ms [CUDA memcpy HtoD]