与 CUDA 占用计算器不同的实验结果

我研究CUDA架构。

我在如下环境中编写了一些并行处理代码。

GPU：GTX580（CC为2.0）

每块线程：16x16 = 256

每线程寄存器：16

每块共享内存：48 字节

我通过编译选项知道寄存器的数量和共享内存的大小：--ptxas-options=-v 另外，网格大小为32x32 = 1024，并且没有额外的共享内存。

因此，我尝试使用 NVIDIA 的 CUDA_Occupancy_Calculator。 然后，它说，

3.) GPU 占用数据显示在此处和图表中： 每个多处理器的活动线程 1536 每个多处理器的活动扭曲 48 每个多处理器的活动线程块 6 每个多处理器占用率 100%

所以，我运行该应用程序。 但是，结果表明块大小比 16x16 快 8x8。

8x8表示块大小，网格大小为64x64。 16x16表示块大小，网格大小为32x32。 因此，线程总数是相同的。它没有改变。

我不知道为什么。请帮我。

以下代码是我的程序的一部分。

void LOAD_VERTEX(){
        MEM[0] = 60;    //y0 
        MEM[1] = 50;    //x0
        MEM[2] = 128;   //r0
        MEM[3] = 0;     //g0
        MEM[4] = 70;    //b0
        MEM[5] = 260;
        MEM[6] = 50;
        MEM[7] = 135;
        MEM[8] = 70;
        MEM[9] = 0;
        MEM[10] = 260;
        MEM[11] = 250;
        MEM[12] = 0;
        MEM[13] = 200;
        MEM[14] = 55;
        MEM[15] = 60;
        MEM[16] = 250;
        MEM[17] = 55;
        MEM[18] = 182;
        MEM[19] = 100;
        MEM[20] = 30;
        MEM[21] = 330;
        MEM[22] = 72;
        MEM[23] = 12;
        MEM[24] = 25;
        MEM[25] = 30;
        MEM[26] = 130;
        MEM[27] = 80;
        MEM[28] = 255;
        MEM[29] = 15;
        MEM[30] = 230; 
        MEM[31] = 330;
        MEM[32] = 56;   
        MEM[33] = 186;  
        MEM[34] = 201;
}

__global__ void PRINT_POLYGON( unsigned char *IMAGEin, int *MEMin, int dev_ID, int a, int b, int c)
{
        int i = blockIdx.x*TILE_WIDTH + threadIdx.x;
        int j = blockIdx.y*TILE_HEIGHT + threadIdx.y;

        float result_a, result_b;
        int temp[15];
        int k;

        for(k = 0; k < 5; k++){
                temp[k] = a*5+k;
                temp[k+5] = b*5+k;
                temp[k+10] = c*5+k;
        }

        int result_a_up = ((MEMin[temp[11]]-MEMin[temp[1]])*(i-MEMin[temp[0]]))-((MEMin[temp[10]]-MEMin[temp[0]])*(j-MEMin[temp[1]]));
        int result_a_down = ((MEMin[temp[11]]-MEMin[temp[1]])*(MEMin[temp[5]]-MEMin[temp[0]]))-((MEMin[temp[6]]-MEMin[temp[1]])*(MEMin[temp[10]]-MEMin[temp[0]]));

        int result_b_up = ((MEMin[temp[6]] -MEMin[temp[1]])*(MEMin[temp[0]]-i))-((MEMin[temp[5]] -MEMin[temp[0]])*(MEMin[temp[1]]-j));
        int result_b_down = ((MEMin[temp[11]]-MEMin[temp[1]])*(MEMin[temp[5]]-MEMin[temp[0]]))-((MEMin[temp[6]]-MEMin[temp[1]])*(MEMin[temp[10]]-MEMin[temp[0]]));

        result_a = float(result_a_up) / float(result_a_down);
        result_b = float(result_b_up) / float(result_b_down);

        int isIn = (0 <= result_a && result_a <=1) && ((0 <= result_b && result_b <= 1)) && ((0 <= (result_a+result_b) && (result_a+result_b) <= 1));

        IMAGEin[(i*HEIGHTs+j)*CHANNELS] += (int)(float(MEMin[temp[2]]) + (float(MEMin[temp[7]])-float(MEMin[temp[2]]))*result_a + (float(MEMin[temp[12]])-float(MEMin[temp[2]]))*result_b) * isIn;      //Red Channel
        IMAGEin[(i*HEIGHTs+j)*CHANNELS+1] += (int)(float(MEMin[temp[3]]) + (float(MEMin[temp[8]])-float(MEMin[temp[3]]))*result_a + (float(MEMin[temp[13]])-float(MEMin[temp[3]]))*result_b) * isIn;    //Green Channel
        IMAGEin[(i*HEIGHTs+j)*CHANNELS+2] += (int)(float(MEMin[temp[4]]) + (float(MEMin[temp[9]])-float(MEMin[temp[4]]))*result_a + (float(MEMin[temp[14]])-float(MEMin[temp[4]]))*result_b) * isIn;    //Blue Channel

}

//The information each device
struct DataStruct {
    int                 deviceID;
    unsigned char       IMAGE_SEG[WIDTH*HEIGHTs*CHANNELS];
};

void* routine( void *pvoidData ) {
        DataStruct  *data = (DataStruct*)pvoidData;
        unsigned char *dev_IMAGE;
        int *dev_MEM;
        unsigned char *IMAGE_SEG = data->IMAGE_SEG;

        HANDLE_ERROR(cudaSetDevice(data->deviceID));

        //initialize array
        memset(IMAGE_SEG, 0, WIDTH*HEIGHTs*CHANNELS);

        printf("Device %d Starting..\n", data->deviceID);

        //Evaluate Time
        cudaEvent_t start, stop;
        cudaEventCreate( &start );
        cudaEventCreate( &stop );

        HANDLE_ERROR( cudaMalloc( (void **)&dev_MEM, sizeof(int)*35) );                                //Creating int array each Block
        HANDLE_ERROR( cudaMalloc( (void **)&dev_IMAGE, sizeof(unsigned char)*WIDTH*HEIGHTs*CHANNELS) ); //output array

        cudaMemcpy(dev_MEM, MEM, sizeof(int)*256, cudaMemcpyHostToDevice);
        cudaMemset(dev_IMAGE, 0, sizeof(unsigned char)*WIDTH*HEIGHTs*CHANNELS);

        dim3    grid(WIDTH/TILE_WIDTH, HEIGHTs/TILE_HEIGHT);            //blocks in a grid
        dim3    block(TILE_WIDTH, TILE_HEIGHT);                         //threads in a block

        cudaEventRecord(start, 0);

        PRINT_POLYGON<<<grid,block>>>( dev_IMAGE, dev_MEM, data->deviceID, 0, 1, 2);                    //Start the Kernel
        PRINT_POLYGON<<<grid,block>>>( dev_IMAGE, dev_MEM, data->deviceID, 0, 2, 3);                    //Start the Kernel
        PRINT_POLYGON<<<grid,block>>>( dev_IMAGE, dev_MEM, data->deviceID, 0, 3, 4);                    //Start the Kernel
        PRINT_POLYGON<<<grid,block>>>( dev_IMAGE, dev_MEM, data->deviceID, 0, 4, 5);                    //Start the Kernel
        PRINT_POLYGON<<<grid,block>>>( dev_IMAGE, dev_MEM, data->deviceID, 3, 2, 4);                    //Start the Kernel
        PRINT_POLYGON<<<grid,block>>>( dev_IMAGE, dev_MEM, data->deviceID, 2, 6, 4);                    //Start the Kernel

        cudaEventRecord(stop, 0);
        cudaEventSynchronize(stop);

        HANDLE_ERROR( cudaMemcpy( IMAGE_SEG, dev_IMAGE, sizeof(unsigned char)*WIDTH*HEIGHTs*CHANNELS, cudaMemcpyDeviceToHost ) );
        HANDLE_ERROR( cudaFree( dev_MEM ) );
        HANDLE_ERROR( cudaFree( dev_IMAGE ) );

        cudaEventElapsedTime( &elapsed_time_ms[data->deviceID], start, stop );          //Calculate elapsed time
        cudaEventDestroy(start);
        cudaEventDestroy(stop);

        printf("Algorithm Elapsed Time : %f ms(Device %d)\n", elapsed_time_ms[data->deviceID], data->deviceID);
        printf("Device %d Complete!\n", data->deviceID);

        return 0;
}

int main( void )
{       
        int i;
        CUTThread thread[7];

        printf("Program Start.\n");                     
        LOAD_VERTEX();

        DataStruct data[DEVICENUM];                     //define device info

        for(i = 0; i < DEVICENUM; i++){
                data[i].deviceID = i;
                thread[i] = start_thread(routine, &(data[i]));
        }

        for(i = 0; i < DEVICENUM; i++){
                end_thread(thread[i]);
        }

        cudaFreeHost(MEM);

    return 0;
}

既然你复制过来了您来自 Nvidia 论坛的问题 https://devtalk.nvidia.com/default/topic/532759/i-39-ve-a-question-about-cuda-occuapncy-calculator-by-nvidia/，我复制一下我的答案 https://devtalk.nvidia.com/default/topic/532759/cuda-programming-and-performance/i-39-ve-a-question-about-cuda-occuapncy-calculator-by-nvidia/post/3750080/#3750080还有：

对于您的内核，您发现性能随占用率升高而降低，这很容易解释为缓存溢出导致更高的占用率。

本地数组temp[]完全占用时需要 1536×15×4=92160 字节的缓存，而 33% 占用时（对于较小的 8×8 块大小）每个 SM 仅需要 512×15×4=30720 字节。通过更大的 48kB 缓存/SM 设置，后者可以完全缓存，从而消除片外内存访问temp[]几乎完全如此，但即使在默认的 16kB 缓存/SM 设置中，缓存命中概率也会显着提高。

As the temp[]无论如何都不需要数组，最快的选择（无论占用哪种情况）都是完全消除它。如果您只插入一个，编译器可能已经能够实现这一点#pragma unroll在初始化循环之前。否则替换所有使用temp[]使用一些宏或内联函数，甚至只是将结果替换到代码中（在这种情况下，我什至会发现更具可读性）。