Синхронизация устройства Shacker DirectX 11?

Question

Фон: выполнить бенчмаркинг/сравнение по платформам GPGPU.

Проблема: Синхронизация устройства при отправке DirectX 11 Compute Shader.

Ищет эквивалент cudaDeviceSynchronize() из clFinish (...) сделать справедливый comparisson как мой алгоритм выполняет.

Функции CUDA и OpenCL более понятны в вопросах блокировки/неблокирования. DirectCompute, тем не менее, больше связан с графическим конвейером (которого я изучаю и очень незнакомым), и поэтому у меня возникли проблемы с обнаружением блокировки вызова Dispatch или если предыдущее выделение/перенос памяти завершено.

Код DX_1:

// Setup 
... 
for (...) { 
    startTimer(); 
    context->Dispatch(number_of_groups, 1, 1); 
    times[i] = stopTimer(); 
} 
// Release 
... 

Код DX_2:

for (...) { 
    // Setup 
    ... 
    startTimer(); 
    context->Dispatch(number_of_groups, 1, 1); 
    times[i] = stopTimer(); 
    // Release 
    ... 
} 

Результаты (среднее время от 2^2 до 2^11 элементов):

DX_1 DX_2 CUDA 
1.6 205.5 24.8 
1.8 133.4 24.8 
29.1 186.5 25.6 
18.6 175.0 25.6 
11.4 187.5 26.6 
85.2 127.7 26.3 
166.4 151.1 28.1 
98.2 149.5 35.2 
26.8 203.5 31.6 

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

Примечание 2: Это очень короткие последовательности (4 - 2048 элементов), интересные тесты выполняются с размером проблем до 2^26 элементов.

Answer 1

Если вас интересует, сколько времени занимает конкретный Draw или Dispatch на GPU, вы должны взглянуть на запросы Timestamp DirectX 11. Вы можете запросить частоту тактового генератора и текущее значение часов до и после работы некоторых графических процессоров и выяснить, сколько времени прошло время стены.

Это, вероятно, хорошая грунтовка/пример того, как это сделать:

https://mynameismjp.wordpress.com/2011/10/13/profiling-in-dx11-with-queries/

Answer 2

Мое новое решение, чтобы избежать синхронизации с устройством. Я посмотрел на некоторые методы retreiting timestamps, а результаты выглядят нормально, и я уверен, что сравнения достаточно справедливы. Я сравнивал время CUDA (Event Record и QPC), и разница небольшая, казалось бы, накладные расходы.

CUDA Event Host QPC 
4,6   30,0 
4,8   30,0 
5,0   31,0 
5,2   32,0 
5,6   34,0 
6,1   34,0 
6,9   31,0 
8,3   47,0 
9,2   34,0 
12,0  39,0 
16,7  46,0 
20,5  55,0 
32,1  69,0 
48,5  111,0 
86,0  134,0 
182,4  237,0 
419,0  473,0 

В случае мой вопрос приносит кому-то в надежде найти, как сделать GPGPU бенчмаркинг я оставлю код за демонстрируя свою текущую стратегию бенчмаркинга.

Примеры кода, CUDA

cudaEvent_t start, stop; 
cudaEventCreate(&start); 
cudaEventCreate(&stop); 
float milliseconds = 0; 
cudaEventRecord(start); 
... 
// Launch my algorithm 
... 
cudaEventRecord(stop); 
cudaEventSynchronize(stop); 
cudaEventElapsedTime(&milliseconds, start, stop);   

OpenCL

cl_event start_event, end_event; 
cl_ulong start = 0, end = 0; 
// Enqueue a dummy kernel for the start event. 
clEnqueueNDRangeKernel(..., &start_event); 
... 
// Launch my algorithm 
... 
// Enqueue a dummy kernel for the end event. 
clEnqueueNDRangeKernel(..., &end_event); 
clWaitForEvents(1, &end_event); 
clGetEventProfilingInfo(start_event, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &start, NULL); 
clGetEventProfilingInfo(end_event, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &end, NULL); 
timeInMS = (double)(end - start)*(double)(1e-06);  

DirectCompute

Здесь я следовал предложение от Адама Майлса и заглянул в этот источник.Будет выглядеть примерно так:

ID3D11Device*    device = nullptr; 
... 
// Setup 
... 
ID3D11QueryPtr disjoint_query; 
ID3D11QueryPtr q_start; 
ID3D11QueryPtr q_end; 
... 
if (disjoint_query == NULL) 
{ 
    D3D11_QUERY_DESC desc; 
    desc.Query = D3D11_QUERY_TIMESTAMP_DISJOINT; 
    desc.MiscFlags = 0; 
    device->CreateQuery(&desc, &disjoint_query); 
    desc.Query = D3D11_QUERY_TIMESTAMP; 
    device->CreateQuery(&desc, &q_start); 
    device->CreateQuery(&desc, &q_end); 
} 
context->Begin(disjoint_query); 
context->End(q_start); 
... 
// Launch my algorithm 
... 
context->End(q_end); 
context->End(disjoint_query); 
UINT64 start, end; 
D3D11_QUERY_DATA_TIMESTAMP_DISJOINT q_freq; 
while (S_OK != context->GetData(q_start, &start, sizeof(UINT64), 0)){}; 
while (S_OK != context->GetData(q_end, &end, sizeof(UINT64), 0)){}; 
while (S_OK != context->GetData(disjoint_query, &q_freq, sizeof(D3D11_QUERY_DATA_TIMESTAMP_DISJOINT), 0)){}; 
timeInMS = (((double)(end - start))/((double)q_freq.Frequency)) * 1000.0; 

C/C++/OpenMP

static LARGE_INTEGER StartingTime, EndingTime, ElapsedMicroseconds, Frequency; 

static void __inline startTimer() 
{ 
    QueryPerformanceFrequency(&Frequency); 
    QueryPerformanceCounter(&StartingTime); 
} 

static double __inline stopTimer() 
{ 
    QueryPerformanceCounter(&EndingTime); 
    ElapsedMicroseconds.QuadPart = EndingTime.QuadPart - StartingTime.QuadPart; 
    ElapsedMicroseconds.QuadPart *= 1000000; 
    ElapsedMicroseconds.QuadPart /= Frequency.QuadPart; 
    return (double)ElapsedMicroseconds.QuadPart; 
} 

Мои примеры кода взяты из контекста, и я пытался сделать некоторые очистке, но ошибки могут присутствовать.