[gfx] Add inline ray tracing support. (#1899)

7 files changed, 811 insertions, 5 deletions

diff --git a/examples/example-base/example-base.cpp b/examples/example-base/example-base.cpp
index 04f938697..c45e3cc27 100644
--- a/examples/example-base/example-base.cpp
+++ b/examples/example-base/example-base.cpp
@@ -9,7 +9,11 @@
 using namespace Slang;
 using namespace gfx;
 
-Slang::Result WindowedAppBase::initializeBase(const char* title, int width, int height)
+Slang::Result WindowedAppBase::initializeBase(
+    const char* title,
+    int width,
+    int height,
+    DeviceType deviceType)
 {
     // Create a window for our application to render into.
     //
@@ -30,7 +34,7 @@ Slang::Result WindowedAppBase::initializeBase(const char* title, int width, int
     gfxEnableDebugLayer();
 #endif
     IDevice::Desc deviceDesc = {};
-    // deviceDesc.slang.targetFlags = SLANG_TARGET_FLAG_DUMP_IR;
+    deviceDesc.deviceType = deviceType;
     gfx::Result res = gfxCreateDevice(&deviceDesc, gDevice.writeRef());
     if (SLANG_FAILED(res))
         return res;
diff --git a/examples/example-base/example-base.h b/examples/example-base/example-base.h
index a1c46aa53..b898f7417 100644
--- a/examples/example-base/example-base.h
+++ b/examples/example-base/example-base.h
@@ -26,7 +26,11 @@ protected:
     Slang::ComPtr<gfx::IRenderPassLayout> gRenderPass;
     Slang::ComPtr<gfx::ICommandQueue> gQueue;
 
-    Slang::Result initializeBase(const char* titile, int width, int height);
+    Slang::Result initializeBase(
+        const char* titile,
+        int width,
+        int height,
+        gfx::DeviceType deviceType = gfx::DeviceType::Default);
     void createSwapchainFramebuffers();
     void mainLoop();
 
diff --git a/examples/gpu-printing/main.cpp b/examples/gpu-printing/main.cpp
index 11a1553a3..03ab5b51f 100644
--- a/examples/gpu-printing/main.cpp
+++ b/examples/gpu-printing/main.cpp
@@ -109,7 +109,7 @@ Result execute()
     printBufferDesc.defaultState = ResourceState::UnorderedAccess;
     printBufferDesc.allowedStates = ResourceStateSet(
         ResourceState::CopySource, ResourceState::CopyDestination, ResourceState::UnorderedAccess);
-    printBufferDesc.cpuAccessFlags = IResource::AccessFlag::Read; // | Resource::AccessFlag::Write;
+    printBufferDesc.cpuAccessFlags = AccessFlag::Read; // | Resource::AccessFlag::Write;
     auto printBuffer = gDevice->createBufferResource(printBufferDesc);
 
     IResourceView::Desc printBufferViewDesc;
diff --git a/examples/ray-tracing/README.md b/examples/ray-tracing/README.md
new file mode 100644
index 000000000..fd7cc1cec
--- /dev/null
+++ b/examples/ray-tracing/README.md
@@ -0,0 +1,9 @@
+Slang "Ray Tracing" Example
+===========================
+
+The goal of this example is to demonstrate how to use hardware ray-tracing in Slang.
+
+The `shaders.slang` file contains a compute shader that traces primary rays from camera and shade intersections with basic lighting + ray-traced shadows. The file also defines a vertex and a fragment shader entry point for displaying the ray-traced image produced by the compute shader.
+
+The `main.cpp` file contains the C++ application code, showing how to use the Slang API to load and compile the shader code, and how to use a graphics API abstraction layer implemented in `tools/gfx` to initiate hardware ray-tracing.
+Note that this abstraction layer is *not* required in order to work with Slang, and it is just there to help us write example and test applications more conveniently.
\ No newline at end of file
diff --git a/examples/ray-tracing/main.cpp b/examples/ray-tracing/main.cpp
new file mode 100644
index 000000000..a5093ad18
--- /dev/null
+++ b/examples/ray-tracing/main.cpp
@@ -0,0 +1,645 @@
+// main.cpp
+
+// This file implements an example of hardware ray-tracing using
+// Slang shaders and the `gfx` graphics API.
+
+#include <slang.h>
+#include "slang-gfx.h"
+#include "gfx-util/shader-cursor.h"
+#include "tools/platform/window.h"
+#include "tools/platform/vector-math.h"
+#include "slang-com-ptr.h"
+#include "source/core/slang-basic.h"
+#include "examples/example-base/example-base.h"
+
+using namespace gfx;
+using namespace Slang;
+
+struct Uniforms
+{
+    float screenWidth, screenHeight;
+    float focalLength = 24.0f, frameHeight = 24.0f;
+    float cameraDir[4];
+    float cameraUp[4];
+    float cameraRight[4];
+    float cameraPosition[4];
+    float lightDir[4];
+};
+
+struct Vertex
+{
+    float position[3];
+};
+
+// Define geometry data for our test scene.
+// The scene contains a floor plane, and a cube placed on top of it at the center.
+static const int kVertexCount = 24;
+static const Vertex kVertexData[kVertexCount] =
+{
+    // Floor plane
+    {{-100.0f, 0, 100.0f}},
+    {{100.0f, 0, 100.0f}},
+    {{100.0f, 0, -100.0f}},
+    {{-100.0f, 0, -100.0f}},
+    // Cube face (+y).
+    {{-1.0f, 2.0, 1.0f}},
+    {{1.0f, 2.0, 1.0f}},
+    {{1.0f, 2.0, -1.0f}},
+    {{-1.0f, 2.0, -1.0f}},
+    // Cube face (+z).
+    {{-1.0f, 0.0, 1.0f}},
+    {{1.0f, 0.0, 1.0f}},
+    {{1.0f, 2.0, 1.0f}},
+    {{-1.0f, 2.0, 1.0f}},
+    // Cube face (-z).
+    {{-1.0f, 0.0, -1.0f}},
+    {{-1.0f, 2.0, -1.0f}},
+    {{1.0f, 2.0, -1.0f}},
+    {{1.0f, 0.0, -1.0f}},
+    // Cube face (-x).
+    {{-1.0f, 0.0, -1.0f}},
+    {{-1.0f, 0.0, 1.0f}},
+    {{-1.0f, 2.0, 1.0f}},
+    {{-1.0f, 2.0, -1.0f}},
+    // Cube face (+x).
+    {{1.0f, 2.0, -1.0f}},
+    {{1.0f, 2.0, 1.0f}},
+    {{1.0f, 0.0, 1.0f}},
+    {{1.0f, 0.0, -1.0f}},
+};
+static const int kIndexCount = 36;
+static const int kIndexData[kIndexCount] =
+{
+    0, 1, 2, 0, 2, 3,
+    4, 5, 6, 4, 6, 7,
+    8, 9, 10, 8, 10, 11,
+    12, 13, 14, 12, 14, 15,
+    16, 17, 18, 16, 18, 19,
+    20, 21, 22, 20, 22, 23
+};
+
+struct Primitive
+{
+    float data[4];
+    float color[4];
+};
+static const int kPrimitiveCount = 12;
+static const Primitive kPrimitiveData[kPrimitiveCount] =
+{
+    {{0.0f, 1.0f, 0.0f, 0.0f}, {0.75f, 0.8f, 0.85f, 1.0f}},
+    {{0.0f, 1.0f, 0.0f, 0.0f}, {0.75f, 0.8f, 0.85f, 1.0f}},
+    {{0.0f, 1.0f, 0.0f, 0.0f}, {0.95f, 0.85f, 0.05f, 1.0f}},
+    {{0.0f, 1.0f, 0.0f, 0.0f}, {0.95f, 0.85f, 0.05f, 1.0f}},
+    {{0.0f, 0.0f, 1.0f, 0.0f}, {0.95f, 0.85f, 0.05f, 1.0f}},
+    {{0.0f, 0.0f, 1.0f, 0.0f}, {0.95f, 0.85f, 0.05f, 1.0f}},
+    {{0.0f, 0.0f, -1.0f, 0.0f}, {0.95f, 0.85f, 0.05f, 1.0f}},
+    {{0.0f, 0.0f, -1.0f, 0.0f}, {0.95f, 0.85f, 0.05f, 1.0f}},
+    {{-1.0f, 0.0f, 0.0f, 0.0f}, {0.95f, 0.85f, 0.05f, 1.0f}},
+    {{-1.0f, 0.0f, 0.0f, 0.0f}, {0.95f, 0.85f, 0.05f, 1.0f}},
+    {{1.0f, 0.0f, 0.0f, 0.0f}, {0.95f, 0.85f, 0.05f, 1.0f}},
+    {{1.0f, 0.0f, 0.0f, 0.0f}, {0.95f, 0.85f, 0.05f, 1.0f}},
+};
+
+
+// We need to use a rasterization pipeline to copy the ray-traced image
+// to the swapchain. To do so we need to render a full-screen triangle.
+// We will define a small helper type that defines the data for such a triangle.
+//
+struct FullScreenTriangle
+{
+    struct Vertex
+    {
+        float position[2];
+    };
+
+    enum
+    {
+        kVertexCount = 3
+    };
+
+    static const Vertex kVertices[kVertexCount];
+};
+const FullScreenTriangle::Vertex FullScreenTriangle::kVertices[FullScreenTriangle::kVertexCount] = {
+    {{-1, -1}},
+    {{-1, 3}},
+    {{3, -1}},
+};
+
+// The example application will be implemented as a `struct`, so that
+// we can scope the resources it allocates without using global variables.
+//
+struct RayTracing : public WindowedAppBase
+{
+
+
+Uniforms gUniforms = {};
+
+
+// Many Slang API functions return detailed diagnostic information
+// (error messages, warnings, etc.) as a "blob" of data, or return
+// a null blob pointer instead if there were no issues.
+//
+// For convenience, we define a subroutine that will dump the information
+// in a diagnostic blob if one is produced, and skip it otherwise.
+//
+void diagnoseIfNeeded(slang::IBlob* diagnosticsBlob)
+{
+    if( diagnosticsBlob != nullptr )
+    {
+        printf("%s", (const char*) diagnosticsBlob->getBufferPointer());
+#ifdef _WIN32
+        _Win32OutputDebugString((const char*)diagnosticsBlob->getBufferPointer());
+#endif
+    }
+}
+
+// Load and compile shader code from souce.
+gfx::Result loadShaderProgram(
+    gfx::IDevice*         device,
+    gfx::PipelineType pipelineType,
+    gfx::IShaderProgram**   outProgram)
+{
+    ComPtr<slang::ISession> slangSession;
+    slangSession = device->getSlangSession();
+
+    ComPtr<slang::IBlob> diagnosticsBlob;
+    slang::IModule* module = slangSession->loadModule("shaders", diagnosticsBlob.writeRef());
+    diagnoseIfNeeded(diagnosticsBlob);
+    if(!module)
+        return SLANG_FAIL;
+
+    Slang::List<slang::IComponentType*> componentTypes;
+    componentTypes.add(module);
+    if (pipelineType == PipelineType::Compute)
+    {
+        ComPtr<slang::IEntryPoint> computeEntryPoint;
+        SLANG_RETURN_ON_FAIL(module->findEntryPointByName("computeMain", computeEntryPoint.writeRef()));
+        componentTypes.add(computeEntryPoint);
+    }
+    else
+    {
+        ComPtr<slang::IEntryPoint> entryPoint;
+        SLANG_RETURN_ON_FAIL(module->findEntryPointByName("vertexMain", entryPoint.writeRef()));
+        componentTypes.add(entryPoint);
+        SLANG_RETURN_ON_FAIL(module->findEntryPointByName("fragmentMain", entryPoint.writeRef()));
+        componentTypes.add(entryPoint);
+    }
+
+    ComPtr<slang::IComponentType> linkedProgram;
+    SlangResult result = slangSession->createCompositeComponentType(
+        componentTypes.getBuffer(),
+        componentTypes.getCount(),
+        linkedProgram.writeRef(),
+        diagnosticsBlob.writeRef());
+    diagnoseIfNeeded(diagnosticsBlob);
+    SLANG_RETURN_ON_FAIL(result);
+
+    gfx::IShaderProgram::Desc programDesc = {};
+    programDesc.pipelineType = pipelineType;
+    programDesc.slangProgram = linkedProgram;
+    SLANG_RETURN_ON_FAIL(device->createProgram(programDesc, outProgram));
+
+    return SLANG_OK;
+}
+
+ComPtr<gfx::IPipelineState> gPresentPipelineState;
+ComPtr<gfx::IPipelineState> gRenderPipelineState;
+ComPtr<gfx::IBufferResource> gFullScreenVertexBuffer;
+ComPtr<gfx::IBufferResource> gVertexBuffer;
+ComPtr<gfx::IBufferResource> gIndexBuffer;
+ComPtr<gfx::IBufferResource> gPrimitiveBuffer;
+ComPtr<gfx::IBufferResource> gTransformBuffer;
+ComPtr<gfx::IResourceView> gPrimitiveBufferSRV;
+ComPtr<gfx::IBufferResource> gInstanceBuffer;
+ComPtr<gfx::IBufferResource> gBLASBuffer;
+ComPtr<gfx::IAccelerationStructure> gBLAS;
+ComPtr<gfx::IBufferResource> gTLASBuffer;
+ComPtr<gfx::IAccelerationStructure> gTLAS;
+ComPtr<gfx::ITextureResource> gResultTexture;
+ComPtr<gfx::IResourceView> gResultTextureUAV;
+
+uint64_t lastTime = 0;
+
+// glm::vec3 lightDir = normalize(glm::vec3(10, 10, 10));
+// glm::vec3 lightColor = glm::vec3(1, 1, 1);
+
+glm::vec3 cameraPosition = glm::vec3(-2.53f, 2.72f, 4.3f);
+float cameraOrientationAngles[2] = {-0.475f, -0.35f}; // Spherical angles (theta, phi).
+
+float translationScale = 0.5f;
+float rotationScale = 0.01f;
+
+// In order to control camera movement, we will
+// use good old WASD
+bool wPressed = false;
+bool aPressed = false;
+bool sPressed = false;
+bool dPressed = false;
+
+bool isMouseDown = false;
+float lastMouseX = 0.0f;
+float lastMouseY = 0.0f;
+
+void setKeyState(platform::KeyCode key, bool state)
+{
+    switch (key)
+    {
+    default:
+        break;
+    case platform::KeyCode::W:
+        wPressed = state;
+        break;
+    case platform::KeyCode::A:
+        aPressed = state;
+        break;
+    case platform::KeyCode::S:
+        sPressed = state;
+        break;
+    case platform::KeyCode::D:
+        dPressed = state;
+        break;
+    }
+}
+void onKeyDown(platform::KeyEventArgs args) { setKeyState(args.key, true); }
+void onKeyUp(platform::KeyEventArgs args) { setKeyState(args.key, false); }
+
+void onMouseDown(platform::MouseEventArgs args)
+{
+    isMouseDown = true;
+    lastMouseX = (float)args.x;
+    lastMouseY = (float)args.y;
+}
+
+void onMouseMove(platform::MouseEventArgs args)
+{
+    if (isMouseDown)
+    {
+        float deltaX = args.x - lastMouseX;
+        float deltaY = args.y - lastMouseY;
+
+        cameraOrientationAngles[0] += -deltaX * rotationScale;
+        cameraOrientationAngles[1] += -deltaY * rotationScale;
+        lastMouseX = (float)args.x;
+        lastMouseY = (float)args.y;
+    }
+}
+void onMouseUp(platform::MouseEventArgs args) { isMouseDown = false; }
+
+Slang::Result initialize()
+{
+    initializeBase("Ray Tracing", 1024, 768, DeviceType::Vulkan);
+    gWindow->events.mouseMove = [this](const platform::MouseEventArgs& e) { onMouseMove(e); };
+    gWindow->events.mouseUp = [this](const platform::MouseEventArgs& e) { onMouseUp(e); };
+    gWindow->events.mouseDown = [this](const platform::MouseEventArgs& e) { onMouseDown(e); };
+    gWindow->events.keyDown = [this](const platform::KeyEventArgs& e) { onKeyDown(e); };
+    gWindow->events.keyUp = [this](const platform::KeyEventArgs& e) { onKeyUp(e); };
+
+    IBufferResource::Desc vertexBufferDesc;
+    vertexBufferDesc.type = IResource::Type::Buffer;
+    vertexBufferDesc.sizeInBytes = kVertexCount * sizeof(Vertex);
+    vertexBufferDesc.defaultState = ResourceState::UnorderedAccess;
+    gVertexBuffer = gDevice->createBufferResource(vertexBufferDesc, &kVertexData[0]);
+    if(!gVertexBuffer) return SLANG_FAIL;
+
+    IBufferResource::Desc indexBufferDesc;
+    indexBufferDesc.type = IResource::Type::Buffer;
+    indexBufferDesc.sizeInBytes = kIndexCount * sizeof(int32_t);
+    indexBufferDesc.defaultState = ResourceState::UnorderedAccess;
+    gIndexBuffer = gDevice->createBufferResource(indexBufferDesc, &kIndexData[0]);
+    if (!gIndexBuffer)
+        return SLANG_FAIL;
+
+    IBufferResource::Desc primitiveBufferDesc;
+    primitiveBufferDesc.type = IResource::Type::Buffer;
+    primitiveBufferDesc.sizeInBytes = kPrimitiveCount * sizeof(Primitive);
+    primitiveBufferDesc.defaultState = ResourceState::ShaderResource;
+    gPrimitiveBuffer = gDevice->createBufferResource(primitiveBufferDesc, &kPrimitiveData[0]);
+    if (!gPrimitiveBuffer)
+        return SLANG_FAIL;
+
+    IResourceView::Desc primitiveSRVDesc = {};
+    primitiveSRVDesc.format = Format::Unknown;
+    primitiveSRVDesc.type = IResourceView::Type::ShaderResource;
+    gPrimitiveBufferSRV = gDevice->createBufferView(gPrimitiveBuffer, primitiveSRVDesc);
+
+    IBufferResource::Desc transformBufferDesc;
+    transformBufferDesc.type = IResource::Type::Buffer;
+    transformBufferDesc.sizeInBytes = sizeof(float) * 16;
+    transformBufferDesc.defaultState = ResourceState::UnorderedAccess;
+    float transformData[12] = {
+        1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f};
+    gTransformBuffer = gDevice->createBufferResource(transformBufferDesc, &transformData);
+    if (!gTransformBuffer)
+        return SLANG_FAIL;
+    // Build bottom level acceleration structure.
+    {
+        IAccelerationStructure::BuildInputs accelerationStructureBuildInputs;
+        IAccelerationStructure::PrebuildInfo accelerationStructurePrebuildInfo;
+        accelerationStructureBuildInputs.descCount = 1;
+        accelerationStructureBuildInputs.kind = IAccelerationStructure::Kind::BottomLevel;
+        accelerationStructureBuildInputs.flags =
+            IAccelerationStructure::BuildFlags::AllowCompaction;
+        IAccelerationStructure::GeometryDesc geomDesc;
+        geomDesc.flags = IAccelerationStructure::GeometryFlags::Opaque;
+        geomDesc.type = IAccelerationStructure::GeometryType::Triangles;
+        geomDesc.content.triangles.indexCount = kIndexCount;
+        geomDesc.content.triangles.indexData = gIndexBuffer->getDeviceAddress();
+        geomDesc.content.triangles.indexFormat = Format::R_UInt32;
+        geomDesc.content.triangles.vertexCount = kVertexCount;
+        geomDesc.content.triangles.vertexData = gVertexBuffer->getDeviceAddress();
+        geomDesc.content.triangles.vertexFormat = Format::RGB_Float32;
+        geomDesc.content.triangles.vertexStride = sizeof(Vertex);
+        geomDesc.content.triangles.transform3x4 = gTransformBuffer->getDeviceAddress();
+        accelerationStructureBuildInputs.geometryDescs = &geomDesc;
+
+        // Query buffer size for acceleration structure build.
+        SLANG_RETURN_ON_FAIL(gDevice->getAccelerationStructurePrebuildInfo(
+            accelerationStructureBuildInputs, &accelerationStructurePrebuildInfo));
+        // Allocate buffers for acceleration structure.
+        IBufferResource::Desc asDraftBufferDesc;
+        asDraftBufferDesc.type = IResource::Type::Buffer;
+        asDraftBufferDesc.defaultState = ResourceState::AccelerationStructure;
+        asDraftBufferDesc.sizeInBytes = accelerationStructurePrebuildInfo.resultDataMaxSize;
+        ComPtr<IBufferResource> draftBuffer = gDevice->createBufferResource(asDraftBufferDesc);
+        IBufferResource::Desc scratchBufferDesc;
+        scratchBufferDesc.type = IResource::Type::Buffer;
+        scratchBufferDesc.defaultState = ResourceState::UnorderedAccess;
+        scratchBufferDesc.sizeInBytes = accelerationStructurePrebuildInfo.scratchDataSize;
+        ComPtr<IBufferResource> scratchBuffer = gDevice->createBufferResource(scratchBufferDesc);
+
+        // Build acceleration structure.
+        ComPtr<IQueryPool> compactedSizeQuery;
+        IQueryPool::Desc queryPoolDesc;
+        queryPoolDesc.count = 1;
+        queryPoolDesc.type = QueryType::AccelerationStructureCompactedSize;
+        SLANG_RETURN_ON_FAIL(
+            gDevice->createQueryPool(queryPoolDesc, compactedSizeQuery.writeRef()));
+
+        ComPtr<IAccelerationStructure> draftAS;
+        IAccelerationStructure::CreateDesc draftCreateDesc;
+        draftCreateDesc.buffer = draftBuffer;
+        draftCreateDesc.kind = IAccelerationStructure::Kind::BottomLevel;
+        draftCreateDesc.offset = 0;
+        draftCreateDesc.size = accelerationStructurePrebuildInfo.resultDataMaxSize;
+        SLANG_RETURN_ON_FAIL(
+            gDevice->createAccelerationStructure(draftCreateDesc, draftAS.writeRef()));
+
+        auto commandBuffer = gTransientHeaps[0]->createCommandBuffer();
+        auto encoder = commandBuffer->encodeRayTracingCommands();
+        IAccelerationStructure::BuildDesc buildDesc = {};
+        buildDesc.dest = draftAS;
+        buildDesc.inputs = accelerationStructureBuildInputs;
+        buildDesc.scratchData = scratchBuffer->getDeviceAddress();
+        AccelerationStructureQueryDesc compactedSizeQueryDesc = {};
+        compactedSizeQueryDesc.queryPool = compactedSizeQuery;
+        compactedSizeQueryDesc.queryType = QueryType::AccelerationStructureCompactedSize;
+        encoder->buildAccelerationStructure(buildDesc, 1, &compactedSizeQueryDesc);
+        encoder->endEncoding();
+        commandBuffer->close();
+        gQueue->executeCommandBuffer(commandBuffer);
+        gQueue->wait();
+
+        uint64_t compactedSize = 0;
+        compactedSizeQuery->getResult(0, 1, &compactedSize);
+        IBufferResource::Desc asBufferDesc;
+        asBufferDesc.type = IResource::Type::Buffer;
+        asBufferDesc.defaultState = ResourceState::AccelerationStructure;
+        asBufferDesc.sizeInBytes = compactedSize;
+        gBLASBuffer = gDevice->createBufferResource(asBufferDesc);
+        IAccelerationStructure::CreateDesc createDesc;
+        createDesc.buffer = gBLASBuffer;
+        createDesc.kind = IAccelerationStructure::Kind::BottomLevel;
+        createDesc.offset = 0;
+        createDesc.size = compactedSize;
+        gDevice->createAccelerationStructure(createDesc, gBLAS.writeRef());
+
+        commandBuffer = gTransientHeaps[0]->createCommandBuffer();
+        encoder = commandBuffer->encodeRayTracingCommands();
+        encoder->copyAccelerationStructure(gBLAS, draftAS, AccelerationStructureCopyMode::Compact);
+        encoder->endEncoding();
+        commandBuffer->close();
+        gQueue->executeCommandBuffer(commandBuffer);
+        gQueue->wait();
+    }
+
+    // Build top level acceleration structure.
+    {
+        List<IAccelerationStructure::InstanceDesc> instanceDescs;
+        instanceDescs.setCount(1);
+        instanceDescs[0].accelerationStructure = gBLAS->getDeviceAddress();
+        instanceDescs[0].flags =
+            IAccelerationStructure::GeometryInstanceFlags::TriangleFacingCullDisable;
+        instanceDescs[0].instanceContributionToHitGroupIndex = 0;
+        instanceDescs[0].instanceID = 0;
+        instanceDescs[0].instanceMask = 0xFF;
+        float transformMatrix[] = {1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f};
+        memcpy(&instanceDescs[0].transform[0][0], transformMatrix, sizeof(float) * 12);
+
+        IBufferResource::Desc instanceBufferDesc;
+        instanceBufferDesc.type = IResource::Type::Buffer;
+        instanceBufferDesc.sizeInBytes =
+            instanceDescs.getCount() * sizeof(IAccelerationStructure::InstanceDesc);
+        instanceBufferDesc.defaultState = ResourceState::UnorderedAccess;
+        gInstanceBuffer = gDevice->createBufferResource(instanceBufferDesc, instanceDescs.getBuffer());
+        if (!gInstanceBuffer)
+            return SLANG_FAIL;
+
+        IAccelerationStructure::BuildInputs accelerationStructureBuildInputs = {};
+        IAccelerationStructure::PrebuildInfo accelerationStructurePrebuildInfo = {};
+        accelerationStructureBuildInputs.descCount = 1;
+        accelerationStructureBuildInputs.kind = IAccelerationStructure::Kind::TopLevel;
+        accelerationStructureBuildInputs.instanceDescs = gInstanceBuffer->getDeviceAddress();
+        
+        // Query buffer size for acceleration structure build.
+        SLANG_RETURN_ON_FAIL(gDevice->getAccelerationStructurePrebuildInfo(
+            accelerationStructureBuildInputs, &accelerationStructurePrebuildInfo));
+
+        IBufferResource::Desc asBufferDesc;
+        asBufferDesc.type = IResource::Type::Buffer;
+        asBufferDesc.defaultState = ResourceState::AccelerationStructure;
+        asBufferDesc.sizeInBytes = accelerationStructurePrebuildInfo.resultDataMaxSize;
+        gTLASBuffer = gDevice->createBufferResource(asBufferDesc);
+
+        IBufferResource::Desc scratchBufferDesc;
+        scratchBufferDesc.type = IResource::Type::Buffer;
+        scratchBufferDesc.defaultState = ResourceState::UnorderedAccess;
+        scratchBufferDesc.sizeInBytes = accelerationStructurePrebuildInfo.scratchDataSize;
+        ComPtr<IBufferResource> scratchBuffer = gDevice->createBufferResource(scratchBufferDesc);
+
+        IAccelerationStructure::CreateDesc createDesc;
+        createDesc.buffer = gTLASBuffer;
+        createDesc.kind = IAccelerationStructure::Kind::TopLevel;
+        createDesc.offset = 0;
+        createDesc.size = accelerationStructurePrebuildInfo.resultDataMaxSize;
+        SLANG_RETURN_ON_FAIL(gDevice->createAccelerationStructure(createDesc, gTLAS.writeRef()));
+
+        auto commandBuffer = gTransientHeaps[0]->createCommandBuffer();
+        auto encoder = commandBuffer->encodeRayTracingCommands();
+        IAccelerationStructure::BuildDesc buildDesc = {};
+        buildDesc.dest = gTLAS;
+        buildDesc.inputs = accelerationStructureBuildInputs;
+        buildDesc.scratchData = scratchBuffer->getDeviceAddress();
+        encoder->buildAccelerationStructure(buildDesc, 0, nullptr);
+        encoder->endEncoding();
+        commandBuffer->close();
+        gQueue->executeCommandBuffer(commandBuffer);
+        gQueue->wait();
+    }
+
+    IBufferResource::Desc fullScreenVertexBufferDesc;
+    fullScreenVertexBufferDesc.type = IResource::Type::Buffer;
+    fullScreenVertexBufferDesc.sizeInBytes =
+        FullScreenTriangle::kVertexCount * sizeof(FullScreenTriangle::Vertex);
+    fullScreenVertexBufferDesc.defaultState = ResourceState::VertexBuffer;
+    gFullScreenVertexBuffer = gDevice->createBufferResource(
+        fullScreenVertexBufferDesc, &FullScreenTriangle::kVertices[0]);
+    if (!gFullScreenVertexBuffer)
+        return SLANG_FAIL;
+
+    InputElementDesc inputElements[] = {
+        {"POSITION", 0, Format::RG_Float32, offsetof(FullScreenTriangle::Vertex, position)},
+    };
+    auto inputLayout = gDevice->createInputLayout(&inputElements[0], SLANG_COUNT_OF(inputElements));
+    if (!inputLayout)
+        return SLANG_FAIL;
+
+    ComPtr<IShaderProgram> shaderProgram;
+    SLANG_RETURN_ON_FAIL(loadShaderProgram(gDevice, PipelineType::Graphics, shaderProgram.writeRef()));
+    GraphicsPipelineStateDesc desc;
+    desc.inputLayout = inputLayout;
+    desc.program = shaderProgram;
+    desc.framebufferLayout = gFramebufferLayout;
+    gPresentPipelineState = gDevice->createGraphicsPipelineState(desc);
+    if (!gPresentPipelineState)
+        return SLANG_FAIL;
+
+    ComPtr<IShaderProgram> computeProgram;
+    SLANG_RETURN_ON_FAIL(
+        loadShaderProgram(gDevice, PipelineType::Compute, computeProgram.writeRef()));
+    ComputePipelineStateDesc computeDesc;
+    computeDesc.program = computeProgram;
+    gRenderPipelineState = gDevice->createComputePipelineState(computeDesc);
+    if (!gRenderPipelineState)
+        return SLANG_FAIL;
+
+    createResultTexture();
+    return SLANG_OK;
+}
+
+void createResultTexture()
+{
+    ITextureResource::Desc resultTextureDesc = {};
+    resultTextureDesc.type = IResource::Type::Texture2D;
+    resultTextureDesc.numMipLevels = 1;
+    resultTextureDesc.size.width = windowWidth;
+    resultTextureDesc.size.height = windowHeight;
+    resultTextureDesc.size.depth = 1;
+    resultTextureDesc.defaultState = ResourceState::UnorderedAccess;
+    resultTextureDesc.format = Format::RGBA_Float16;
+    gResultTexture = gDevice->createTextureResource(resultTextureDesc);
+    IResourceView::Desc resultUAVDesc = {};
+    resultUAVDesc.format = resultTextureDesc.format;
+    resultUAVDesc.type = IResourceView::Type::UnorderedAccess;
+    gResultTextureUAV = gDevice->createTextureView(gResultTexture, resultUAVDesc);
+}
+
+virtual void windowSizeChanged() override
+{
+    WindowedAppBase::windowSizeChanged();
+    createResultTexture();
+}
+
+glm::vec3 getVectorFromSphericalAngles(float theta, float phi)
+{
+    auto sinTheta = sin(theta);
+    auto cosTheta = cos(theta);
+    auto sinPhi = sin(phi);
+    auto cosPhi = cos(phi);
+    return glm::vec3(-sinTheta * cosPhi, sinPhi, -cosTheta * cosPhi);
+}
+void updateUniforms()
+{
+    gUniforms.screenWidth = (float)windowWidth;
+    gUniforms.screenHeight = (float)windowHeight;
+    if (!lastTime)
+        lastTime = getCurrentTime();
+    uint64_t currentTime = getCurrentTime();
+    float deltaTime = float(double(currentTime - lastTime) / double(getTimerFrequency()));
+    lastTime = currentTime;
+
+    auto camDir =
+        getVectorFromSphericalAngles(cameraOrientationAngles[0], cameraOrientationAngles[1]);
+    auto camUp = getVectorFromSphericalAngles(
+        cameraOrientationAngles[0], cameraOrientationAngles[1] + glm::pi<float>() * 0.5f);
+    auto camRight = glm::cross(camDir, camUp);
+
+    glm::vec3 movement = glm::vec3(0);
+    if (wPressed)
+        movement += camDir;
+    if (sPressed)
+        movement -= camDir;
+    if (aPressed)
+        movement -= camRight;
+    if (dPressed)
+        movement += camRight;
+
+    cameraPosition += deltaTime * translationScale * movement;
+
+    memcpy(gUniforms.cameraDir, &camDir, sizeof(float) * 3);
+    memcpy(gUniforms.cameraUp, &camUp, sizeof(float) * 3);
+    memcpy(gUniforms.cameraRight, &camRight, sizeof(float) * 3);
+    memcpy(gUniforms.cameraPosition, &cameraPosition, sizeof(float) * 3);
+    auto lightDir = glm::normalize(glm::vec3(1.0f, 3.0f, 2.0f));
+    memcpy(gUniforms.lightDir, &lightDir, sizeof(float) * 3);
+}
+
+virtual void renderFrame(int frameBufferIndex) override
+{
+    updateUniforms();
+    {
+        ComPtr<ICommandBuffer> renderCommandBuffer =
+            gTransientHeaps[frameBufferIndex]->createCommandBuffer();
+        auto renderEncoder = renderCommandBuffer->encodeComputeCommands();
+        auto rootObject = renderEncoder->bindPipeline(gRenderPipelineState);
+        auto cursor = ShaderCursor(rootObject->getEntryPoint(0));
+        cursor["resultTexture"].setResource(gResultTextureUAV);
+        cursor["uniforms"].setData(&gUniforms, sizeof(Uniforms));
+        cursor["sceneBVH"].setResource(gTLAS);
+        cursor["primitiveBuffer"].setResource(gPrimitiveBufferSRV);
+        renderEncoder->dispatchCompute((windowWidth + 15) / 16, (windowHeight + 15) / 16, 1);
+        renderEncoder->endEncoding();
+        renderCommandBuffer->close();
+        gQueue->executeCommandBuffer(renderCommandBuffer);
+    }
+
+    {
+        ComPtr<ICommandBuffer> presentCommandBuffer =
+            gTransientHeaps[frameBufferIndex]->createCommandBuffer();
+        auto presentEncoder = presentCommandBuffer->encodeRenderCommands(
+            gRenderPass, gFramebuffers[frameBufferIndex]);
+        gfx::Viewport viewport = {};
+        viewport.maxZ = 1.0f;
+        viewport.extentX = (float)windowWidth;
+        viewport.extentY = (float)windowHeight;
+        presentEncoder->setViewportAndScissor(viewport);
+        auto rootObject = presentEncoder->bindPipeline(gPresentPipelineState);
+        auto cursor = ShaderCursor(rootObject->getEntryPoint(1));
+        cursor["t"].setResource(gResultTextureUAV);
+        presentEncoder->setVertexBuffer(
+            0, gFullScreenVertexBuffer, sizeof(FullScreenTriangle::Vertex));
+        presentEncoder->setPrimitiveTopology(PrimitiveTopology::TriangleList);
+        presentEncoder->draw(3);
+        presentEncoder->endEncoding();
+        presentCommandBuffer->close();
+        gQueue->executeCommandBuffer(presentCommandBuffer);
+    }
+    // With that, we are done drawing for one frame, and ready for the next.
+    //
+    gSwapchain->present();
+}
+
+};
+
+// This macro instantiates an appropriate main function to
+// run the application defined above.
+PLATFORM_UI_MAIN(innerMain<RayTracing>)
diff --git a/examples/ray-tracing/shaders.slang b/examples/ray-tracing/shaders.slang
new file mode 100644
index 000000000..0aff55435
--- /dev/null
+++ b/examples/ray-tracing/shaders.slang
@@ -0,0 +1,144 @@
+// shaders.slang
+
+struct Uniforms
+{
+    float screenWidth, screenHeight;
+    float focalLength, frameHeight;
+    float4 cameraDir;
+    float4 cameraUp;
+    float4 cameraRight;
+    float4 cameraPosition;
+    float4 lightDir;
+};
+
+struct Primitive
+{
+    float4 data0;
+    float4 color;
+    float3 getNormal() { return data0.xyz; }
+    float3 getColor() { return color.xyz; }
+};
+
+bool traceRayFirstHit(
+    RaytracingAccelerationStructure sceneBVH,
+    float3 rayOrigin,
+    float3 rayDir,
+    out float t,
+    out int primitiveIndex)
+{
+    RayDesc ray;
+    ray.Origin = rayOrigin;
+    ray.TMin = 0.01f;
+    ray.Direction = rayDir;
+    ray.TMax = 1e4f;
+    RayQuery<RAY_FLAG_SKIP_PROCEDURAL_PRIMITIVES |
+             RAY_FLAG_ACCEPT_FIRST_HIT_AND_END_SEARCH> q;
+    let rayFlags = RAY_FLAG_SKIP_PROCEDURAL_PRIMITIVES |
+             RAY_FLAG_ACCEPT_FIRST_HIT_AND_END_SEARCH;
+
+    q.TraceRayInline(
+        sceneBVH,
+        rayFlags,
+        0xff,
+        ray);
+    q.Proceed();
+
+    if(q.CommittedStatus() == COMMITTED_TRIANGLE_HIT)
+    {
+        t = q.CommittedRayT();
+        primitiveIndex = q.CommittedPrimitiveIndex();
+        return true;
+    }
+    return false;
+}
+
+bool traceRayNearestHit(
+    RaytracingAccelerationStructure sceneBVH,
+    float3 rayOrigin,
+    float3 rayDir,
+    out float t,
+    out int primitiveIndex)
+{
+    RayDesc ray;
+    ray.Origin = rayOrigin;
+    ray.TMin = 0.01f;
+    ray.Direction = rayDir;
+    ray.TMax = 1e4f;
+    RayQuery<RAY_FLAG_NONE> q;
+    let rayFlags = RAY_FLAG_NONE;
+
+    q.TraceRayInline(
+        sceneBVH,
+        rayFlags,
+        0xff,
+        ray);
+
+    q.Proceed();
+    if(q.CommittedStatus() == COMMITTED_TRIANGLE_HIT)
+    {
+        t = q.CommittedRayT();
+        primitiveIndex = q.CommittedPrimitiveIndex();
+        return true;
+    }
+    return false;
+}
+
+[shader("compute")]
+[numthreads(16,16,1)]
+void computeMain(
+    uint3 threadIdx : SV_DispatchThreadID,
+    uniform RWTexture2D resultTexture,
+    uniform RaytracingAccelerationStructure sceneBVH,
+    uniform StructuredBuffer<Primitive> primitiveBuffer,
+    uniform Uniforms uniforms)
+{
+    if (threadIdx.x >= (int)uniforms.screenWidth) return;
+    if (threadIdx.y >= (int)uniforms.screenHeight) return;
+
+    float frameWidth = uniforms.screenWidth / uniforms.screenHeight * uniforms.frameHeight;
+    float imageY = (threadIdx.y / uniforms.screenHeight - 0.5f) * uniforms.frameHeight;
+    float imageX = (threadIdx.x / uniforms.screenWidth - 0.5f) * frameWidth;
+    float imageZ = uniforms.focalLength;
+    float3 rayDir = normalize(uniforms.cameraDir.xyz*imageZ - uniforms.cameraUp.xyz * imageY + uniforms.cameraRight.xyz * imageX);
+
+    float4 resultColor = 0;
+
+    int primitiveIndex;
+    float intersectionT;
+    if (traceRayNearestHit(sceneBVH, uniforms.cameraPosition.xyz, rayDir, intersectionT, primitiveIndex))
+    {
+        float3 hitLocation = uniforms.cameraPosition.xyz + rayDir * intersectionT;
+        float3 shadowRayDir = uniforms.lightDir.xyz;
+        float shadow = 1.0;
+        float shadowIntersectionT;
+        int shadowPrimitiveIndex;
+        if (traceRayFirstHit(sceneBVH, hitLocation, shadowRayDir, shadowIntersectionT, shadowPrimitiveIndex))
+        {
+            shadow = 0.0f;
+        }
+        float3 normal = primitiveBuffer[primitiveIndex].getNormal();
+        float3 color = primitiveBuffer[primitiveIndex].getColor();
+        float ndotl = max(0.0, shadow * dot(normal, uniforms.lightDir.xyz));
+        float intensity = ndotl * 0.7 + 0.3;
+        resultColor = float4(color * intensity, 1.0f);
+    }
+    resultTexture[threadIdx.xy] = resultColor;
+}
+
+/// Vertex and fragment shader for displaying the final image.
+
+[shader("vertex")]
+float4 vertexMain(float2 position : POSITION)
+    : SV_Position
+{
+    return float4(position, 0.5, 1.0);
+}
+
+[shader("fragment")]
+float4 fragmentMain(
+    float4 sv_position : SV_Position,
+    uniform RWTexture2D t)
+    : SV_Target
+{
+    return t.Load(sv_position.xy);
+}
diff --git a/examples/shader-object/main.cpp b/examples/shader-object/main.cpp
index 71c5de983..8240aa63b 100644
--- a/examples/shader-object/main.cpp
+++ b/examples/shader-object/main.cpp
@@ -171,7 +171,7 @@ int main()
         ResourceState::CopyDestination,
         ResourceState::CopySource);
     bufferDesc.defaultState = ResourceState::UnorderedAccess;
-    bufferDesc.cpuAccessFlags = IResource::AccessFlag::Write | IResource::AccessFlag::Read;
+    bufferDesc.cpuAccessFlags = AccessFlag::Write | AccessFlag::Read;
 
     ComPtr<gfx::IBufferResource> numbersBuffer;
     SLANG_RETURN_ON_FAIL(device->createBufferResource(