如何优化这个CUDA内核

问题描述

我已经对我的模型进行了分析，似乎这个内核占我的总运行时间的大约2/3。我在寻找建议来优化它。代码如下。

  __ global__ void calcFlux（double * concs，double * fluxes，double * dt）
 {
 int idx = blockIdx.x * blockDim.x + threadIdx.x; 
 fluxes [idx] = knowles_flux（idx，concs） 
 // fluxes [idx] = flux（idx，concs）; 
} 
 
 __device__ double knowles_flux（int r，double * conc）
 {
 double frag_term = 0; 
 double flux = 0; 
 if（r ==（（maxlength）-1））
 {
 //计算类型：Max
 flux = -km *（r）* conc [r ] + 2 *（ka）* conc [r-1] * conc [0] 
} 
 else if（r>（（nc）-1））
 {
 //计算类型：F
 // arrSum3（conc， & frag_term，r + 1，maxlength-1）; 
 for（int s = r + 1; s<（maxlength）; s ++）
 {
 frag_term + = conc [s] 
} 
 flux =  - （km）*（r）* conc [r] + 2 *（km）* frag_term_2 *（ka）* conc [r] * conc [0] + 2 *（ka）* conc [r-1] * conc [0]; 
} 
 else if（r ==（（nc）-1））
 {
 //计算类型：N
 // arrSum3 & frag_term，r + 1，maxlength-1）; 
 for（int s = r + 1; s<（maxlength）; s ++）
 {
 frag_term + = conc [s] 
} 
 flux =（kn）* pow（conc [0]，（nc））+ 2 *（km）* frag_term_2 *（ka）* conc [r] * conc [0] ; 
} 
 else if（r  {
 //计算类型：O
 flux = 0; 
} 
 return flux; 
} 
  

只是为了让你了解为什么for循环是一个问题，内核在大约maxlength = 9000个元素的数组上启动。为了我们现在的目的，nc在2-6的范围内。这里是一个插图说明这个内核如何处理传入数组（conc）。对于此数组，需要对不同的元素组应用五种不同类型的计算。

 数组元素：0 1 2 3 4 5 6 7 8 9 ... 8955 8956 8957 8958 8959 8960 
计算类型：MOOOOONFFF ... FFFFF Max 
  

我现在想处理的潜在问题是分支分歧从四元组if-else和for循环。

我处理分支分歧的想法是将这个内核分成四个独立的设备函数或内核，分别处理每个区域，并同时启动。我不知道这是明显好于只是让分支发生发生，如果我没有错误，会导致四种计算类型串行运行。

要处理for循环，您会注意到有一个注释掉的arrSum3函数，我写的基于我以前（也可能很差）写并行还原内核。使用它代替for循环显着增加了我的运行时。我觉得有一个聪明的方式来完成我试图用for循环，但我不是那么聪明，我的顾问厌倦了我浪费时间的思考。

感谢任何帮助。

EDIT >完整代码位于此处：