Результат эксперимента отличается от калькулятора занятости CUDA

Question

Я изучаю архитектуру CUDA.

Я сделал некоторый код параллельной обработки в среде, как показано ниже.

ГПУ: GTX580 (CC 2,0)

нитей на Блок: 16x16 = 256

Регистры на тему: 16

общей памяти для каждого блока: 48 байт

Я знаю, количество регистров и объем общей памяти по параметру компиляции: --ptxas-options = -v Кроме того, размер сетки 32x32 = 1024, и нет дополнительной общей памяти.

Итак, я попытался использовать CUDA_Occupancy_Calculator от NVIDIA. тогда, он сказал, отображается

3.) GPU Заполняемость данные здесь и в графах: Активные темы на многопроцессорных 1536 Активные Warps на многопроцессорных 48 Активные темы Блоки в многопроцессорных 6 размещения каждого многопроцессорных 100%

Итак, я запускаю приложение. Но результат показал, что размер блока составляет 8x8 быстрее, чем 16x16.

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

Answer 1

Поскольку вы скопировали your question from the Nvidia forum, я скопирую my answer, а также:

Для вашего ядра вашей ознакомительной пониженной производительности с более высокой заполняемости легко объясняется кэш перелива для повышения занятости.

Для локального массива temp[] при полном заполнении требуется 1536 × 15 × 4 = 92160 байт кеша, тогда как при 33% заполнении (при меньшем размере блока размером 8 × 8) требуется только 512 × 15 × 4 = 30720 байт за SM. С более высокой настройкой кэша/SMB 48 КБ последний может быть полностью кэширован, что позволяет полностью исключить внекорпусные обращения к памяти для temp[], но даже в настройках кеша/SM по умолчанию 16 КБ вероятность попадания в кеш существенно выше.

Поскольку массив temp[] в любом случае не нужен, самым быстрым вариантом (при любом заполнении) было бы полное его устранение. Возможно, компилятор смог достичь этого, если вы просто вставляете #pragma unroll перед циклом инициализации. В противном случае замените все виды использования temp[] на небольшую макрокоманду или встроенную функцию или просто замените результат на код (который в этом случае я бы даже нашел более удобочитаемым).