DirectX 11 计算着色器设备同步？

Question

背景：对GPGPU平台进行基准测试/比较。

问题：在调度DirectX 11 Compute Shader时，设备同步的问题。

寻找与我的算法性能相当的cudaDeviceSynchronize()或clFinish(...)的等价物来进行公平比较。

CUDA和OpenCL函数更清楚地说明了阻塞/非阻塞问题。然而，DirectCompute更多地涉及图形管道（我正在学习并且非常不熟悉），因此我很难确定调度调用是否是阻塞的，以及之前的内存分配/传输是否已经完成。

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 

注意：这些时间是在连接屏幕的台式GPU上运行的，可能会出现一些不稳定的计时。时间不应包括主机到设备缓冲区的传输。

注意2：这些是非常短的序列（4-2048个元素），有趣的测试是针对多达2^26个元素的问题规模进行的。

- thorbear

Answer 1

我的新解决方案是避免与设备同步。我已经研究了一些检索时间戳的方法，结果看起来不错，我相当确定比较是足够公平的。我比较了我的CUDA时间（事件记录 vs. 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

这里我遵循了Adam Miles的建议，并研究了该源代码。看起来会是这样的：

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

我的代码示例是摘自上下文，我尝试进行了一些清理，但可能存在错误。

- thorbear

---

可能我会把它设为答案，但我会等待我的OpenGL解决方案和任何可能的建议。 - thorbear

Answer 2

如果您对GPU上特定绘制或调度所需的时间感兴趣，那么您应该查看DirectX 11的时间戳查询。您可以在执行某些GPU工作之前和之后查询GPU的时钟频率和当前时钟值，并计算出它所花费的墙上时间。

这可能是如何执行此操作的良好入门/示例：

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

- Adam Miles

---

谢谢！我已经编写了使用DirectX 11的时间戳查询功能的代码，而且它看起来很好用。我认识到我的主要问题将会是公平比较。到目前为止，我已经在主机上使用Windows QueryPerformanceCounter（QPC）API，但这包括很多开销，而在使用设备时间戳时不包括这些开销。 - thorbear

正如您在原始问题中所提到的，计时实际发出API调用本身所需的时间并不是您感兴趣的。除非您能确保在调用Dispatch时GPU处于空闲状态（可能），立即开始执行工作（很可能不是）并且工作在停止计时器之前已经完成（它不会），否则使用QPC将无法提供您想要的信息。 - Adam Miles