在 CUDA 中计算 CGMA答案

【问题标题】：Calculation of CGMA in CUDA在 CUDA 中计算 CGMA
【发布时间】：2013-10-15 20:56:06
【问题描述】：

我对 CGMA 的计算感到困惑。我知道 CGMA = # of operations / # of memory fetches。首先，当x = g_A[idx]时，我应该将写操作计入x还是因为它存储在寄存器中而忽略它？同样，在z = (x*y) + (y/x) + (y-x); 中，我是否应该将x 和y 的读取计算为计算CGMA 时的内存读取？最后，我是不是应该把Kernel函数中的所有操作（那五行）都数一数？

__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;
    }
}

【问题讨论】：

对寄存器的读取和写入不计入 GMA。 2 从 g_A 和 g_B 读取计数。写入g_C 很重要。为了使算术计算正确，您需要查看 SASS 代码，或者您可以简单地从源代码或 PTX 代码进行估计。您应该计算所有算术运算，包括与idx相关的算术运算
<运算符算不算操作？
是的。在内核中执行的任何算术都算数（例如，即使是非显而易见的地址计算算术）。
@RobertCrovella CGMA 似乎被定义为每次访问 CUDA 程序区域内的全局内存时执行的浮点计算数。 < 操作符好像被反汇编成ISETP，也就是整数比较，所以我说应该不算吧？我看到的唯一浮点操作是那些与反汇编代码中的x/y 相关的操作。
可能你是对的。我正在考虑分析驱动的优化，并试图发现内核是计算受限还是内存受限。在这种情况下，我们要考虑在内核中完成的所有计算操作。

标签： c cuda

【解决方案1】：

你的内核对应的反汇编代码（编译为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 转换算作浮点运算）。

【讨论】：

您为哪种架构编译？有趣的是它没有整数除法硬件。
@RogerDahl Fermi：compute_20,sm_20.

【解决方案2】：

看起来CGMA代表Compute to Global Memory Access，定义为数字每次访问全局内存时执行的浮点计算 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

【讨论】：