часы () в Opencl - PullRequest
Новая
       17

часы () в Opencl

4 голосов
/ 13 января 2012

Я знаю, что в CUDA есть функция clock (), где вы можете вставить код ядра и запросить время GPU.Но мне интересно, существует ли такая вещь в OpenCL?Есть ли способ запросить время GPU в OpenCL?(Я использую набор инструментов NVIDIA).

Ответы [ 5 ]

5 голосов
/ 13 января 2012

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

3 голосов
/ 22 октября 2012

Попробуйте это (конечно, работать только с NVidia OpenCL): 

uint clock_time()
{
    uint clock_time;
    asm("mov.u32 %0, %%clock;" : "=r"(clock_time));
    return clock_time;
}
3 голосов
/ 17 октября 2012

Просто для других, которые приходят к ней за помощью: краткое введение в профилирование времени выполнения ядра с OpenCL

Включить режим профилирования: 

cmdQueue = clCreateCommandQueue(context, *devices, CL_QUEUE_PROFILING_ENABLE, &err);

Профилирование ядра: 

cl_event prof_event; 
clEnqueueNDRangeKernel(cmdQueue, kernel, 1 , 0, globalWorkSize, NULL, 0, NULL, &prof_event);

Считать данные профилирования в: 

cl_ulong ev_start_time=(cl_ulong)0;     
cl_ulong ev_end_time=(cl_ulong)0;   

clFinish(cmdQueue);
err = clWaitForEvents(1, &prof_event);
err |= clGetEventProfilingInfo(prof_event, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &ev_start_time, NULL);
err |= clGetEventProfilingInfo(prof_event, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &ev_end_time, NULL);

Рассчитать время выполнения ядра: 

float run_time_gpu = (float)(ev_end_time - ev_start_time)/1000; // in usec

Профилирование отдельных рабочих элементов / рабочих групп пока невозможно. Вы можете установить globalWorkSize = localWorkSize для профилирования. Тогда у вас есть только одна рабочая группа.

Кстати: профилирование одного рабочего элемента (некоторых рабочих элементов) не очень полезно. Только с некоторыми рабочими элементами вы не сможете скрыть задержки памяти и накладные расходы, ведущие к бессмысленным измерениям.

2 голосов
/ 18 октября 2012

В NVIDIA OpenCL SDK есть пример Использование встроенного PTX с OpenCL . Регистр часов доступен через встроенный PTX как специальный регистр% clock. % clock описывается в PTX: руководство по параллельному выполнению потоков ISA . Вы должны быть в состоянии заменить %% laneid на %% clock.

Я никогда не проверял это с OpenCL, но использую его в CUDA.

Обратите внимание, что компилятор может переупорядочить или удалить чтение регистра.

1 голос
/ 13 декабря 2015

В NVIDIA вы можете использовать следующее: 

typedef unsigned long uint64_t; // if you haven't done so earlier
inline uint64_t n_nv_Clock()
{
    uint64_t n_clock;
    asm volatile("mov.u64 %0, %%clock64;" : "=l" (n_clock)); // make sure the compiler will not reorder this
    return n_clock;
}

Ключевое слово volatile сообщает оптимизатору, что вы действительно это имеете в виду, и не хотите, чтобы его убирали / оптимизировали.Это стандартный способ сделать это как в PTX , так и, например, в gcc .

. Обратите внимание, что возвращается тактовая частота , а не наносекунды.Вам необходимо запросить тактовую частоту устройства (используя clGetDeviceInfo(device, CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof(freq), &freq, 0))).Также обратите внимание, что на старых устройствах есть две частоты (или три, если вы посчитаете частоту памяти, которая не имеет значения в этом случае): часы устройства и часы шейдера.Вам нужны шейдерные часы.

С 64-битной версией регистра вам не нужно беспокоиться о переполнении, поскольку на это обычно уходят сотни лет.С другой стороны, 32-разрядная версия может переполняться довольно часто (вы все равно можете восстановить результат - если он не переполнится дважды).

...