6

我想通过 BS_x*BS_Y 线程将内容移动到共享内存来读取 (BS_X+1)*(BS_Y+1) 全局内存位置，并且我开发了以下代码。

int i       = threadIdx.x;
int j       = threadIdx.y;
int idx     = blockIdx.x*BLOCK_SIZE_X + threadIdx.x;
int idy     = blockIdx.y*BLOCK_SIZE_Y + threadIdx.y;

int index1  = j*BLOCK_SIZE_Y+i;

int i1      = (index1)%(BLOCK_SIZE_X+1);
int j1      = (index1)/(BLOCK_SIZE_Y+1);

int i2      = (BLOCK_SIZE_X*BLOCK_SIZE_Y+index1)%(BLOCK_SIZE_X+1);
int j2      = (BLOCK_SIZE_X*BLOCK_SIZE_Y+index1)/(BLOCK_SIZE_Y+1);

__shared__ double Ezx_h_shared_ext[BLOCK_SIZE_X+1][BLOCK_SIZE_Y+1];     

Ezx_h_shared_ext[i1][j1]=Ezx_h[(blockIdx.y*BLOCK_SIZE_Y+j1)*xdim+(blockIdx.x*BLOCK_SIZE_X+i1)];

if ((i2<(BLOCK_SIZE_X+1))&&(j2<(BLOCK_SIZE_Y+1))) 
Ezx_h_shared_ext[i2][j2]=Ezx_h[(blockIdx.y*BLOCK_SIZE_Y+j2)*xdim+(blockIdx.x*BLOCK_SIZE_X+i2)];

在我的理解中，合并是顺序处理的连续内存读取的并行等价物。我现在如何检测全局内存访问是否已合并？我注意到有一个从 (i1,j1) 到 (i2,j2) 的索引跳转。提前致谢。

4

3 回答 3

5

我已经使用手写的合并分析器评估了您的代码的内存访问。评估显示代码较少利用合并。这是您可能会发现有用的合并分析器：

#include <stdio.h>
#include <malloc.h>

typedef struct dim3_t{
    int x;
    int y;
} dim3;


// KERNEL LAUNCH PARAMETERS
#define GRIDDIMX 4
#define GRIDDIMY 4
#define BLOCKDIMX 16
#define BLOCKDIMY 16


// ARCHITECTURE DEPENDENT
// number of threads aggregated for coalescing
#define COALESCINGWIDTH 32
// number of bytes in one coalesced transaction
#define CACHEBLOCKSIZE 128
#define CACHE_BLOCK_ADDR(addr,size)  (addr*size)&(~(CACHEBLOCKSIZE-1))


int main(){
    // fixed dim3 variables
    // grid and block size
    dim3 blockDim,gridDim;
    blockDim.x=BLOCKDIMX;
    blockDim.y=BLOCKDIMY;
    gridDim.x=GRIDDIMX;
    gridDim.y=GRIDDIMY;

    // counters
    int unq_accesses=0;
    int *unq_addr=(int*)malloc(sizeof(int)*COALESCINGWIDTH);
    int total_unq_accesses=0;

    // iter over total number of threads
    // and count the number of memory requests (the coalesced requests)
    int I, II, III;
    for(I=0; I<GRIDDIMX*GRIDDIMY; I++){
        dim3 blockIdx;
        blockIdx.x = I%GRIDDIMX;
        blockIdx.y = I/GRIDDIMX;
        for(II=0; II<BLOCKDIMX*BLOCKDIMY; II++){
            if(II%COALESCINGWIDTH==0){
                // new coalescing bunch
                total_unq_accesses+=unq_accesses;
                unq_accesses=0;
            }
            dim3 threadIdx;
            threadIdx.x=II%BLOCKDIMX;
            threadIdx.y=II/BLOCKDIMX;

            ////////////////////////////////////////////////////////
            // Change this section to evaluate different accesses //
            ////////////////////////////////////////////////////////
            // do your indexing here
            #define BLOCK_SIZE_X BLOCKDIMX
            #define BLOCK_SIZE_Y BLOCKDIMY
            #define xdim 32
            int i       = threadIdx.x;
            int j       = threadIdx.y;
            int idx     = blockIdx.x*BLOCK_SIZE_X + threadIdx.x;
            int idy     = blockIdx.y*BLOCK_SIZE_Y + threadIdx.y;

            int index1  = j*BLOCK_SIZE_Y+i;

            int i1      = (index1)%(BLOCK_SIZE_X+1);
            int j1      = (index1)/(BLOCK_SIZE_Y+1);

            int i2      = (BLOCK_SIZE_X*BLOCK_SIZE_Y+index1)%(BLOCK_SIZE_X+1);
            int j2      = (BLOCK_SIZE_X*BLOCK_SIZE_Y+index1)/(BLOCK_SIZE_Y+1);
            // calculate the accessed location and offset here
            // change the line "Ezx_h[(blockIdx.y*BLOCK_SIZE_Y+j1)*xdim+(blockIdx.x*BLOCK_SIZE_X+i1)];" to
            int addr = (blockIdx.y*BLOCK_SIZE_Y+j1)*xdim+(blockIdx.x*BLOCK_SIZE_X+i1);
            int size = sizeof(double);
            //////////////////////////
            // End of modifications //
            //////////////////////////

            printf("tid (%d,%d) from blockid (%d,%d) accessing to block %d\n",threadIdx.x,threadIdx.y,blockIdx.x,blockIdx.y,CACHE_BLOCK_ADDR(addr,size));
            // check whether it can be merged with existing requests or not
            short merged=0;
            for(III=0; III<unq_accesses; III++){
                if(CACHE_BLOCK_ADDR(addr,size)==CACHE_BLOCK_ADDR(unq_addr[III],size)){
                    merged=1;
                    break;
                }
            }
            if(!merged){
                // new cache block accessed over this coalescing width
                unq_addr[unq_accesses]=CACHE_BLOCK_ADDR(addr,size);
                unq_accesses++;
            }
        }
    }
    printf("%d threads make %d memory transactions\n",GRIDDIMX*GRIDDIMY*BLOCKDIMX*BLOCKDIMY, total_unq_accesses);
}

该代码将为网格的每个线程运行并计算合并请求的数量，内存访问合并的度量。

要使用分析器，请将代码的索引计算部分粘贴到指定区域，并将内存访问（数组）分解为“地址”和“大小”。我已经为您的代码完成了此操作，其中索引为：

int i       = threadIdx.x;
int j       = threadIdx.y;
int idx     = blockIdx.x*BLOCK_SIZE_X + threadIdx.x;
int idy     = blockIdx.y*BLOCK_SIZE_Y + threadIdx.y;

int index1  = j*BLOCK_SIZE_Y+i;

int i1      = (index1)%(BLOCK_SIZE_X+1);
int j1      = (index1)/(BLOCK_SIZE_Y+1);

int i2      = (BLOCK_SIZE_X*BLOCK_SIZE_Y+index1)%(BLOCK_SIZE_X+1);
int j2      = (BLOCK_SIZE_X*BLOCK_SIZE_Y+index1)/(BLOCK_SIZE_Y+1);

并且内存访问是：

Ezx_h_shared_ext[i1][j1]=Ezx_h[(blockIdx.y*BLOCK_SIZE_Y+j1)*xdim+(blockIdx.x*BLOCK_SIZE_X+i1)];

分析器报告 4096 个线程访问 4064 个缓存块。为您的实际网格和块大小运行代码并分析合并行为。

于 2012-12-07T23:08:06.893 回答
3

随着 GPU 的发展，获得合并访问的要求变得不那么严格。对于早期的 GPU 架构，您对合并访问的描述比较新的架构更准确。特别是，费米（计算能力 2.0）显着放宽了要求。在 Fermi 和更高版本上，连续访问内存位置并不重要。相反，重点已经转移到使用尽可能少的内存事务访问内存。在 Fermi 上，全局内存事务为 128 字节宽。因此，当一个 warp 中的 32 个线程遇到执行加载或存储的指令时，将安排 128 字节事务来为 warp 中的所有线程提供服务。然后，性能取决于需要多少事务。如果所有线程都访问与 128 字节对齐的 128 字节区域内的值，单笔交易是必要的。如果所有线程访问不同 128 字节区域中的值，则需要 32 个事务。这将是服务于扭曲中单个指令的请求的最坏情况。

您使用 CUDA 分析器之一来确定服务请求所需事务的平均值。该数字应尽可能接近 1。较高的数字意味着您应该查看是否有机会优化内核中的内存访问。

于 2012-12-07T22:49:09.463 回答
1

视觉分析器是检查工作的绝佳工具。在您拥有一段功能正确的代码后，然后从可视化分析器中运行它。例如，在 linux 上，假设您有一个 X 会话，只需从终端窗口运行 nvvp。然后，您将获得一个向导，该向导将提示您对应用程序进行概要分析以及任何命令行参数。

然后，分析器将对您的应用程序进行基本运行以收集统计信息。您还可以选择更高级的统计信息收集（需要额外运行），其中之一是内存利用率统计信息。它将以峰值百分比的形式报告内存利用率，并且还会针对它认为值得您注意的低利用率标记警告。

如果您的使用率高于 50%，那么您的应用程序可能正在按照您期望的方式运行。如果您的数字较低，您可能错过了一些合并细节。它将分别报告内存读取和内存写入的统计信息。要获得 100% 或接近它，您还需要确保来自 warp 的合并读取和写入在 128 字节边界上对齐。

在这些情况下的一个常见错误是使用基于 threadIdx.y 的变量作为变化最快的索引。在我看来，您并没有犯这个错误。例如，这是一个常见的错误，shared[threadIdx.x][threadIdx.y]因为这通常是我们在 C 中思考它的方式。但是线程首先在 x 轴上分组在一起，所以我们想要使用shared[threadIdx.y][threadIdx.x]或类似的东西。如果你确实犯了这个错误，你的代码在功能上仍然是正确的，但是你会在分析器中得到低百分比的利用率数字，比如大约 12% 甚至 3%。

如前所述，要达到 50% 以上并接近 100%，您需要确保不仅所有线程请求都是相邻的，而且它们在 128B 边界上对齐。由于 L1/L2 缓存，这些不是硬性规则，而是指南。缓存可以在一定程度上减轻一些错误。

于 2012-12-07T22:46:16.073 回答