Вопрос о деталях раздачи с блоков на СМ в CUDA - PullRequest
Новая
       68

Вопрос о деталях раздачи с блоков на СМ в CUDA

2 голосов
/ 23 августа 2011

Позвольте мне в качестве примера взять аппаратное обеспечение с вычислительной способностью 1.3.

30 SM.Тогда максимум 240 блоков могут быть запущены одновременно (с учетом ограничения регистров и общей памяти ограничение на количество блоков может быть намного ниже).Эти блоки за пределами 240 должны ждать доступных аппаратных ресурсов.

Мой вопрос заключается в том, когда эти блоки за пределами 240 будут назначены SM.Как только несколько блоков первых 240 будут завершены?Или когда все из первых 240 блоков завершены?

Я написал такой фрагмент кода. 

#include<stdio.h>
#include<string.h>
#include<cuda_runtime.h>
#include<cutil_inline.h>

const int BLOCKNUM = 1024;
const int N=240;
__global__ void kernel ( volatile int* mark ) {
    if ( blockIdx.x == 0 ) while ( mark[N] == 0 );
    if ( threadIdx.x == 0 ) mark[blockIdx.x] = 1;
}

int main() {
    int * mark;
    cudaMalloc ( ( void** ) &mark, sizeof ( int ) *BLOCKNUM );
    cudaMemset ( mark, 0, sizeof ( int ) *BLOCKNUM );
    kernel <<< BLOCKNUM, 1>>> ( mark );
    cudaFree ( mark );
    return 0;
}

Этот код вызывает тупик и не завершается.Но если я изменю N с 240 на 239, код сможет завершиться.Поэтому я хочу узнать некоторые подробности о планировании блоков.

Ответы [ 4 ]

2 голосов
/ 23 августа 2011

На GT200 с помощью микро-бенчмаркинга было продемонстрировано, что новые блоки планируются всякий раз, когда SM удаляет все текущие активные блоки, которые он выполнял.Таким образом, ответ - когда некоторые блоки завершены, и гранулярность планирования - уровень SM.Кажется, существует консенсус в отношении того, что графические процессоры Fermi имеют более тонкую детализацию планирования, чем аппаратные средства предыдущих поколений.

1 голос
/ 24 августа 2011

Я не могу найти ссылку на это для вычислительных возможностей <1.3.</p>

В архитектурах Fermi представлен новый диспетчер блоков, называемый механизмом GigaThread.
GigaThread обеспечивает немедленную замену блоков на SM по завершении выполнения, а также позволяет выполнять параллельное выполнение ядра.

0 голосов
/ 08 сентября 2011

В Fermi я уверен, что на SM запланирован блок, как только для него будет комната . Т.е. всякий раз, когда SM заканчивает выполнение одного блока, он выполняет другой блок, если остался какой-либо блок. (Однако фактический заказ не является детерминированным).

В старых версиях я не знаю. Но вы можете проверить это с помощью встроенной функции clock ().

Например, я использовал следующий код ядра OpenCL (вы можете легко преобразовать его в CUDA): 

   __kernel void test(uint* start, uint* end, float* buffer);
   {
       int id = get_global_id(0);
       start[id] = clock();
       __do_something_here;
       end[id] = clock();
   }

Затем выведите его в файл и постройте график. Вы увидите, насколько это наглядно.

0 голосов
/ 25 августа 2011

Хотя на этот вопрос нет официального ответа, вы можете измерить с помощью атомарных операций, когда ваши блоки начинают вашу работу и когда они заканчиваются.

Попробуйте поиграть со следующим кодом: 

#include <stdio.h>

const int maxBlocks=60; //Number of blocks of size 512 threads on current device required to achieve full occupancy

__global__ void emptyKernel() {}


__global__ void myKernel(int *control, int *output) {
        if (threadIdx.x==1) {
                //register that we enter
                int enter=atomicAdd(control,1);
                output[blockIdx.x]=enter;

                //some intensive and long task
                int &var=output[blockIdx.x+gridDim.x]; //var references global memory
                var=1;
                for (int i=0; i<12345678; ++i) {
                        var+=1+tanhf(var);
                }

                //register that we quit
                var=atomicAdd(control,1);
        }
}


int main() {

        int *gpuControl;
        cudaMalloc((void**)&gpuControl, sizeof(int));
        int cpuControl=0;
        cudaMemcpy(gpuControl,&cpuControl,sizeof(int),cudaMemcpyHostToDevice);


        int *gpuOutput;
        cudaMalloc((void**)&gpuOutput, sizeof(int)*maxBlocks*2);
        int cpuOutput[maxBlocks*2];

        for (int i=0; i<maxBlocks*2; ++i) //clear the host array just to be on the safe side
                cpuOutput[i]=-1;

        // play with these values
        const int thr=479;
        const int p=13;
        const int q=maxBlocks;

        //I found that this may actually affect the scheduler! Try with and without this call.
        emptyKernel<<<p,thr>>>();

        cudaEvent_t timerStart;
        cudaEvent_t timerStop;
        cudaEventCreate(&timerStart);
        cudaEventCreate(&timerStop);

        cudaThreadSynchronize();

        cudaEventRecord(timerStart,0);

        myKernel<<<q,512>>>(gpuControl, gpuOutput);

        cudaEventRecord(timerStop,0);
        cudaEventSynchronize(timerStop);

        cudaMemcpy(cpuOutput,gpuOutput,sizeof(int)*maxBlocks*2,cudaMemcpyDeviceToHost);

        cudaThreadSynchronize();
        float thisTime;
        cudaEventElapsedTime(&thisTime,timerStart,timerStop);

        cudaEventDestroy(timerStart);
        cudaEventDestroy(timerStop);
        printf("Elapsed time: %f\n",thisTime);

        for (int i=0; i<q; ++i)
                printf("%d: %d-%d\n",i,cpuOutput[i],cpuOutput[i+q]);
}

То, что вы получаете на выходе - это идентификатор блока, за которым следуют ввод "время" и выход "время". Таким образом, вы можете узнать, в каком порядке произошли эти события.

...