summaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
Diffstat
-rw-r--r--slang-gfx.h27
-rw-r--r--tools/gfx/command-writer.h13
-rw-r--r--tools/gfx/cpu/render-cpu.cpp44
-rw-r--r--tools/gfx/cuda/render-cuda.cpp83
-rw-r--r--tools/gfx/d3d11/render-d3d11.cpp107
-rw-r--r--tools/gfx/d3d12/render-d3d12.cpp144
-rw-r--r--tools/gfx/debug-layer.cpp37
-rw-r--r--tools/gfx/debug-layer.h17
-rw-r--r--tools/gfx/immediate-renderer-base.cpp26
-rw-r--r--tools/gfx/immediate-renderer-base.h8
-rw-r--r--tools/gfx/open-gl/render-gl.cpp12
-rw-r--r--tools/gfx/renderer-shared.cpp2
-rw-r--r--tools/gfx/renderer-shared.h1
-rw-r--r--tools/gfx/vulkan/render-vk.cpp106
-rw-r--r--tools/gfx/vulkan/vk-api.h8
15 files changed, 630 insertions, 5 deletions
diff --git a/slang-gfx.h b/slang-gfx.h
index 83ecd0b29..380167fc9 100644
--- a/slang-gfx.h
+++ b/slang-gfx.h
@@ -849,10 +849,31 @@ public:
0xdaab0b1a, 0xf45d, 0x4ae9, { 0xbf, 0x2c, 0xe0, 0xbb, 0x76, 0x7d, 0xfa, 0xd1 } \
}
+enum class QueryType
+{
+ Timestamp,
+};
+
+class IQueryPool : public ISlangUnknown
+{
+public:
+ struct Desc
+ {
+ QueryType type;
+ SlangInt count;
+ };
+public:
+ virtual SLANG_NO_THROW Result SLANG_MCALL getResult(SlangInt queryIndex, SlangInt count, uint64_t* data) = 0;
+};
+#define SLANG_UUID_IQueryPool \
+ { 0xc2cc3784, 0x12da, 0x480a, { 0xa8, 0x74, 0x8b, 0x31, 0x96, 0x1c, 0xa4, 0x36 } }
+
+
class ICommandEncoder : public ISlangUnknown
{
public:
virtual SLANG_NO_THROW void SLANG_MCALL endEncoding() = 0;
+ virtual SLANG_NO_THROW void SLANG_MCALL writeTimestamp(IQueryPool* queryPool, SlangInt queryIndex) = 0;
};
#define SLANG_UUID_ICommandEncoder \
{ \
@@ -1104,6 +1125,9 @@ struct DeviceInfo
/// The name of the graphics adapter.
const char* adapterName = nullptr;
+
+ /// The clock frequency used in timestamp queries.
+ uint64_t timestampFrequency = 0;
};
enum class DebugMessageType
@@ -1378,6 +1402,9 @@ public:
/// Get the type of this renderer
virtual SLANG_NO_THROW const DeviceInfo& SLANG_MCALL getDeviceInfo() const = 0;
+
+ virtual SLANG_NO_THROW Result SLANG_MCALL createQueryPool(
+ const IQueryPool::Desc& desc, IQueryPool** outPool) = 0;
};
#define SLANG_UUID_IDevice \
diff --git a/tools/gfx/command-writer.h b/tools/gfx/command-writer.h
index e93244ed1..adbb53d7a 100644
--- a/tools/gfx/command-writer.h
+++ b/tools/gfx/command-writer.h
@@ -23,7 +23,8 @@ enum class CommandName
SetStencilReference,
DispatchCompute,
UploadBufferData,
- CopyBuffer
+ CopyBuffer,
+ WriteTimestamp,
};
const uint8_t kMaxCommandOperands = 5;
@@ -82,6 +83,7 @@ public:
Slang::List<Command> m_commands;
Slang::List<Slang::ComPtr<ISlangUnknown>> m_objects;
Slang::List<uint8_t> m_data;
+ bool m_hasWriteTimestamps = false;
public:
void clear()
@@ -91,6 +93,7 @@ public:
obj = nullptr;
m_objects.clear();
m_data.clear();
+ m_hasWriteTimestamps = false;
}
// Copies user data into `m_data` buffer and returns the offset to retrieve the data.
@@ -248,5 +251,13 @@ public:
m_commands.add(
Command(CommandName::DispatchCompute, (uint32_t)x, (uint32_t)y, (uint32_t)z));
}
+
+ void writeTimestamp(IQueryPool* pool, SlangInt index)
+ {
+ auto poolOffset = encodeObject(pool);
+ m_commands.add(
+ Command(CommandName::WriteTimestamp, poolOffset, (uint32_t)index));
+ m_hasWriteTimestamps = true;
+ }
};
}
diff --git a/tools/gfx/cpu/render-cpu.cpp b/tools/gfx/cpu/render-cpu.cpp
index 03628c166..ac8b612fb 100644
--- a/tools/gfx/cpu/render-cpu.cpp
+++ b/tools/gfx/cpu/render-cpu.cpp
@@ -1,6 +1,8 @@
// render-cpu.cpp
#include "render-cpu.h"
+#include <chrono>
+
#include "slang.h"
#include "slang-com-ptr.h"
#include "slang-com-helper.h"
@@ -999,6 +1001,34 @@ public:
}
};
+class CPUQueryPool : public IQueryPool, public ComObject
+{
+public:
+ SLANG_COM_OBJECT_IUNKNOWN_ALL;
+ IQueryPool* getInterface(const Guid& guid)
+ {
+ if (guid == GfxGUID::IID_ISlangUnknown || guid == GfxGUID::IID_IQueryPool)
+ return static_cast<IQueryPool*>(this);
+ return nullptr;
+ }
+public:
+ List<uint64_t> m_queries;
+ Result init(const IQueryPool::Desc& desc)
+ {
+ m_queries.setCount(desc.count);
+ return SLANG_OK;
+ }
+ virtual SLANG_NO_THROW Result SLANG_MCALL getResult(
+ SlangInt queryIndex, SlangInt count, uint64_t* data) override
+ {
+ for (SlangInt i = 0; i < count; i++)
+ {
+ data[i] = m_queries[queryIndex + i];
+ }
+ return SLANG_OK;
+ }
+};
+
class CPUDevice : public ImmediateComputeDeviceBase
{
private:
@@ -1225,6 +1255,20 @@ public:
return Result();
}
+ virtual SLANG_NO_THROW Result SLANG_MCALL createQueryPool(
+ const IQueryPool::Desc& desc, IQueryPool** outPool) override
+ {
+ RefPtr<CPUQueryPool> pool = new CPUQueryPool();
+ returnComPtr(outPool, pool);
+ return SLANG_OK;
+ }
+
+ virtual void writeTimestamp(IQueryPool* pool, SlangInt index) override
+ {
+ static_cast<CPUQueryPool*>(pool)->m_queries[index] =
+ std::chrono::high_resolution_clock::now().time_since_epoch().count();
+ }
+
virtual SLANG_NO_THROW const DeviceInfo& SLANG_MCALL getDeviceInfo() const override
{
return m_info;
diff --git a/tools/gfx/cuda/render-cuda.cpp b/tools/gfx/cuda/render-cuda.cpp
index ed7f44ed2..3e93c090a 100644
--- a/tools/gfx/cuda/render-cuda.cpp
+++ b/tools/gfx/cuda/render-cuda.cpp
@@ -707,6 +707,58 @@ public:
}
};
+class CUDAQueryPool : public IQueryPool, public ComObject
+{
+public:
+ SLANG_COM_OBJECT_IUNKNOWN_ALL;
+ IQueryPool* getInterface(const Guid& guid)
+ {
+ if (guid == GfxGUID::IID_ISlangUnknown || guid == GfxGUID::IID_IQueryPool)
+ return static_cast<IQueryPool*>(this);
+ return nullptr;
+ }
+public:
+ // The event object for each query. Owned by the pool.
+ List<CUevent> m_events;
+
+ // The event that marks the starting point.
+ CUevent m_startEvent;
+
+ Result init(const IQueryPool::Desc& desc)
+ {
+ cuEventCreate(&m_startEvent, 0);
+ cuEventRecord(m_startEvent, 0);
+ m_events.setCount(desc.count);
+ for (SlangInt i = 0; i < m_events.getCount(); i++)
+ {
+ cuEventCreate(&m_events[i], 0);
+ }
+ return SLANG_OK;
+ }
+
+ ~CUDAQueryPool()
+ {
+ for (auto& e : m_events)
+ {
+ cuEventDestroy(e);
+ }
+ cuEventDestroy(m_startEvent);
+ }
+
+ virtual SLANG_NO_THROW Result SLANG_MCALL getResult(
+ SlangInt queryIndex, SlangInt count, uint64_t* data) override
+ {
+ for (SlangInt i = 0; i < count; i++)
+ {
+ float time = 0.0f;
+ cuEventSynchronize(m_events[i + queryIndex]);
+ cuEventElapsedTime(&time, m_startEvent, m_events[i + queryIndex]);
+ data[i] = (uint64_t)((double)time * 1000.0f);
+ }
+ return SLANG_OK;
+ }
+};
+
class CUDADevice : public RendererBase
{
private:
@@ -906,6 +958,11 @@ public:
m_writer->bindRootShaderObject(m_rootObject);
m_writer->dispatchCompute(x, y, z);
}
+
+ virtual SLANG_NO_THROW void SLANG_MCALL writeTimestamp(IQueryPool* pool, SlangInt index) override
+ {
+ m_writer->writeTimestamp(pool, index);
+ }
};
ComputeCommandEncoderImpl m_computeCommandEncoder;
@@ -959,6 +1016,11 @@ public:
{
m_writer->uploadBufferData(dst, offset, size, data);
}
+
+ virtual SLANG_NO_THROW void SLANG_MCALL writeTimestamp(IQueryPool* pool, SlangInt index) override
+ {
+ m_writer->writeTimestamp(pool, index);
+ }
};
ResourceCommandEncoderImpl m_resourceCommandEncoder;
@@ -1139,6 +1201,12 @@ public:
cudaMemcpy((uint8_t*)dstImpl->m_cudaMemory + offset, data, size, cudaMemcpyDefault);
}
+ void writeTimestamp(IQueryPool* pool, SlangInt index)
+ {
+ auto poolImpl = static_cast<CUDAQueryPool*>(pool);
+ cuEventRecord(poolImpl->m_events[index], stream);
+ }
+
void execute(CommandBufferImpl* commandBuffer)
{
for (auto& cmd : commandBuffer->m_commands)
@@ -1171,6 +1239,10 @@ public:
cmd.operands[2],
commandBuffer->getData<uint8_t>(cmd.operands[3]));
break;
+ case CommandName::WriteTimestamp:
+ writeTimestamp(
+ commandBuffer->getObject<IQueryPool>(cmd.operands[0]),
+ (SlangInt)cmd.operands[1]);
}
}
}
@@ -1218,6 +1290,7 @@ public:
cudaGetDeviceProperties(&deviceProperties, m_deviceIndex);
m_adapterName = deviceProperties.name;
m_info.adapterName = m_adapterName.begin();
+ m_info.timestampFrequency = 1000000;
}
return SLANG_OK;
@@ -1691,6 +1764,16 @@ public:
return SLANG_OK;
}
+ virtual SLANG_NO_THROW Result SLANG_MCALL createQueryPool(
+ const IQueryPool::Desc& desc,
+ IQueryPool** outPool) override
+ {
+ RefPtr<CUDAQueryPool> pool = new CUDAQueryPool();
+ SLANG_RETURN_ON_FAIL(pool->init(desc));
+ returnComPtr(outPool, pool);
+ return SLANG_OK;
+ }
+
virtual Result createShaderObjectLayout(
slang::TypeLayoutReflection* typeLayout,
ShaderObjectLayoutBase** outLayout) override
diff --git a/tools/gfx/d3d11/render-d3d11.cpp b/tools/gfx/d3d11/render-d3d11.cpp
index 3e71ff3d4..c4cd3f655 100644
--- a/tools/gfx/d3d11/render-d3d11.cpp
+++ b/tools/gfx/d3d11/render-d3d11.cpp
@@ -100,6 +100,9 @@ public:
UInt inputElementCount,
IInputLayout** outLayout) override;
+ virtual SLANG_NO_THROW Result SLANG_MCALL createQueryPool(
+ const IQueryPool::Desc& desc, IQueryPool** outPool) override;
+
virtual Result createShaderObjectLayout(
slang::TypeLayoutReflection* typeLayout,
ShaderObjectLayoutBase** outLayout) override;
@@ -143,11 +146,33 @@ public:
virtual void drawIndexed(UInt indexCount, UInt startIndex, UInt baseVertex) override;
virtual void dispatchCompute(int x, int y, int z) override;
virtual void submitGpuWork() override {}
- virtual void waitForGpu() override {}
+ virtual void waitForGpu() override
+ {
+
+ }
virtual SLANG_NO_THROW const DeviceInfo& SLANG_MCALL getDeviceInfo() const override
{
return m_info;
}
+ virtual void beginCommandBuffer(const CommandBufferInfo& info) override
+ {
+ if (info.hasWriteTimestamps)
+ {
+ m_immediateContext->Begin(m_disjointQuery);
+ }
+ }
+ virtual void endCommandBuffer(const CommandBufferInfo& info) override
+ {
+ if (info.hasWriteTimestamps)
+ {
+ m_immediateContext->End(m_disjointQuery);
+ }
+ }
+ virtual void writeTimestamp(IQueryPool* pool, SlangInt index) override
+ {
+ auto poolImpl = static_cast<QueryPoolImpl*>(pool);
+ m_immediateContext->End(poolImpl->getQuery(index));
+ }
protected:
@@ -337,6 +362,62 @@ protected:
ComPtr<ID3D11InputLayout> m_layout;
};
+ class QueryPoolImpl : public IQueryPool, public ComObject
+ {
+ public:
+ SLANG_COM_OBJECT_IUNKNOWN_ALL;
+ IQueryPool* getInterface(const Guid& guid)
+ {
+ if (guid == GfxGUID::IID_ISlangUnknown || guid == GfxGUID::IID_IQueryPool)
+ return static_cast<IQueryPool*>(this);
+ return nullptr;
+ }
+ public:
+ List<ComPtr<ID3D11Query>> m_queries;
+ RefPtr<D3D11Device> m_device;
+ D3D11_QUERY_DESC m_queryDesc;
+ Result init(const IQueryPool::Desc& desc, D3D11Device* device)
+ {
+ m_device = device;
+ m_queryDesc.MiscFlags = 0;
+ switch (desc.type)
+ {
+ case QueryType::Timestamp:
+ m_queryDesc.Query = D3D11_QUERY_TIMESTAMP;
+ break;
+ default:
+ return SLANG_E_INVALID_ARG;
+ }
+ m_queries.setCount(desc.count);
+ return SLANG_OK;
+ }
+ ID3D11Query* getQuery(SlangInt index)
+ {
+ if (!m_queries[index])
+ m_device->m_device->CreateQuery(&m_queryDesc, m_queries[index].writeRef());
+ return m_queries[index].get();
+ }
+
+ virtual SLANG_NO_THROW Result SLANG_MCALL getResult(
+ SlangInt queryIndex, SlangInt count, uint64_t* data) override
+ {
+ D3D11_QUERY_DATA_TIMESTAMP_DISJOINT disjointData;
+ while (S_OK != m_device->m_immediateContext->GetData(
+ m_device->m_disjointQuery, &disjointData, sizeof(D3D11_QUERY_DATA_TIMESTAMP_DISJOINT), 0))
+ {
+ Sleep(1);
+ }
+ m_device->m_info.timestampFrequency = disjointData.Frequency;
+
+ for (SlangInt i = 0; i < count; i++)
+ {
+ SLANG_RETURN_ON_FAIL(m_device->m_immediateContext->GetData(
+ m_queries[queryIndex + i], data + i, sizeof(uint64_t), 0));
+ }
+ return SLANG_OK;
+ }
+ };
+
class PipelineStateImpl : public PipelineStateBase
{
public:
@@ -1937,11 +2018,12 @@ protected:
ComPtr<ID3D11DeviceContext> m_immediateContext;
ComPtr<ID3D11Texture2D> m_backBufferTexture;
ComPtr<IDXGIFactory> m_dxgiFactory;
-
RefPtr<FramebufferImpl> m_currentFramebuffer;
RefPtr<PipelineStateImpl> m_currentPipelineState;
+ ComPtr<ID3D11Query> m_disjointQuery;
+
uint32_t m_stencilRef = 0;
bool m_depthStencilStateDirty = true;
@@ -2189,6 +2271,19 @@ SlangResult D3D11Device::initialize(const Desc& desc)
m_nvapi = true;
#endif
}
+
+ {
+ // Create a TIMESTAMP_DISJOINT query object to query/update frequency info.
+ D3D11_QUERY_DESC disjointQueryDesc = {};
+ disjointQueryDesc.Query = D3D11_QUERY_TIMESTAMP_DISJOINT;
+ SLANG_RETURN_ON_FAIL(m_device->CreateQuery(
+ &disjointQueryDesc, m_disjointQuery.writeRef()));
+ m_immediateContext->Begin(m_disjointQuery);
+ m_immediateContext->End(m_disjointQuery);
+ D3D11_QUERY_DATA_TIMESTAMP_DISJOINT disjointData = {};
+ m_immediateContext->GetData(m_disjointQuery, &disjointData, sizeof(disjointData), 0);
+ m_info.timestampFrequency = disjointData.Frequency;
+ }
return SLANG_OK;
}
@@ -2971,6 +3066,14 @@ Result D3D11Device::createInputLayout(const InputElementDesc* inputElementsIn, U
return SLANG_OK;
}
+Result D3D11Device::createQueryPool(const IQueryPool::Desc& desc, IQueryPool** outPool)
+{
+ RefPtr<QueryPoolImpl> result = new QueryPoolImpl();
+ SLANG_RETURN_ON_FAIL(result->init(desc, this));
+ returnComPtr(outPool, result);
+ return SLANG_OK;
+}
+
void* D3D11Device::map(IBufferResource* bufferIn, MapFlavor flavor)
{
BufferResourceImpl* bufferResource = static_cast<BufferResourceImpl*>(bufferIn);
diff --git a/tools/gfx/d3d12/render-d3d12.cpp b/tools/gfx/d3d12/render-d3d12.cpp
index 427e98125..e05d68661 100644
--- a/tools/gfx/d3d12/render-d3d12.cpp
+++ b/tools/gfx/d3d12/render-d3d12.cpp
@@ -124,6 +124,9 @@ public:
virtual SLANG_NO_THROW Result SLANG_MCALL createComputePipelineState(
const ComputePipelineStateDesc& desc, IPipelineState** outState) override;
+ virtual SLANG_NO_THROW Result SLANG_MCALL createQueryPool(
+ const IQueryPool::Desc& desc, IQueryPool** outState) override;
+
virtual SLANG_NO_THROW SlangResult SLANG_MCALL readTextureResource(
ITextureResource* resource,
ResourceState state,
@@ -316,6 +319,56 @@ public:
}
};
+ class QueryPoolImpl : public IQueryPool, public ComObject
+ {
+ public:
+ SLANG_COM_OBJECT_IUNKNOWN_ALL
+ IQueryPool* getInterface(const Guid& guid)
+ {
+ if (guid == GfxGUID::IID_ISlangUnknown || guid == GfxGUID::IID_IQueryPool)
+ return static_cast<IQueryPool*>(this);
+ return nullptr;
+ }
+ public:
+ Result init(const IQueryPool::Desc& desc, D3D12Device* device);
+
+ virtual SLANG_NO_THROW Result SLANG_MCALL getResult(SlangInt queryIndex, SlangInt count, uint64_t* data) override
+ {
+ m_commandList->Reset(m_commandAllocator, nullptr);
+ m_commandList->ResolveQueryData(m_queryHeap, m_queryType, (UINT)queryIndex, (UINT)count, m_readBackBuffer, 0);
+ m_commandList->Close();
+ ID3D12CommandList* cmdList = m_commandList;
+ m_commandQueue->ExecuteCommandLists(1, &cmdList);
+ m_eventValue++;
+ m_fence->SetEventOnCompletion(m_eventValue, m_waitEvent);
+ m_commandQueue->Signal(m_fence, m_eventValue);
+ WaitForSingleObject(m_waitEvent, INFINITE);
+
+ int8_t* mappedData = nullptr;
+ D3D12_RANGE readRange = { sizeof(uint64_t) * queryIndex,sizeof(uint64_t) * (queryIndex + count) };
+ m_readBackBuffer.getResource()->Map(0, &readRange, (void**)&mappedData);
+ memcpy(data, mappedData + sizeof(uint64_t) * queryIndex, sizeof(uint64_t) * count);
+ m_readBackBuffer.getResource()->Unmap(0, nullptr);
+ return SLANG_OK;
+ }
+
+ void writeTimestamp(ID3D12GraphicsCommandList* cmdList, SlangInt index)
+ {
+ cmdList->EndQuery(m_queryHeap, D3D12_QUERY_TYPE_TIMESTAMP, (UINT)index);
+ }
+
+ public:
+ D3D12_QUERY_TYPE m_queryType;
+ ComPtr<ID3D12QueryHeap> m_queryHeap;
+ D3D12Resource m_readBackBuffer;
+ ComPtr<ID3D12CommandAllocator> m_commandAllocator;
+ ComPtr<ID3D12GraphicsCommandList> m_commandList;
+ ComPtr<ID3D12Fence> m_fence;
+ ComPtr<ID3D12CommandQueue> m_commandQueue;
+ HANDLE m_waitEvent;
+ UINT64 m_eventValue = 0;
+ };
+
struct BoundVertexBuffer
{
RefPtr<BufferResourceImpl> m_buffer;
@@ -3391,6 +3444,11 @@ public:
m_framebuffer = nullptr;
}
+ virtual SLANG_NO_THROW void SLANG_MCALL writeTimestamp(IQueryPool* pool, SlangInt index) override
+ {
+ static_cast<QueryPoolImpl*>(pool)->writeTimestamp(m_d3dCmdList, index);
+ }
+
virtual SLANG_NO_THROW void SLANG_MCALL
setStencilReference(uint32_t referenceValue) override
{
@@ -3438,6 +3496,10 @@ public:
{
PipelineCommandEncoder::endEncodingImpl();
}
+ virtual SLANG_NO_THROW void SLANG_MCALL writeTimestamp(IQueryPool* pool, SlangInt index) override
+ {
+ static_cast<QueryPoolImpl*>(pool)->writeTimestamp(m_d3dCmdList, index);
+ }
void init(
D3D12Device* renderer,
TransientResourceHeapImpl* transientHeap,
@@ -3533,6 +3595,10 @@ public:
data);
}
virtual SLANG_NO_THROW void SLANG_MCALL endEncoding() {}
+ virtual SLANG_NO_THROW void SLANG_MCALL writeTimestamp(IQueryPool* pool, SlangInt index) override
+ {
+ static_cast<QueryPoolImpl*>(pool)->writeTimestamp(m_commandBuffer->m_cmdList, index);
+ }
};
ResourceCommandEncoderImpl m_resourceCommandEncoder;
@@ -4422,6 +4488,8 @@ Result D3D12Device::initialize(const Desc& desc)
// `CommandQueueImpl` holds a back reference to `D3D12Device`, make it a weak reference here
// since this object is already owned by `D3D12Device`.
m_resourceCommandQueue->breakStrongReferenceToDevice();
+ // Retrieve timestamp frequency.
+ m_resourceCommandQueue->m_d3dQueue->GetTimestampFrequency(&m_info.timestampFrequency);
SLANG_RETURN_ON_FAIL(createTransientResourceHeapImpl(0, 8, 4, m_resourceCommandTransientHeap.writeRef()));
// `TransientResourceHeap` holds a back reference to `D3D12Device`, make it a weak reference here
@@ -5520,4 +5588,80 @@ Result D3D12Device::createComputePipelineState(const ComputePipelineStateDesc& i
return SLANG_OK;
}
+Result D3D12Device::QueryPoolImpl::init(const IQueryPool::Desc& desc, D3D12Device* device)
+{
+ // Translate query type.
+ D3D12_QUERY_HEAP_DESC heapDesc = {};
+ heapDesc.Count = (UINT)desc.count;
+ heapDesc.NodeMask = 1;
+ switch (desc.type)
+ {
+ case QueryType::Timestamp:
+ heapDesc.Type = D3D12_QUERY_HEAP_TYPE_TIMESTAMP;
+ m_queryType = D3D12_QUERY_TYPE_TIMESTAMP;
+ break;
+ default:
+ return SLANG_E_INVALID_ARG;
+ }
+
+ // Create query heap.
+ auto d3dDevice = device->m_device;
+ SLANG_RETURN_ON_FAIL(d3dDevice->CreateQueryHeap(
+ &heapDesc, IID_PPV_ARGS(m_queryHeap.writeRef())));
+
+ // Create readback buffer.
+ D3D12_HEAP_PROPERTIES heapProps;
+ heapProps.Type = D3D12_HEAP_TYPE_READBACK;
+ heapProps.CPUPageProperty = D3D12_CPU_PAGE_PROPERTY_UNKNOWN;
+ heapProps.MemoryPoolPreference = D3D12_MEMORY_POOL_UNKNOWN;
+ heapProps.CreationNodeMask = 1;
+ heapProps.VisibleNodeMask = 1;
+ D3D12_RESOURCE_DESC resourceDesc = {};
+ _initBufferResourceDesc(sizeof(uint64_t) * desc.count, resourceDesc);
+ SLANG_RETURN_ON_FAIL(m_readBackBuffer.initCommitted(
+ d3dDevice,
+ heapProps,
+ D3D12_HEAP_FLAG_NONE,
+ resourceDesc,
+ D3D12_RESOURCE_STATE_COPY_DEST,
+ nullptr));
+
+ // Create command allocator.
+ SLANG_RETURN_ON_FAIL(d3dDevice->CreateCommandAllocator(
+ D3D12_COMMAND_LIST_TYPE_DIRECT, IID_PPV_ARGS(m_commandAllocator.writeRef())));
+
+ // Create command list.
+ SLANG_RETURN_ON_FAIL(d3dDevice->CreateCommandList(
+ 0,
+ D3D12_COMMAND_LIST_TYPE_DIRECT,
+ m_commandAllocator,
+ nullptr,
+ IID_PPV_ARGS(m_commandList.writeRef())));
+ m_commandList->Close();
+
+ // Create fence.
+ SLANG_RETURN_ON_FAIL(d3dDevice->CreateFence(
+ 0, D3D12_FENCE_FLAG_NONE, IID_PPV_ARGS(m_fence.writeRef())));
+
+ // Get command queue from device.
+ m_commandQueue = device->m_resourceCommandQueue->m_d3dQueue;
+
+ // Create wait event.
+ m_waitEvent = CreateEventEx(
+ nullptr,
+ false,
+ 0,
+ EVENT_ALL_ACCESS);
+
+ return SLANG_OK;
+}
+
+Result D3D12Device::createQueryPool(const IQueryPool::Desc& desc, IQueryPool** outState)
+{
+ RefPtr<QueryPoolImpl> queryPoolImpl = new QueryPoolImpl();
+ SLANG_RETURN_ON_FAIL(queryPoolImpl->init(desc, this));
+ returnComPtr(outState, queryPoolImpl);
+ return SLANG_OK;
+}
+
} // renderer_test
diff --git a/tools/gfx/debug-layer.cpp b/tools/gfx/debug-layer.cpp
index 26e55ca7e..13875d6ec 100644
--- a/tools/gfx/debug-layer.cpp
+++ b/tools/gfx/debug-layer.cpp
@@ -140,6 +140,8 @@ SLANG_GFX_DEBUG_GET_INTERFACE_IMPL(ShaderObject)
SLANG_GFX_DEBUG_GET_INTERFACE_IMPL(ShaderProgram)
SLANG_GFX_DEBUG_GET_INTERFACE_IMPL(Swapchain)
SLANG_GFX_DEBUG_GET_INTERFACE_IMPL(TransientResourceHeap)
+SLANG_GFX_DEBUG_GET_INTERFACE_IMPL(QueryPool)
+
#undef SLANG_GFX_DEBUG_GET_INTERFACE_IMPL
#undef SLANG_GFX_DEBUG_GET_INTERFACE_IMPL_PARENT
@@ -477,6 +479,16 @@ const DeviceInfo& DebugDevice::getDeviceInfo() const
return baseObject->getDeviceInfo();
}
+Result DebugDevice::createQueryPool(const IQueryPool::Desc& desc, IQueryPool** outPool)
+{
+ SLANG_GFX_API_FUNC;
+ RefPtr<DebugQueryPool> result = new DebugQueryPool();
+ result->desc = desc;
+ SLANG_RETURN_ON_FAIL(baseObject->createQueryPool(desc, result->baseObject.writeRef()));
+ returnComPtr(outPool, result);
+ return SLANG_OK;
+}
+
IResource::Type DebugBufferResource::getType()
{
SLANG_GFX_API_FUNC;
@@ -624,6 +636,12 @@ void DebugComputeCommandEncoder::dispatchCompute(int x, int y, int z)
baseObject->dispatchCompute(x, y, z);
}
+void DebugComputeCommandEncoder::writeTimestamp(IQueryPool* pool, SlangInt index)
+{
+ SLANG_GFX_API_FUNC;
+ baseObject->writeTimestamp(static_cast<DebugQueryPool*>(pool)->baseObject, index);
+}
+
void DebugRenderCommandEncoder::endEncoding()
{
SLANG_GFX_API_FUNC;
@@ -706,12 +724,24 @@ void DebugRenderCommandEncoder::setStencilReference(uint32_t referenceValue)
return baseObject->setStencilReference(referenceValue);
}
+void DebugRenderCommandEncoder::writeTimestamp(IQueryPool* pool, SlangInt index)
+{
+ SLANG_GFX_API_FUNC;
+ baseObject->writeTimestamp(static_cast<DebugQueryPool*>(pool)->baseObject, index);
+}
+
void DebugResourceCommandEncoder::endEncoding()
{
SLANG_GFX_API_FUNC;
baseObject->endEncoding();
}
+void DebugResourceCommandEncoder::writeTimestamp(IQueryPool* pool, SlangInt index)
+{
+ SLANG_GFX_API_FUNC;
+ baseObject->writeTimestamp(static_cast<DebugQueryPool*>(pool)->baseObject, index);
+}
+
void DebugResourceCommandEncoder::copyBuffer(
IBufferResource* dst,
size_t dstOffset,
@@ -977,4 +1007,11 @@ Result DebugRootShaderObject::setSpecializationArgs(
return baseObject->setSpecializationArgs(offset, args, count);
}
+Result DebugQueryPool::getResult(SlangInt index, SlangInt count, uint64_t* data)
+{
+ if (index < 0 || index + count >= desc.count)
+ GFX_DIAGNOSE_ERROR("index is out of bounds.");
+ return baseObject->getResult(index, count, data);
+}
+
} // namespace gfx
diff --git a/tools/gfx/debug-layer.h b/tools/gfx/debug-layer.h
index 89ee9d837..0fb8e681b 100644
--- a/tools/gfx/debug-layer.h
+++ b/tools/gfx/debug-layer.h
@@ -99,6 +99,20 @@ public:
size_t size,
ISlangBlob** outBlob) override;
virtual SLANG_NO_THROW const DeviceInfo& SLANG_MCALL getDeviceInfo() const override;
+ virtual SLANG_NO_THROW Result SLANG_MCALL createQueryPool(
+ const IQueryPool::Desc& desc,
+ IQueryPool** outPool) override;
+};
+
+class DebugQueryPool : public DebugObject<IQueryPool>
+{
+public:
+ SLANG_COM_OBJECT_IUNKNOWN_ALL;
+
+ IQueryPool::Desc desc;
+public:
+ IQueryPool* getInterface(const Slang::Guid& guid);
+ virtual SLANG_NO_THROW Result SLANG_MCALL getResult(SlangInt index, SlangInt count, uint64_t* data) override;
};
class DebugBufferResource : public DebugObject<IBufferResource>
@@ -227,6 +241,7 @@ public:
virtual SLANG_NO_THROW Result SLANG_MCALL
bindPipeline(IPipelineState* state, IShaderObject** outRootShaderObject) override;
virtual SLANG_NO_THROW void SLANG_MCALL dispatchCompute(int x, int y, int z) override;
+ virtual SLANG_NO_THROW void SLANG_MCALL writeTimestamp(IQueryPool* pool, SlangInt index) override;
public:
DebugCommandBuffer* commandBuffer;
@@ -264,6 +279,7 @@ public:
virtual SLANG_NO_THROW void SLANG_MCALL
drawIndexed(UInt indexCount, UInt startIndex = 0, UInt baseVertex = 0) override;
virtual SLANG_NO_THROW void SLANG_MCALL setStencilReference(uint32_t referenceValue) override;
+ virtual SLANG_NO_THROW void SLANG_MCALL writeTimestamp(IQueryPool* pool, SlangInt index) override;
public:
DebugCommandBuffer* commandBuffer;
@@ -289,6 +305,7 @@ public:
size_t size) override;
virtual SLANG_NO_THROW void SLANG_MCALL
uploadBufferData(IBufferResource* dst, size_t offset, size_t size, void* data) override;
+ virtual SLANG_NO_THROW void SLANG_MCALL writeTimestamp(IQueryPool* pool, SlangInt index) override;
public:
DebugCommandBuffer* commandBuffer;
diff --git a/tools/gfx/immediate-renderer-base.cpp b/tools/gfx/immediate-renderer-base.cpp
index 19a16eac1..8fffbfdfa 100644
--- a/tools/gfx/immediate-renderer-base.cpp
+++ b/tools/gfx/immediate-renderer-base.cpp
@@ -32,6 +32,7 @@ public:
public:
CommandWriter m_writer;
+ bool m_hasWriteTimestamps = false;
RefPtr<ImmediateRendererBase> m_renderer;
RefPtr<ShaderObjectBase> m_rootShaderObject;
@@ -163,6 +164,11 @@ public:
{
m_writer->setStencilReference(referenceValue);
}
+
+ virtual SLANG_NO_THROW void SLANG_MCALL writeTimestamp(IQueryPool* pool, SlangInt index) override
+ {
+ m_writer->writeTimestamp(pool, index);
+ }
};
RenderCommandEncoderImpl m_renderCommandEncoder;
@@ -227,6 +233,11 @@ public:
m_writer->bindRootShaderObject(m_commandBuffer->m_rootShaderObject);
m_writer->dispatchCompute(x, y, z);
}
+
+ virtual SLANG_NO_THROW void SLANG_MCALL writeTimestamp(IQueryPool* pool, SlangInt index) override
+ {
+ m_writer->writeTimestamp(pool, index);
+ }
};
ComputeCommandEncoderImpl m_computeCommandEncoder;
@@ -280,6 +291,11 @@ public:
{
m_writer->uploadBufferData(dst, offset, size, data);
}
+
+ virtual SLANG_NO_THROW void SLANG_MCALL writeTimestamp(IQueryPool* pool, SlangInt index) override
+ {
+ m_writer->writeTimestamp(pool, index);
+ }
};
ResourceCommandEncoderImpl m_resourceCommandEncoder;
@@ -375,6 +391,9 @@ public:
cmd.operands[3],
cmd.operands[4]);
break;
+ case CommandName::WriteTimestamp:
+ m_renderer->writeTimestamp(m_writer.getObject<IQueryPool>(cmd.operands[0]), (SlangInt)cmd.operands[1]);
+ break;
default:
assert(!"unknown command");
break;
@@ -411,10 +430,17 @@ public:
virtual SLANG_NO_THROW void SLANG_MCALL
executeCommandBuffers(uint32_t count, ICommandBuffer* const* commandBuffers) override
{
+ CommandBufferInfo info = {};
+ for (uint32_t i = 0; i < count; i++)
+ {
+ info.hasWriteTimestamps |= static_cast<CommandBufferImpl*>(commandBuffers[i])->m_writer.m_hasWriteTimestamps;
+ }
+ static_cast<ImmediateRendererBase*>(m_renderer.get())->beginCommandBuffer(info);
for (uint32_t i = 0; i < count; i++)
{
static_cast<CommandBufferImpl*>(commandBuffers[i])->execute();
}
+ static_cast<ImmediateRendererBase*>(m_renderer.get())->endCommandBuffer(info);
}
virtual SLANG_NO_THROW void SLANG_MCALL wait() override { getRenderer()->waitForGpu(); }
diff --git a/tools/gfx/immediate-renderer-base.h b/tools/gfx/immediate-renderer-base.h
index 3acceed15..b1e867edc 100644
--- a/tools/gfx/immediate-renderer-base.h
+++ b/tools/gfx/immediate-renderer-base.h
@@ -40,6 +40,11 @@ public:
void establishStrongReferenceToDevice() { m_renderer.establishStrongReference(); }
};
+struct CommandBufferInfo
+{
+ bool hasWriteTimestamps;
+};
+
class ImmediateRendererBase : public RendererBase
{
public:
@@ -75,6 +80,9 @@ public:
virtual void waitForGpu() = 0;
virtual void* map(IBufferResource* buffer, MapFlavor flavor) = 0;
virtual void unmap(IBufferResource* buffer, size_t offsetWritten, size_t sizeWritten) = 0;
+ virtual void writeTimestamp(IQueryPool* pool, SlangInt index) = 0;
+ virtual void beginCommandBuffer(const CommandBufferInfo&) {}
+ virtual void endCommandBuffer(const CommandBufferInfo&) {}
public:
Slang::RefPtr<ImmediateCommandQueueBase> m_queue;
diff --git a/tools/gfx/open-gl/render-gl.cpp b/tools/gfx/open-gl/render-gl.cpp
index cd4811911..524c28858 100644
--- a/tools/gfx/open-gl/render-gl.cpp
+++ b/tools/gfx/open-gl/render-gl.cpp
@@ -170,6 +170,18 @@ public:
virtual void dispatchCompute(int x, int y, int z) override;
virtual void submitGpuWork() override {}
virtual void waitForGpu() override {}
+ virtual void writeTimestamp(IQueryPool* pool, SlangInt index) override
+ {
+ SLANG_UNUSED(pool);
+ SLANG_UNUSED(index);
+ }
+ virtual SLANG_NO_THROW Result SLANG_MCALL createQueryPool(
+ const IQueryPool::Desc& desc, IQueryPool** pool) override
+ {
+ SLANG_UNUSED(desc);
+ *pool = nullptr;
+ return SLANG_E_NOT_IMPLEMENTED;
+ }
virtual SLANG_NO_THROW const DeviceInfo& SLANG_MCALL getDeviceInfo() const override
{
return m_info;
diff --git a/tools/gfx/renderer-shared.cpp b/tools/gfx/renderer-shared.cpp
index 1b384f8eb..9d6c85309 100644
--- a/tools/gfx/renderer-shared.cpp
+++ b/tools/gfx/renderer-shared.cpp
@@ -31,6 +31,8 @@ const Slang::Guid GfxGUID::IID_IComputeCommandEncoder = SLANG_UUID_IComputeComma
const Slang::Guid GfxGUID::IID_IResourceCommandEncoder = SLANG_UUID_IResourceCommandEncoder;
const Slang::Guid GfxGUID::IID_ICommandBuffer = SLANG_UUID_ICommandBuffer;
const Slang::Guid GfxGUID::IID_ICommandQueue = SLANG_UUID_ICommandQueue;
+const Slang::Guid GfxGUID::IID_IQueryPool = SLANG_UUID_IQueryPool;
+
StageType translateStage(SlangStage slangStage)
{
diff --git a/tools/gfx/renderer-shared.h b/tools/gfx/renderer-shared.h
index 0dc0f75ae..127987726 100644
--- a/tools/gfx/renderer-shared.h
+++ b/tools/gfx/renderer-shared.h
@@ -35,6 +35,7 @@ struct GfxGUID
static const Slang::Guid IID_IResourceCommandEncoder;
static const Slang::Guid IID_ICommandBuffer;
static const Slang::Guid IID_ICommandQueue;
+ static const Slang::Guid IID_IQueryPool;
};
// We use a `BreakableReference` to avoid the cyclic reference situation in gfx implementation.
diff --git a/tools/gfx/vulkan/render-vk.cpp b/tools/gfx/vulkan/render-vk.cpp
index d040fcebf..78054b4d5 100644
--- a/tools/gfx/vulkan/render-vk.cpp
+++ b/tools/gfx/vulkan/render-vk.cpp
@@ -104,6 +104,9 @@ public:
virtual SLANG_NO_THROW Result SLANG_MCALL createComputePipelineState(
const ComputePipelineStateDesc& desc,
IPipelineState** outState) override;
+ virtual SLANG_NO_THROW Result SLANG_MCALL createQueryPool(
+ const IQueryPool::Desc& desc,
+ IQueryPool** outPool) override;
virtual SLANG_NO_THROW SlangResult SLANG_MCALL readTextureResource(
ITextureResource* texture,
@@ -3498,7 +3501,7 @@ public:
VkCommandBufferBeginInfo beginInfo = {
VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO,
nullptr,
- VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT};
+ VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT };
api.vkBeginCommandBuffer(m_commandBuffer, &beginInfo);
if (m_preCommandBuffer)
{
@@ -3520,7 +3523,7 @@ public:
VkCommandBufferBeginInfo beginInfo = {
VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO,
nullptr,
- VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT};
+ VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT };
api.vkBeginCommandBuffer(m_preCommandBuffer, &beginInfo);
return SLANG_OK;
}
@@ -3534,6 +3537,20 @@ public:
return m_preCommandBuffer;
}
+ static void _writeTimestamp(
+ VulkanApi* api,
+ VkCommandBuffer vkCmdBuffer,
+ IQueryPool* queryPool,
+ SlangInt index)
+ {
+ auto queryPoolImpl = static_cast<QueryPoolImpl*>(queryPool);
+ api->vkCmdResetQueryPool(vkCmdBuffer, queryPoolImpl->m_pool, (uint32_t)index, 1);
+ api->vkCmdWriteTimestamp(vkCmdBuffer,
+ VK_PIPELINE_STAGE_BOTTOM_OF_PIPE_BIT,
+ queryPoolImpl->m_pool,
+ (uint32_t)index);
+ }
+
public:
class RenderCommandEncoder
: public IRenderCommandEncoder
@@ -3592,6 +3609,11 @@ public:
endEncodingImpl();
}
+ virtual SLANG_NO_THROW void SLANG_MCALL writeTimestamp(IQueryPool* queryPool, SlangInt index) override
+ {
+ _writeTimestamp(m_api, m_vkCommandBuffer, queryPool, index);
+ }
+
virtual SLANG_NO_THROW Result SLANG_MCALL
bindPipeline(IPipelineState* pipelineState, IShaderObject** outRootObject) override
{
@@ -3843,6 +3865,11 @@ public:
flushBindingState(VK_PIPELINE_BIND_POINT_COMPUTE);
m_api->vkCmdDispatch(m_vkCommandBuffer, x, y, z);
}
+
+ virtual SLANG_NO_THROW void SLANG_MCALL writeTimestamp(IQueryPool* queryPool, SlangInt index) override
+ {
+ _writeTimestamp(m_api, m_vkCommandBuffer, queryPool, index);
+ }
};
RefPtr<ComputeCommandEncoder> m_computeCommandEncoder;
@@ -3940,6 +3967,15 @@ public:
nullptr);
}
+ virtual SLANG_NO_THROW void SLANG_MCALL writeTimestamp(IQueryPool* queryPool, SlangInt index) override
+ {
+ _writeTimestamp(
+ &m_commandBuffer->m_renderer->m_api,
+ m_commandBuffer->m_commandBuffer,
+ queryPool,
+ index);
+ }
+
void init(CommandBufferImpl* commandBuffer)
{
m_commandBuffer = commandBuffer;
@@ -4120,6 +4156,59 @@ public:
virtual SLANG_NO_THROW Result SLANG_MCALL synchronizeAndReset() override;
};
+ class QueryPoolImpl
+ : public IQueryPool
+ , public ComObject
+ {
+ public:
+ SLANG_COM_OBJECT_IUNKNOWN_ALL
+ IQueryPool* getInterface(const Guid& guid)
+ {
+ if (guid == GfxGUID::IID_ISlangUnknown || guid == GfxGUID::IID_IQueryPool)
+ return static_cast<IQueryPool*>(this);
+ return nullptr;
+ }
+ public:
+ Result init(const IQueryPool::Desc& desc, VKDevice* device)
+ {
+ m_device = device;
+ VkQueryPoolCreateInfo createInfo = {};
+ createInfo.sType = VK_STRUCTURE_TYPE_QUERY_POOL_CREATE_INFO;
+ createInfo.queryCount = (uint32_t)desc.count;
+ switch (desc.type)
+ {
+ case QueryType::Timestamp:
+ createInfo.queryType = VK_QUERY_TYPE_TIMESTAMP;
+ break;
+ default:
+ return SLANG_E_INVALID_ARG;
+ }
+ SLANG_VK_RETURN_ON_FAIL(m_device->m_api.vkCreateQueryPool(
+ m_device->m_api.m_device, &createInfo, nullptr, &m_pool));
+ return SLANG_OK;
+ }
+ ~QueryPoolImpl()
+ {
+ m_device->m_api.vkDestroyQueryPool(m_device->m_api.m_device, m_pool, nullptr);
+ }
+ public:
+ virtual SLANG_NO_THROW Result SLANG_MCALL getResult(SlangInt index, SlangInt count, uint64_t* data) override
+ {
+ SLANG_VK_RETURN_ON_FAIL(m_device->m_api.vkGetQueryPoolResults(
+ m_device->m_api.m_device,
+ m_pool,
+ (uint32_t)index,
+ (uint32_t)count,
+ sizeof(uint64_t) * count,
+ data,
+ sizeof(uint64_t), 0));
+ return SLANG_OK;
+ }
+ public:
+ VkQueryPool m_pool;
+ RefPtr<VKDevice> m_device;
+ };
+
class SwapchainImpl
: public ISwapchain
, public ComObject
@@ -4935,6 +5024,9 @@ Result VKDevice::initVulkanInstanceAndDevice(bool useValidationLayer)
VkPhysicalDeviceProperties basicProps = {};
m_api.vkGetPhysicalDeviceProperties(m_api.m_physicalDevice, &basicProps);
+ // Compute timestamp frequency.
+ m_info.timestampFrequency = uint64_t(1e9 / basicProps.limits.timestampPeriod);
+
// Get the API version
const uint32_t majorVersion = VK_VERSION_MAJOR(basicProps.apiVersion);
const uint32_t minorVersion = VK_VERSION_MINOR(basicProps.apiVersion);
@@ -6473,4 +6565,14 @@ Result VKDevice::createComputePipelineState(const ComputePipelineStateDesc& inDe
return SLANG_OK;
}
+Result VKDevice::createQueryPool(
+ const IQueryPool::Desc& desc,
+ IQueryPool** outPool)
+{
+ RefPtr<QueryPoolImpl> result = new QueryPoolImpl();
+ SLANG_RETURN_ON_FAIL(result->init(desc, this));
+ returnComPtr(outPool, result);
+ return SLANG_OK;
+}
+
} // renderer_test
diff --git a/tools/gfx/vulkan/vk-api.h b/tools/gfx/vulkan/vk-api.h
index c6537ff8f..75b88dd96 100644
--- a/tools/gfx/vulkan/vk-api.h
+++ b/tools/gfx/vulkan/vk-api.h
@@ -91,6 +91,10 @@ namespace gfx {
x(vkCmdCopyBufferToImage)\
x(vkCmdPushConstants) \
x(vkCmdSetStencilReference) \
+ x(vkCmdWriteTimestamp) \
+ x(vkCmdBeginQuery) \
+ x(vkCmdEndQuery) \
+ x(vkCmdResetQueryPool) \
\
x(vkCreateFence) \
x(vkDestroyFence) \
@@ -116,6 +120,10 @@ namespace gfx {
\
x(vkBindImageMemory) \
x(vkBindBufferMemory) \
+ \
+ x(vkCreateQueryPool) \
+ x(vkGetQueryPoolResults) \
+ x(vkDestroyQueryPool) \
/* */
#if SLANG_WINDOWS_FAMILY
generated by cgit v1.2.3 (git 2.39.1) at 2026-06-01 21:30:17 +0000