CUDA中CGMA的計算

我對CGMA的計算感到困惑。我知道CGMA =運算次數/內存提取次數。首先，當x = g_A[idx]時，我應該將寫操作計數到x還是忽略它，因爲它存儲在寄存器中？同樣，在z = (x*y) + (y/x) + (y-x);中，我應該將讀取的x和y作爲計算CGMA中的內存讀數嗎？最後，我是否應該計算內核函數中的所有操作（這五行）？CUDA中CGMA的計算

__global__ void PerformSomeOperations(int* g_A,int* g_B,int* g_C, int Size) 
{ 
    const int idx = threadIdx.x + (blockIdx.x*blockDim.x); 
    if(idx < Size) 
    { 
     int x = g_A[idx]; 
     int y = g_B[idx]; 
     int z = 0; 
     z = (x*y) + (y/x) + (y-x); 
     g_C[idx] = z; 
    } 
}

來源

2013-10-15 Shibli

讀取和寫入寄存器不計入GMA。 2從'g_A'和'g_B'計數中讀取。寫入'g_C'是很重要的。爲了使算術計算正確，您需要查看SASS代碼，或者您可以簡單地從源代碼或PTX代碼進行估算。您應該計算所有算術運算，包括那些與'idx'有關的運算操作 –

操作符是否被算作操作？ – Shibli

是的。在內核中執行的任何算術計數（例如甚至不明顯的地址計算算法）。 –

貌似CGMA代表計算全局存儲器訪問和被定義爲用於一個CUDA程序的區域內的每個訪問全局存儲器執行浮點計算的數量。

計算比率的最佳方法是在CUDA探查器中運行程序，並使用性能計數器訪問內存和浮點操作。根據我發現的定義，你的內核的CGMA爲零，因爲它執行的是整數運算，而不是浮點運算。如果更改了定義，那麼x = g_A[idx]是一個讀取操作並且沒有寫入操作。這是因爲寄存器文件沒有存儲在全局存儲器中（CGMA中的「G」）。 z = (x*y) + (y/x) + (y-x);中沒有全局內存讀取，因此將其計爲5次操作。如果所有線程都以idx < Size運行，那麼您有3個全局內存訪問和8個操作。但請注意，在CUDA中，全局內存訪問的性能取決於它們是否合併。許多合併的內存訪問的運行速度可能比少數幾個未聚合的訪問速度快得多。所以CGMA不一定能夠準確地描述內核的性能潛力。

參考文獻：

http://www.greatlakesconsortium.org/events/GPUMulticore/Chapter4-CudaMemoryModel.pdf

http://cs.nyu.edu/courses/spring12/CSCI-GA.3033-012/lecture6.pdf

來源

2013-10-15 04:10:30

對應你的內核（編譯爲compute_20,sm_20）的反彙編代碼如下

/*0000*/  MOV R1, c[0x1][0x100];      
/*0008*/  S2R R0, SR_CTAID.X;       
/*0010*/  S2R R2, SR_TID.X;        
/*0018*/  IMAD R0, R0, c[0x0][0x8], R2;    
/*0020*/  ISETP.GE.AND P0, PT, R0, c[0x0][0x2c], PT; 
/*0028*/ @P0 EXIT ;          
/*0030*/  SHL R0, R0, 0x2;       
/*0038*/  IADD R2, R0, c[0x0][0x20];     
/*0040*/  IADD R3, R0, c[0x0][0x24];     
/*0048*/  IADD R0, R0, c[0x0][0x28];     
/*0050*/  LD R2, [R2];        R2 = x = g_A[idx] 
/*0058*/  LD R3, [R3];        R3 = y = g_B[idx] 
/*0060*/  I2I.S32.S32 R5, |R2|;      
/*0068*/  I2F.F32.U32.RP R4, R5;      R4 = (float)x 
/*0070*/  MUFU.RCP R4, R4;       R4 = 1/R4 
/*0078*/  IADD32I R4, R4, 0xffffffe;     
/*0080*/  F2I.FTZ.U32.F32.TRUNC R4, R4;    
/*0088*/  IMUL.U32.U32 R6, R5, R4;     R6 = x * (1/y) 
/*0090*/  I2I.S32.S32 R7, -R6;      
/*0098*/  I2I.S32.S32 R6, |R3|;      
/*00a0*/  IMAD.U32.U32.HI R7, R4, R7, R4;    
/*00a8*/  IMUL.U32.U32.HI R4, R7, R6;    
/*00b0*/  LOP.XOR R7, R3, R2;       
/*00b8*/  IMAD.U32.U32 R6, -R5, R4, R6;    
/*00c0*/  ISETP.GE.AND P1, PT, R7, RZ, PT;   
/*00c8*/  ISETP.LE.U32.AND P0, PT, R5, R6, PT;  
/*00d0*/ @P0 ISUB R6, R6, R5;       
/*00d8*/ @P0 IADD R4, R4, 0x1;       
/*00e0*/  ISETP.GE.U32.AND P0, PT, R6, R5, PT;  
/*00e8*/  LOP.PASS_B R6, RZ, ~R2;      
/*00f0*/  ISUB R5, R3, R2;       
/*00f8*/ @P0 IADD R4, R4, 0x1;       
/*0100*/ @!P1 I2I.S32.S32 R4, -R4;      
/*0108*/  ICMP.EQ R4, R6, R4, R2;      
/*0110*/  IADD R4, R5, R4;       
/*0118*/  IMAD R2, R3, R2, R4;      
/*0120*/  ST [R0], R2;        
/*0128*/  EXIT ;

從上面的代碼，也有遵循浮點運算

I2F.F32.U32.RP R4, R5;      Integer to Float conversion 
MUFU.RCP R4, R4;       Multifunction Floating Point Operation (Reciprocal) 
F2I.FTZ.U32.F32.TRUNC R4, R4;    Float to Integer conversion

這些操作看起來與(x/y)有關，它是兩個整數之間的分界，但需要轉換爲浮點。我並沒有真正意識到轉換是否被算作浮點運算。我在代碼中看不到任何其他浮點操作。

全局存儲器操作如下所示3

LD R2, [R2];        
LD R3, [R3];        
ST [R0], R2;

我要說的是CGMA = 3/3 = 1對於你的情況（計數int2float和float2int轉換爲浮點運算）。

來源

2013-10-15 10:11:23 JackOLantern

你編譯了哪個體系結構？有趣的是，它沒有整數分頻硬件。 –

@RogerDahl Fermi：'compute_20，sm_20'。 – JackOLantern

CUDA中CGMA的計算

回答

相關問題