爲什麼CPU和GPU之間的結果不一樣？

Question

這是在GPU上運行

tid=threadidx%x 
bid=blockidx%x 
bdim=blockdim%x 

isec = mesh_sec_1(lev)+bid-1 
if (isec .le. mesh_sec_0(lev)) then 
    if(.not. sec_is_int(isec)) return 

    do iele = tid, sec_n_ele(isec), bdim 

     idx = n_ele_idx(isec)+iele 

     u(1:5) = fv_u(1:5,idx) 
     u(6 ) = fv_t( idx) 
     g  = 0.0d0 
     do j= sec_iA_ls(idx), sec_iA_ls(idx+1)-1 
      ss = sec_jA_ls(1,j) 
      ee = sec_jA_ls(2,j) 
      tem = n_ele_idx(ss)+ee 
      du(1:5) = fv_u(1:5, n_ele_idx(ss)+ee)-u(1:5) 
      du(6 ) = fv_t( n_ele_idx(ss)+ee)-u(6 ) 
      coe(1:3) = sec_coe_ls(1:3,j) 
      do k=1,6 
       g(1:3,k)=g(1:3,k)+du(k)*sec_coe_ls(1:3,j) 
      end do 
     end do 
     do j=1,6 
     do i=1,3 
      fv_gra(i+(j-1)*3,idx)=g(i,j) 
     end do 
     end do 
    end do 
end if 

和明年我的代碼是CPU上運行

do isec = h_mesh_sec_1(lev),h_mesh_sec_0(lev) 
    if(.not. h_sec_is_int(isec)) cycle 
    do iele=1,h_sec_n_ele(isec) 

     idx = h_n_ele_idx(isec)+iele 

     u(1:5) = h_fv_u(1:5,idx) 
     u(6 ) = h_fv_t( idx) 
     g  = 0.0d0 
     do j= h_sec_iA_ls(idx),h_sec_iA_ls(idx+1)-1 
      ss = h_sec_jA_ls(1,j) 
      ee = h_sec_jA_ls(2,j) 
      du(1:5) = h_fv_u(1:5,h_n_ele_idx(ss)+ee)-u(1:5) 
      du(6 ) = h_fv_t( h_n_ele_idx(ss)+ee)-u(6 ) 
      do k=1,6 
       g(1:3,k)= g(1:3,k) + du(k)*h_sec_coe_ls(1:3,j) 
      end do 
     end do 
     do j=1,6 
     do i=1,3 
      h_fv_gra(i+(j-1)*3,idx) = g(i,j) 
     enddo 
     enddo 

    end do 
end do 

可變我的代碼H_之間*和*表明它屬於CPU和GPU分別。 結果在許多方面都是一樣的，但在某些方面它們有點不同。我像這樣添加校驗碼。

do i =1,size(h_fv_gra,1) 
    do j = 1,size(h_fv_gra,2) 
     if(hd_fv_gra(i,j)-h_fv_gra(i,j) .ge. 1.0d-9) then 
      print *,hd_fv_gra(i,j)-h_fv_gra(i,j),i,j 
     end if 
    end do 
end do 

hd_ *是gpu結果的副本。我們可以看到的區別：

1.8626451492309570E-009 13 14306

1.8626451492309570E-009 13 14465

1.8626451492309570E-009 13 14472

1.8626451492309570E-009 14 14128

1.8626451492309570E-009 14 14146

1.8626451492309570E-009 14 14150

1.8626451492309570E-009 14 14153

1.8626451492309570E-009 14 14155

1.8626451492309570E-009 14 14156

所以我很困惑有關。 Cuda的精度不應該如此大。任何答覆都會受到歡迎。 另外，我不知道如何在GPU代碼中打印變量，這可以幫助我進行調試。

Answer 1

在您的代碼中，計算g值最可能受益於CUDA中的Fused Multiply Add（fma）優化。

g(1:3,k)=g(1:3,k)+du(k)*sec_coe_ls(1:3,j) 

在CPU方面，這不是不可能的，但在很大程度上依賴於編譯器選項（和運行代碼的CPU實際應該把它落實FMA）。

要強制單獨乘的使用和添加，要使用內部函數從CUDA，定義here，如：

__device__ ​ double __dadd_rn (double x, double y)在舍入到最近的偶數模式添加兩個浮點值。

和

__device__ ​ double __dmul_rn (double x, double y)乘法中舍入到最近的偶數模式下的兩個浮點值。

舍入模式與CPU上定義的舍入模式相同（取決於CPU架構，無論是Power還是Intel x86或其他）。

另一種方法是在編譯cuda時，使用選項（從nvcc詳述here）將--fmad false選項傳遞到ptxas。