Scarfski/bgfx
1
0
Fork 0
mirror of https://github.com/bkaradzic/bgfx.git synced 2026-11-10 18:15:32 +00:00
Code Issues Packages Projects Releases Wiki Activity

Updated Tint.

Browse Source
This commit is contained in:
Бранимир Караџић
2026-02-21 19:54:05 -08:00
committed by Branimir Karadžić
parent 651c7711b1
commit de0cdd0740
205 changed files with 39298 additions and 20377 deletions
Download Patch File Download Diff File Expand all files Collapse all files
3rdparty/dawn/include/tint/tint.h
Vendored
+13 -8
View File
@@ -34,14 +34,19 @@
// TODO(tint:88): When implementing support for an install target, all of these
// headers will need to be moved to include/tint/.
#include "src/tint/api/common/subgroup_matrix.h"
#include "src/tint/api/common/substitute_overrides_config.h"
#include "src/tint/api/common/vertex_pulling_config.h"
#include "src/tint/api/tint.h"
#include "src/tint/lang/core/type/manager.h"
#include "src/tint/lang/wgsl/inspector/inspector.h"
#include "src/tint/utils/diagnostic/formatter.h"
#include "src/tint/utils/text/styled_text.h"
#include "src/tint/api/common/binding_point.h" // IWYU pragma: export
#include "src/tint/api/common/resource_type.h" // IWYU pragma: export
#include "src/tint/api/common/subgroup_matrix.h" // IWYU pragma: export
#include "src/tint/api/common/substitute_overrides_config.h" // IWYU pragma: export
#include "src/tint/api/common/vertex_pulling_config.h" // IWYU pragma: export
#include "src/tint/api/common/workgroup_info.h" // IWYU pragma: export
#include "src/tint/api/tint.h" // IWYU pragma: export
#include "src/tint/lang/core/type/manager.h" // IWYU pragma: export
#include "src/tint/lang/wgsl/enums.h" // IWYU pragma: export
#include "src/tint/lang/wgsl/feature_status.h" // IWYU pragma: export
#include "src/tint/lang/wgsl/inspector/inspector.h" // IWYU pragma: export
#include "src/tint/utils/diagnostic/formatter.h" // IWYU pragma: export
#include "src/tint/utils/text/styled_text.h" // IWYU pragma: export
///////////////
// NOTE if adding a new guard include here, it must also appear in src/tint/api/tint.cc for the
3rdparty/dawn/include/webgpu/webgpu.h
Vendored
+59 -45
View File
@@ -124,6 +124,7 @@ typedef struct WGPUBindGroupImpl* WGPUBindGroup WGPU_OBJECT_ATTRIBUTE;
typedef struct WGPUBindGroupLayoutImpl* WGPUBindGroupLayout WGPU_OBJECT_ATTRIBUTE;
typedef struct WGPUBufferImpl* WGPUBuffer WGPU_OBJECT_ATTRIBUTE;
typedef struct WGPUCommandBufferImpl* WGPUCommandBuffer WGPU_OBJECT_ATTRIBUTE;
typedef struct WGPUCommandEncoderImpl* WGPUCommandEncoder WGPU_OBJECT_ATTRIBUTE;
typedef struct WGPUComputePipelineImpl* WGPUComputePipeline WGPU_OBJECT_ATTRIBUTE;
typedef struct WGPUDeviceImpl* WGPUDevice WGPU_OBJECT_ATTRIBUTE;
typedef struct WGPUExternalTextureImpl* WGPUExternalTexture WGPU_OBJECT_ATTRIBUTE;
@@ -143,13 +144,13 @@ typedef struct WGPUSurfaceImpl* WGPUSurface WGPU_OBJECT_ATTRIBUTE;
typedef struct WGPUTexelBufferViewImpl* WGPUTexelBufferView WGPU_OBJECT_ATTRIBUTE;
typedef struct WGPUTextureImpl* WGPUTexture WGPU_OBJECT_ATTRIBUTE;
typedef struct WGPUTextureViewImpl* WGPUTextureView WGPU_OBJECT_ATTRIBUTE;
typedef struct WGPUCommandEncoderImpl* WGPUCommandEncoder WGPU_OBJECT_ATTRIBUTE;
typedef struct WGPUComputePassEncoderImpl* WGPUComputePassEncoder WGPU_OBJECT_ATTRIBUTE;
typedef struct WGPURenderBundleEncoderImpl* WGPURenderBundleEncoder WGPU_OBJECT_ATTRIBUTE;
typedef struct WGPURenderPassEncoderImpl* WGPURenderPassEncoder WGPU_OBJECT_ATTRIBUTE;
// Structure forward declarations
struct WGPUAdapterPropertiesD3D;
struct WGPUAdapterPropertiesExplicitComputeSubgroupSizeConfigs;
struct WGPUAdapterPropertiesVk;
struct WGPUAdapterPropertiesWGPU;
struct WGPUBindingResource;
@@ -174,7 +175,7 @@ struct WGPUDawnFakeBufferOOMForTesting;
struct WGPUDawnFakeDeviceInitializeErrorForTesting;
struct WGPUDawnHostMappedPointerLimits;
struct WGPUDawnInjectedInvalidSType;
struct WGPUDawnRenderPassColorAttachmentRenderToSingleSampled;
struct WGPUDawnRenderPassSampleCount;
struct WGPUDawnShaderModuleSPIRVOptionsDescriptor;
struct WGPUDawnTexelCopyBufferRowAlignmentLimits;
struct WGPUDawnTextureInternalUsageDescriptor;
@@ -208,7 +209,6 @@ struct WGPURenderPassMaxDrawCount;
struct WGPURequestAdapterWebGPUBackendOptions;
struct WGPURequestAdapterWebXROptions;
struct WGPUResourceTableDescriptor;
struct WGPUResourceTableLimits;
struct WGPUSamplerBindingLayout;
struct WGPUShaderModuleCompilationOptions;
struct WGPUShaderSourceSPIRV;
@@ -612,6 +612,8 @@ typedef enum WGPUFeatureName {
WGPUFeatureName_SharedTextureMemoryD3D12Resource = 0x00050039,
WGPUFeatureName_ChromiumExperimentalSamplingResourceTable = 0x0005003A,
WGPUFeatureName_ChromiumExperimentalSubgroupSizeControl = 0x0005003B,
WGPUFeatureName_AtomicVec2uMinMax = 0x0005003C,
WGPUFeatureName_Unorm16FormatsForExternalTexture = 0x0005003D,
WGPUFeatureName_Force32 = 0x7FFFFFFF
} WGPUFeatureName WGPU_ENUM_ATTRIBUTE;
@@ -833,7 +835,7 @@ typedef enum WGPUSType {
WGPUSType_RequestAdapterOptionsLUID = 0x0005000C,
WGPUSType_RequestAdapterOptionsGetGLProc = 0x0005000D,
WGPUSType_RequestAdapterOptionsD3D11Device = 0x0005000E,
WGPUSType_DawnRenderPassColorAttachmentRenderToSingleSampled = 0x0005000F,
WGPUSType_DawnRenderPassSampleCount = 0x0005000F,
WGPUSType_RenderPassPixelLocalStorage = 0x00050010,
WGPUSType_PipelineLayoutPixelLocalStorage = 0x00050011,
WGPUSType_BufferHostMappedPointer = 0x00050012,
@@ -897,8 +899,8 @@ typedef enum WGPUSType {
WGPUSType_SharedBufferMemoryD3D12SharedMemoryFileMappingHandleDescriptor = 0x0005004D,
WGPUSType_SharedTextureMemoryD3D12ResourceDescriptor = 0x0005004E,
WGPUSType_RequestAdapterOptionsAngleVirtualizationGroup = 0x0005004F,
WGPUSType_ResourceTableLimits = 0x00050050,
WGPUSType_PipelineLayoutResourceTable = 0x00050051,
WGPUSType_PipelineLayoutResourceTable = 0x00050050,
WGPUSType_AdapterPropertiesExplicitComputeSubgroupSizeConfigs = 0x00050051,
WGPUSType_Force32 = 0x7FFFFFFF
} WGPUSType WGPU_ENUM_ATTRIBUTE;
@@ -1051,14 +1053,15 @@ typedef enum WGPUTextureFormat {
WGPUTextureFormat_ASTC12x10UnormSrgb = 0x00000063,
WGPUTextureFormat_ASTC12x12Unorm = 0x00000064,
WGPUTextureFormat_ASTC12x12UnormSrgb = 0x00000065,
WGPUTextureFormat_R8BG8Biplanar420Unorm = 0x00050006,
WGPUTextureFormat_R10X6BG10X6Biplanar420Unorm = 0x00050007,
WGPUTextureFormat_R8BG8A8Triplanar420Unorm = 0x00050008,
WGPUTextureFormat_R8BG8Biplanar422Unorm = 0x00050009,
WGPUTextureFormat_R8BG8Biplanar444Unorm = 0x0005000A,
WGPUTextureFormat_R10X6BG10X6Biplanar422Unorm = 0x0005000B,
WGPUTextureFormat_R10X6BG10X6Biplanar444Unorm = 0x0005000C,
WGPUTextureFormat_External = 0x0005000D,
WGPUTextureFormat_R8BG8Biplanar420Unorm = 0x00050000,
WGPUTextureFormat_R10X6BG10X6Biplanar420Unorm = 0x00050001,
WGPUTextureFormat_R8BG8A8Triplanar420Unorm = 0x00050002,
WGPUTextureFormat_R8BG8Biplanar422Unorm = 0x00050003,
WGPUTextureFormat_R8BG8Biplanar444Unorm = 0x00050004,
WGPUTextureFormat_R10X6BG10X6Biplanar422Unorm = 0x00050005,
WGPUTextureFormat_R10X6BG10X6Biplanar444Unorm = 0x00050006,
WGPUTextureFormat_External = 0x00050007,
WGPUTextureFormat_OpaqueYCbCrAndroid = 0x00050007,
WGPUTextureFormat_Force32 = 0x7FFFFFFF
} WGPUTextureFormat WGPU_ENUM_ATTRIBUTE;
@@ -1156,19 +1159,22 @@ typedef enum WGPUWGSLLanguageFeatureName {
WGPUWGSLLanguageFeatureName_PointerCompositeAccess = 0x00000004,
WGPUWGSLLanguageFeatureName_UniformBufferStandardLayout = 0x00000005,
WGPUWGSLLanguageFeatureName_SubgroupId = 0x00000006,
WGPUWGSLLanguageFeatureName_SizedBindingArray = 0x00050007,
WGPUWGSLLanguageFeatureName_TexelBuffers = 0x00050008,
WGPUWGSLLanguageFeatureName_ChromiumPrint = 0x00050009,
WGPUWGSLLanguageFeatureName_FragmentDepth = 0x0005000A,
WGPUWGSLLanguageFeatureName_ImmediateAddressSpace = 0x0005000B,
WGPUWGSLLanguageFeatureName_SubgroupUniformity = 0x0005000C,
WGPUWGSLLanguageFeatureName_TextureAndSamplerLet = 0x0005000D,
WGPUWGSLLanguageFeatureName_BufferView = 0x0005000E,
WGPUWGSLLanguageFeatureName_TextureAndSamplerLet = 0x00000007,
WGPUWGSLLanguageFeatureName_ChromiumTestingUnimplemented = 0x00050000,
WGPUWGSLLanguageFeatureName_ChromiumTestingUnsafeExperimental = 0x00050001,
WGPUWGSLLanguageFeatureName_ChromiumTestingExperimental = 0x00050002,
WGPUWGSLLanguageFeatureName_ChromiumTestingShippedWithKillswitch = 0x00050003,
WGPUWGSLLanguageFeatureName_ChromiumTestingShipped = 0x00050004,
WGPUWGSLLanguageFeatureName_SizedBindingArray = 0x00050005,
WGPUWGSLLanguageFeatureName_TexelBuffers = 0x00050006,
WGPUWGSLLanguageFeatureName_ChromiumPrint = 0x00050007,
WGPUWGSLLanguageFeatureName_FragmentDepth = 0x00050008,
WGPUWGSLLanguageFeatureName_ImmediateAddressSpace = 0x00050009,
WGPUWGSLLanguageFeatureName_SubgroupUniformity = 0x0005000A,
WGPUWGSLLanguageFeatureName_BufferView = 0x0005000B,
WGPUWGSLLanguageFeatureName_FilteringParameters = 0x0005000C,
WGPUWGSLLanguageFeatureName_SwizzleAssignment = 0x0005000D,
WGPUWGSLLanguageFeatureName_LinearIndexing = 0x0005000E,
WGPUWGSLLanguageFeatureName_Force32 = 0x7FFFFFFF
} WGPUWGSLLanguageFeatureName WGPU_ENUM_ATTRIBUTE;
@@ -1442,6 +1448,24 @@ typedef struct WGPUAdapterPropertiesD3D {
/*.shaderModel=*/0 _wgpu_COMMA \
})
// Can be chained in WGPUAdapterInfo
typedef struct WGPUAdapterPropertiesExplicitComputeSubgroupSizeConfigs {
WGPUChainedStruct chain;
uint32_t minExplicitComputeSubgroupSize;
uint32_t maxExplicitComputeSubgroupSize;
uint32_t maxComputeWorkgroupSubgroups;
} WGPUAdapterPropertiesExplicitComputeSubgroupSizeConfigs WGPU_STRUCTURE_ATTRIBUTE;
#define WGPU_ADAPTER_PROPERTIES_EXPLICIT_COMPUTE_SUBGROUP_SIZE_CONFIGS_INIT _wgpu_MAKE_INIT_STRUCT(WGPUAdapterPropertiesExplicitComputeSubgroupSizeConfigs, { \
/*.chain=*/_wgpu_MAKE_INIT_STRUCT(WGPUChainedStruct, { \
/*.next=*/NULL _wgpu_COMMA \
/*.sType=*/WGPUSType_AdapterPropertiesExplicitComputeSubgroupSizeConfigs _wgpu_COMMA \
}) _wgpu_COMMA \
/*.minExplicitComputeSubgroupSize=*/0 _wgpu_COMMA \
/*.maxExplicitComputeSubgroupSize=*/0 _wgpu_COMMA \
/*.maxComputeWorkgroupSubgroups=*/0 _wgpu_COMMA \
})
// Can be chained in WGPUAdapterInfo
typedef struct WGPUAdapterPropertiesVk {
WGPUChainedStruct chain;
@@ -1801,18 +1825,18 @@ typedef struct WGPUDawnInjectedInvalidSType {
/*.invalidSType=*/_wgpu_ENUM_ZERO_INIT(WGPUSType) _wgpu_COMMA \
})
// Can be chained in WGPURenderPassColorAttachment
typedef struct WGPUDawnRenderPassColorAttachmentRenderToSingleSampled {
// Can be chained in WGPURenderPassDescriptor
typedef struct WGPUDawnRenderPassSampleCount {
WGPUChainedStruct chain;
uint32_t implicitSampleCount;
} WGPUDawnRenderPassColorAttachmentRenderToSingleSampled WGPU_STRUCTURE_ATTRIBUTE;
uint32_t sampleCount;
} WGPUDawnRenderPassSampleCount WGPU_STRUCTURE_ATTRIBUTE;
#define WGPU_DAWN_RENDER_PASS_COLOR_ATTACHMENT_RENDER_TO_SINGLE_SAMPLED_INIT _wgpu_MAKE_INIT_STRUCT(WGPUDawnRenderPassColorAttachmentRenderToSingleSampled, { \
#define WGPU_DAWN_RENDER_PASS_SAMPLE_COUNT_INIT _wgpu_MAKE_INIT_STRUCT(WGPUDawnRenderPassSampleCount, { \
/*.chain=*/_wgpu_MAKE_INIT_STRUCT(WGPUChainedStruct, { \
/*.next=*/NULL _wgpu_COMMA \
/*.sType=*/WGPUSType_DawnRenderPassColorAttachmentRenderToSingleSampled _wgpu_COMMA \
/*.sType=*/WGPUSType_DawnRenderPassSampleCount _wgpu_COMMA \
}) _wgpu_COMMA \
/*.implicitSampleCount=*/1 _wgpu_COMMA \
/*.sampleCount=*/1 _wgpu_COMMA \
})
// Can be chained in WGPUShaderModuleDescriptor
@@ -2283,20 +2307,6 @@ typedef struct WGPUResourceTableDescriptor {
/*.size=*/0 _wgpu_COMMA \
})
// Can be chained in WGPULimits
typedef struct WGPUResourceTableLimits {
WGPUChainedStruct chain;
uint32_t maxResourceTableSize;
} WGPUResourceTableLimits WGPU_STRUCTURE_ATTRIBUTE;
#define WGPU_RESOURCE_TABLE_LIMITS_INIT _wgpu_MAKE_INIT_STRUCT(WGPUResourceTableLimits, { \
/*.chain=*/_wgpu_MAKE_INIT_STRUCT(WGPUChainedStruct, { \
/*.next=*/NULL _wgpu_COMMA \
/*.sType=*/WGPUSType_ResourceTableLimits _wgpu_COMMA \
}) _wgpu_COMMA \
/*.maxResourceTableSize=*/WGPU_LIMIT_U32_UNDEFINED _wgpu_COMMA \
})
typedef struct WGPUSamplerBindingLayout {
WGPUChainedStruct * nextInChain;
WGPUSamplerBindingType type;
@@ -4249,7 +4259,6 @@ typedef void (*WGPUProcCommandEncoderPopDebugGroup)(WGPUCommandEncoder commandEn
typedef void (*WGPUProcCommandEncoderPushDebugGroup)(WGPUCommandEncoder commandEncoder, WGPUStringView groupLabel) WGPU_FUNCTION_ATTRIBUTE;
typedef void (*WGPUProcCommandEncoderResolveQuerySet)(WGPUCommandEncoder commandEncoder, WGPUQuerySet querySet, uint32_t firstQuery, uint32_t queryCount, WGPUBuffer destination, uint64_t destinationOffset) WGPU_FUNCTION_ATTRIBUTE;
typedef void (*WGPUProcCommandEncoderSetLabel)(WGPUCommandEncoder commandEncoder, WGPUStringView label) WGPU_FUNCTION_ATTRIBUTE;
typedef void (*WGPUProcCommandEncoderSetResourceTable)(WGPUCommandEncoder commandEncoder, WGPU_NULLABLE WGPUResourceTable table) WGPU_FUNCTION_ATTRIBUTE;
typedef void (*WGPUProcCommandEncoderWriteBuffer)(WGPUCommandEncoder commandEncoder, WGPUBuffer buffer, uint64_t bufferOffset, uint8_t const * data, uint64_t size) WGPU_FUNCTION_ATTRIBUTE;
typedef void (*WGPUProcCommandEncoderWriteTimestamp)(WGPUCommandEncoder commandEncoder, WGPUQuerySet querySet, uint32_t queryIndex) WGPU_FUNCTION_ATTRIBUTE;
typedef void (*WGPUProcCommandEncoderAddRef)(WGPUCommandEncoder commandEncoder) WGPU_FUNCTION_ATTRIBUTE;
@@ -4266,6 +4275,7 @@ typedef void (*WGPUProcComputePassEncoderSetBindGroup)(WGPUComputePassEncoder co
typedef void (*WGPUProcComputePassEncoderSetImmediates)(WGPUComputePassEncoder computePassEncoder, uint32_t offset, void const * data, size_t size) WGPU_FUNCTION_ATTRIBUTE;
typedef void (*WGPUProcComputePassEncoderSetLabel)(WGPUComputePassEncoder computePassEncoder, WGPUStringView label) WGPU_FUNCTION_ATTRIBUTE;
typedef void (*WGPUProcComputePassEncoderSetPipeline)(WGPUComputePassEncoder computePassEncoder, WGPUComputePipeline pipeline) WGPU_FUNCTION_ATTRIBUTE;
typedef void (*WGPUProcComputePassEncoderSetResourceTable)(WGPUComputePassEncoder computePassEncoder, WGPU_NULLABLE WGPUResourceTable table) WGPU_FUNCTION_ATTRIBUTE;
typedef void (*WGPUProcComputePassEncoderWriteTimestamp)(WGPUComputePassEncoder computePassEncoder, WGPUQuerySet querySet, uint32_t queryIndex) WGPU_FUNCTION_ATTRIBUTE;
typedef void (*WGPUProcComputePassEncoderAddRef)(WGPUComputePassEncoder computePassEncoder) WGPU_FUNCTION_ATTRIBUTE;
typedef void (*WGPUProcComputePassEncoderRelease)(WGPUComputePassEncoder computePassEncoder) WGPU_FUNCTION_ATTRIBUTE;
@@ -4384,6 +4394,7 @@ typedef void (*WGPUProcRenderBundleEncoderSetImmediates)(WGPURenderBundleEncoder
typedef void (*WGPUProcRenderBundleEncoderSetIndexBuffer)(WGPURenderBundleEncoder renderBundleEncoder, WGPUBuffer buffer, WGPUIndexFormat format, uint64_t offset, uint64_t size) WGPU_FUNCTION_ATTRIBUTE;
typedef void (*WGPUProcRenderBundleEncoderSetLabel)(WGPURenderBundleEncoder renderBundleEncoder, WGPUStringView label) WGPU_FUNCTION_ATTRIBUTE;
typedef void (*WGPUProcRenderBundleEncoderSetPipeline)(WGPURenderBundleEncoder renderBundleEncoder, WGPURenderPipeline pipeline) WGPU_FUNCTION_ATTRIBUTE;
typedef void (*WGPUProcRenderBundleEncoderSetResourceTable)(WGPURenderBundleEncoder renderBundleEncoder, WGPU_NULLABLE WGPUResourceTable table) WGPU_FUNCTION_ATTRIBUTE;
typedef void (*WGPUProcRenderBundleEncoderSetVertexBuffer)(WGPURenderBundleEncoder renderBundleEncoder, uint32_t slot, WGPU_NULLABLE WGPUBuffer buffer, uint64_t offset, uint64_t size) WGPU_FUNCTION_ATTRIBUTE;
typedef void (*WGPUProcRenderBundleEncoderAddRef)(WGPURenderBundleEncoder renderBundleEncoder) WGPU_FUNCTION_ATTRIBUTE;
typedef void (*WGPUProcRenderBundleEncoderRelease)(WGPURenderBundleEncoder renderBundleEncoder) WGPU_FUNCTION_ATTRIBUTE;
@@ -4409,6 +4420,7 @@ typedef void (*WGPUProcRenderPassEncoderSetImmediates)(WGPURenderPassEncoder ren
typedef void (*WGPUProcRenderPassEncoderSetIndexBuffer)(WGPURenderPassEncoder renderPassEncoder, WGPUBuffer buffer, WGPUIndexFormat format, uint64_t offset, uint64_t size) WGPU_FUNCTION_ATTRIBUTE;
typedef void (*WGPUProcRenderPassEncoderSetLabel)(WGPURenderPassEncoder renderPassEncoder, WGPUStringView label) WGPU_FUNCTION_ATTRIBUTE;
typedef void (*WGPUProcRenderPassEncoderSetPipeline)(WGPURenderPassEncoder renderPassEncoder, WGPURenderPipeline pipeline) WGPU_FUNCTION_ATTRIBUTE;
typedef void (*WGPUProcRenderPassEncoderSetResourceTable)(WGPURenderPassEncoder renderPassEncoder, WGPU_NULLABLE WGPUResourceTable table) WGPU_FUNCTION_ATTRIBUTE;
typedef void (*WGPUProcRenderPassEncoderSetScissorRect)(WGPURenderPassEncoder renderPassEncoder, uint32_t x, uint32_t y, uint32_t width, uint32_t height) WGPU_FUNCTION_ATTRIBUTE;
typedef void (*WGPUProcRenderPassEncoderSetStencilReference)(WGPURenderPassEncoder renderPassEncoder, uint32_t reference) WGPU_FUNCTION_ATTRIBUTE;
typedef void (*WGPUProcRenderPassEncoderSetVertexBuffer)(WGPURenderPassEncoder renderPassEncoder, uint32_t slot, WGPU_NULLABLE WGPUBuffer buffer, uint64_t offset, uint64_t size) WGPU_FUNCTION_ATTRIBUTE;
@@ -4602,7 +4614,6 @@ WGPU_EXPORT void wgpuCommandEncoderPopDebugGroup(WGPUCommandEncoder commandEncod
WGPU_EXPORT void wgpuCommandEncoderPushDebugGroup(WGPUCommandEncoder commandEncoder, WGPUStringView groupLabel) WGPU_FUNCTION_ATTRIBUTE;
WGPU_EXPORT void wgpuCommandEncoderResolveQuerySet(WGPUCommandEncoder commandEncoder, WGPUQuerySet querySet, uint32_t firstQuery, uint32_t queryCount, WGPUBuffer destination, uint64_t destinationOffset) WGPU_FUNCTION_ATTRIBUTE;
WGPU_EXPORT void wgpuCommandEncoderSetLabel(WGPUCommandEncoder commandEncoder, WGPUStringView label) WGPU_FUNCTION_ATTRIBUTE;
WGPU_EXPORT void wgpuCommandEncoderSetResourceTable(WGPUCommandEncoder commandEncoder, WGPU_NULLABLE WGPUResourceTable table) WGPU_FUNCTION_ATTRIBUTE;
WGPU_EXPORT void wgpuCommandEncoderWriteBuffer(WGPUCommandEncoder commandEncoder, WGPUBuffer buffer, uint64_t bufferOffset, uint8_t const * data, uint64_t size) WGPU_FUNCTION_ATTRIBUTE;
WGPU_EXPORT void wgpuCommandEncoderWriteTimestamp(WGPUCommandEncoder commandEncoder, WGPUQuerySet querySet, uint32_t queryIndex) WGPU_FUNCTION_ATTRIBUTE;
WGPU_EXPORT void wgpuCommandEncoderAddRef(WGPUCommandEncoder commandEncoder) WGPU_FUNCTION_ATTRIBUTE;
@@ -4619,6 +4630,7 @@ WGPU_EXPORT void wgpuComputePassEncoderSetBindGroup(WGPUComputePassEncoder compu
WGPU_EXPORT void wgpuComputePassEncoderSetImmediates(WGPUComputePassEncoder computePassEncoder, uint32_t offset, void const * data, size_t size) WGPU_FUNCTION_ATTRIBUTE;
WGPU_EXPORT void wgpuComputePassEncoderSetLabel(WGPUComputePassEncoder computePassEncoder, WGPUStringView label) WGPU_FUNCTION_ATTRIBUTE;
WGPU_EXPORT void wgpuComputePassEncoderSetPipeline(WGPUComputePassEncoder computePassEncoder, WGPUComputePipeline pipeline) WGPU_FUNCTION_ATTRIBUTE;
WGPU_EXPORT void wgpuComputePassEncoderSetResourceTable(WGPUComputePassEncoder computePassEncoder, WGPU_NULLABLE WGPUResourceTable table) WGPU_FUNCTION_ATTRIBUTE;
WGPU_EXPORT void wgpuComputePassEncoderWriteTimestamp(WGPUComputePassEncoder computePassEncoder, WGPUQuerySet querySet, uint32_t queryIndex) WGPU_FUNCTION_ATTRIBUTE;
WGPU_EXPORT void wgpuComputePassEncoderAddRef(WGPUComputePassEncoder computePassEncoder) WGPU_FUNCTION_ATTRIBUTE;
WGPU_EXPORT void wgpuComputePassEncoderRelease(WGPUComputePassEncoder computePassEncoder) WGPU_FUNCTION_ATTRIBUTE;
@@ -4737,6 +4749,7 @@ WGPU_EXPORT void wgpuRenderBundleEncoderSetImmediates(WGPURenderBundleEncoder re
WGPU_EXPORT void wgpuRenderBundleEncoderSetIndexBuffer(WGPURenderBundleEncoder renderBundleEncoder, WGPUBuffer buffer, WGPUIndexFormat format, uint64_t offset, uint64_t size) WGPU_FUNCTION_ATTRIBUTE;
WGPU_EXPORT void wgpuRenderBundleEncoderSetLabel(WGPURenderBundleEncoder renderBundleEncoder, WGPUStringView label) WGPU_FUNCTION_ATTRIBUTE;
WGPU_EXPORT void wgpuRenderBundleEncoderSetPipeline(WGPURenderBundleEncoder renderBundleEncoder, WGPURenderPipeline pipeline) WGPU_FUNCTION_ATTRIBUTE;
WGPU_EXPORT void wgpuRenderBundleEncoderSetResourceTable(WGPURenderBundleEncoder renderBundleEncoder, WGPU_NULLABLE WGPUResourceTable table) WGPU_FUNCTION_ATTRIBUTE;
WGPU_EXPORT void wgpuRenderBundleEncoderSetVertexBuffer(WGPURenderBundleEncoder renderBundleEncoder, uint32_t slot, WGPU_NULLABLE WGPUBuffer buffer, uint64_t offset, uint64_t size) WGPU_FUNCTION_ATTRIBUTE;
WGPU_EXPORT void wgpuRenderBundleEncoderAddRef(WGPURenderBundleEncoder renderBundleEncoder) WGPU_FUNCTION_ATTRIBUTE;
WGPU_EXPORT void wgpuRenderBundleEncoderRelease(WGPURenderBundleEncoder renderBundleEncoder) WGPU_FUNCTION_ATTRIBUTE;
@@ -4762,6 +4775,7 @@ WGPU_EXPORT void wgpuRenderPassEncoderSetImmediates(WGPURenderPassEncoder render
WGPU_EXPORT void wgpuRenderPassEncoderSetIndexBuffer(WGPURenderPassEncoder renderPassEncoder, WGPUBuffer buffer, WGPUIndexFormat format, uint64_t offset, uint64_t size) WGPU_FUNCTION_ATTRIBUTE;
WGPU_EXPORT void wgpuRenderPassEncoderSetLabel(WGPURenderPassEncoder renderPassEncoder, WGPUStringView label) WGPU_FUNCTION_ATTRIBUTE;
WGPU_EXPORT void wgpuRenderPassEncoderSetPipeline(WGPURenderPassEncoder renderPassEncoder, WGPURenderPipeline pipeline) WGPU_FUNCTION_ATTRIBUTE;
WGPU_EXPORT void wgpuRenderPassEncoderSetResourceTable(WGPURenderPassEncoder renderPassEncoder, WGPU_NULLABLE WGPUResourceTable table) WGPU_FUNCTION_ATTRIBUTE;
WGPU_EXPORT void wgpuRenderPassEncoderSetScissorRect(WGPURenderPassEncoder renderPassEncoder, uint32_t x, uint32_t y, uint32_t width, uint32_t height) WGPU_FUNCTION_ATTRIBUTE;
WGPU_EXPORT void wgpuRenderPassEncoderSetStencilReference(WGPURenderPassEncoder renderPassEncoder, uint32_t reference) WGPU_FUNCTION_ATTRIBUTE;
WGPU_EXPORT void wgpuRenderPassEncoderSetVertexBuffer(WGPURenderPassEncoder renderPassEncoder, uint32_t slot, WGPU_NULLABLE WGPUBuffer buffer, uint64_t offset, uint64_t size) WGPU_FUNCTION_ATTRIBUTE;
3rdparty/dawn/src/tint/api/common/override_id.h
Vendored
+1
View File
@@ -29,6 +29,7 @@
#define SRC_TINT_API_COMMON_OVERRIDE_ID_H_
#include <stdint.h>
#include <functional>
#include "src/tint/utils/math/hash.h"
3rdparty/dawn/src/tint/api/common/resource_table_config.h
Vendored
+2
View File
@@ -58,6 +58,7 @@ struct ResourceTableConfig {
BindingPoint resource_table_binding;
// The binding point for the supporting storage buffer. This is a post-remapping binding point.
// TODO(crbug.com/435317394): Rename to metadata_buffer_binding
BindingPoint storage_buffer_binding;
// The ordering of default bindings which are placed after the user bindings. These will be used
@@ -91,6 +92,7 @@ struct ResourceTableConfig {
resource_table_binding,
storage_buffer_binding,
default_binding_type_order);
TINT_REFLECT_EQUALS(ResourceTableConfig);
};
} // namespace tint
3rdparty/dawn/src/tint/api/common/resource_type.h
Vendored
+16 -6
View File
@@ -35,22 +35,28 @@ namespace tint {
enum class ResourceType : uint32_t {
kEmpty,
kTexture1d_f32,
kTexture1d_f32_filterable,
kTexture1d_f32_unfilterable,
kTexture1d_i32,
kTexture1d_u32,
kTexture2d_f32,
kTexture2d_f32_filterable,
kTexture2d_f32_unfilterable,
kTexture2d_i32,
kTexture2d_u32,
kTexture2dArray_f32,
kTexture2dArray_f32_filterable,
kTexture2dArray_f32_unfilterable,
kTexture2dArray_i32,
kTexture2dArray_u32,
kTexture3d_f32,
kTexture3d_f32_filterable,
kTexture3d_f32_unfilterable,
kTexture3d_i32,
kTexture3d_u32,
kTextureCube_f32,
kTextureCube_f32_filterable,
kTextureCube_f32_unfilterable,
kTextureCube_i32,
kTextureCube_u32,
kTextureCubeArray_f32,
kTextureCubeArray_f32_filterable,
kTextureCubeArray_f32_unfilterable,
kTextureCubeArray_i32,
kTextureCubeArray_u32,
@@ -63,6 +69,10 @@ enum class ResourceType : uint32_t {
kTextureDepthCube,
kTextureDepthCubeArray,
kTextureDepthMultisampled2d,
kSampler_filtering,
kSampler_non_filtering,
kSampler_comparison,
};
TINT_REFLECT_ENUM_RANGE(tint::ResourceType, kEmpty, kTextureDepthMultisampled2d);
3rdparty/dawn/src/tint/api/common/subgroup_matrix.h
Vendored
-1
View File
@@ -31,7 +31,6 @@
#include <unordered_set>
#include "src/tint/utils/math/hash.h"
#include "src/tint/utils/reflection.h"
namespace tint {
3rdparty/dawn/src/tint/lang/core/ir/transform/decompose_uniform_access.h → 3rdparty/dawn/src/tint/api/common/workgroup_info.h
Vendored
+23 -17
View File
@@ -1,4 +1,4 @@
// Copyright 2024 The Dawn & Tint Authors
// Copyright 2026 The Dawn & Tint Authors
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are met:
@@ -25,25 +25,31 @@
// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
#ifndef SRC_TINT_LANG_CORE_IR_TRANSFORM_DECOMPOSE_UNIFORM_ACCESS_H_
#define SRC_TINT_LANG_CORE_IR_TRANSFORM_DECOMPOSE_UNIFORM_ACCESS_H_
#ifndef SRC_TINT_API_COMMON_WORKGROUP_INFO_H_
#define SRC_TINT_API_COMMON_WORKGROUP_INFO_H_
#include "src/tint/utils/result.h"
#include <cstddef>
#include <cstdint>
#include <optional>
// Forward declarations.
namespace tint::core::ir {
class Module;
} // namespace tint::core::ir
namespace tint {
namespace tint::core::ir::transform {
/// Workgroup size information, which is part of the output of shader generation.
struct WorkgroupInfo {
/// The x-component
uint32_t x = 0;
/// The y-component
uint32_t y = 0;
/// The z-component
uint32_t z = 0;
/// DecomposeUniformAccess is a transform used to replace uniform buffer accesses with a
/// the correct load of the HLSL primitive objects
///
/// @param module the module to transform
/// @returns success or failure
Result<SuccessType> DecomposeUniformAccess(core::ir::Module& module);
/// The needed workgroup storage size
size_t storage_size = 0;
} // namespace tint::core::ir::transform
/// The `@subgroup_size` attribute
std::optional<uint32_t> subgroup_size;
};
#endif // SRC_TINT_LANG_CORE_IR_TRANSFORM_DECOMPOSE_UNIFORM_ACCESS_H_
} // namespace tint
#endif // SRC_TINT_API_COMMON_WORKGROUP_INFO_H_
3rdparty/dawn/src/tint/lang/core/constant/invalid.h
Vendored
+1
View File
@@ -29,6 +29,7 @@
#define SRC_TINT_LANG_CORE_CONSTANT_INVALID_H_
#include <variant>
#include "src/tint/lang/core/constant/value.h"
#include "src/tint/lang/core/number.h"
#include "src/tint/lang/core/type/invalid.h"
3rdparty/dawn/src/tint/lang/core/constant/manager.cc
Vendored
+6
View File
@@ -42,6 +42,7 @@
#include "src/tint/lang/core/type/i8.h"
#include "src/tint/lang/core/type/manager.h"
#include "src/tint/lang/core/type/matrix.h"
#include "src/tint/lang/core/type/u16.h"
#include "src/tint/lang/core/type/u32.h"
#include "src/tint/lang/core/type/u64.h"
#include "src/tint/lang/core/type/u8.h"
@@ -115,6 +116,10 @@ const Scalar<u8>* Manager::Get(u8 value) {
return Get<Scalar<u8>>(types.u8(), value);
}
const Scalar<u16>* Manager::Get(u16 value) {
return Get<Scalar<u16>>(types.u16(), value);
}
const Scalar<f32>* Manager::Get(f32 value) {
return Get<Scalar<f32>>(types.f32(), value);
}
@@ -182,6 +187,7 @@ const Value* Manager::Zero(const core::type::Type* type) {
[&](const core::type::U32*) { return Get(u32(0)); }, //
[&](const core::type::I8*) { return Get(i8(0)); }, //
[&](const core::type::U8*) { return Get(u8(0)); }, //
[&](const core::type::U16*) { return Get(u16(0)); }, //
[&](const core::type::U64*) { return Get(u64(0)); }, //
[&](const core::type::F32*) { return Get(f32(0)); }, //
[&](const core::type::F16*) { return Get(f16(0)); }, //
3rdparty/dawn/src/tint/lang/core/constant/manager.h
Vendored
+4
View File
@@ -118,6 +118,10 @@ class Manager final {
/// @return a Scalar holding the u8 value @p value
const Scalar<u8>* Get(u8 value);
/// @param value the constant value
/// @return a Scalar holding the u16 value @p value
const Scalar<u16>* Get(u16 value);
/// @param value the constant value
/// @return a Scalar holding the f32 value @p value
const Scalar<f32>* Get(f32 value);
3rdparty/dawn/src/tint/lang/core/constant/scalar.cc
Vendored
+1
View File
@@ -35,6 +35,7 @@ TINT_INSTANTIATE_TYPEINFO(tint::core::constant::Scalar<tint::core::i8>);
TINT_INSTANTIATE_TYPEINFO(tint::core::constant::Scalar<tint::core::u32>);
TINT_INSTANTIATE_TYPEINFO(tint::core::constant::Scalar<tint::core::u64>);
TINT_INSTANTIATE_TYPEINFO(tint::core::constant::Scalar<tint::core::u8>);
TINT_INSTANTIATE_TYPEINFO(tint::core::constant::Scalar<tint::core::u16>);
TINT_INSTANTIATE_TYPEINFO(tint::core::constant::Scalar<tint::core::f16>);
TINT_INSTANTIATE_TYPEINFO(tint::core::constant::Scalar<tint::core::f32>);
TINT_INSTANTIATE_TYPEINFO(tint::core::constant::Scalar<bool>);
3rdparty/dawn/src/tint/lang/core/enums.cc
Vendored
+82 -5
View File
@@ -192,6 +192,54 @@ std::string_view ToString(InterpolationType value) {
return "<unknown>";
}
/// ParseSamplerFiltering parses a SamplerFiltering from a string.
/// @param str the string to parse
/// @returns the parsed enum, or SamplerFiltering::kUndefined if the string could not be parsed.
SamplerFiltering ParseSamplerFiltering(std::string_view str) {
if (str == "filtering") {
return SamplerFiltering::kFiltering;
}
if (str == "non_filtering") {
return SamplerFiltering::kNonFiltering;
}
return SamplerFiltering::kUndefined;
}
std::string_view ToString(SamplerFiltering value) {
switch (value) {
case SamplerFiltering::kUndefined:
return "undefined";
case SamplerFiltering::kFiltering:
return "filtering";
case SamplerFiltering::kNonFiltering:
return "non_filtering";
}
return "<unknown>";
}
/// ParseTextureFilterable parses a TextureFilterable from a string.
/// @param str the string to parse
/// @returns the parsed enum, or TextureFilterable::kUndefined if the string could not be parsed.
TextureFilterable ParseTextureFilterable(std::string_view str) {
if (str == "filterable") {
return TextureFilterable::kFilterable;
}
if (str == "unfilterable") {
return TextureFilterable::kUnfilterable;
}
return TextureFilterable::kUndefined;
}
std::string_view ToString(TextureFilterable value) {
switch (value) {
case TextureFilterable::kUndefined:
return "undefined";
case TextureFilterable::kFilterable:
return "filterable";
case TextureFilterable::kUnfilterable:
return "unfilterable";
}
return "<unknown>";
}
/// ParseSubgroupMatrixKind parses a SubgroupMatrixKind from a string.
/// @param str the string to parse
/// @returns the parsed enum, or SubgroupMatrixKind::kUndefined if the string could not be parsed.
@@ -700,6 +748,9 @@ BuiltinType ParseBuiltinType(std::string_view str) {
if (str == "texture_storage_3d") {
return BuiltinType::kTextureStorage3D;
}
if (str == "u16") {
return BuiltinType::kU16;
}
if (str == "u32") {
return BuiltinType::kU32;
}
@@ -931,6 +982,8 @@ std::string_view ToString(BuiltinType value) {
return "texture_storage_2d_array";
case BuiltinType::kTextureStorage3D:
return "texture_storage_3d";
case BuiltinType::kU16:
return "u16";
case BuiltinType::kU32:
return "u32";
case BuiltinType::kU8:
@@ -988,6 +1041,9 @@ BuiltinValue ParseBuiltinValue(std::string_view str) {
if (str == "global_invocation_id") {
return BuiltinValue::kGlobalInvocationId;
}
if (str == "global_invocation_index") {
return BuiltinValue::kGlobalInvocationIndex;
}
if (str == "instance_index") {
return BuiltinValue::kInstanceIndex;
}
@@ -1030,6 +1086,9 @@ BuiltinValue ParseBuiltinValue(std::string_view str) {
if (str == "workgroup_id") {
return BuiltinValue::kWorkgroupId;
}
if (str == "workgroup_index") {
return BuiltinValue::kWorkgroupIndex;
}
return BuiltinValue::kUndefined;
}
std::string_view ToString(BuiltinValue value) {
@@ -1050,6 +1109,8 @@ std::string_view ToString(BuiltinValue value) {
return "front_facing";
case BuiltinValue::kGlobalInvocationId:
return "global_invocation_id";
case BuiltinValue::kGlobalInvocationIndex:
return "global_invocation_index";
case BuiltinValue::kInstanceIndex:
return "instance_index";
case BuiltinValue::kLocalInvocationId:
@@ -1078,6 +1139,8 @@ std::string_view ToString(BuiltinValue value) {
return "vertex_index";
case BuiltinValue::kWorkgroupId:
return "workgroup_id";
case BuiltinValue::kWorkgroupIndex:
return "workgroup_index";
}
return "<unknown>";
}
@@ -1232,6 +1295,8 @@ std::string_view ToString(ParameterUsage usage) {
return "bias";
case ParameterUsage::kBits:
return "bits";
case ParameterUsage::kColMajor:
return "col_major";
case ParameterUsage::kCompareValue:
return "compare_value";
case ParameterUsage::kComponent:
@@ -1314,6 +1379,8 @@ std::string_view ToString(ParameterUsage usage) {
return "scope";
case ParameterUsage::kSourceLaneIndex:
return "sourceLaneIndex";
case ParameterUsage::kStride:
return "stride";
case ParameterUsage::kTexel:
return "texel";
case ParameterUsage::kTexture:
@@ -1582,9 +1649,6 @@ BuiltinFn ParseBuiltinFn(std::string_view name) {
if (name == "step") {
return BuiltinFn::kStep;
}
if (name == "storageBarrier") {
return BuiltinFn::kStorageBarrier;
}
if (name == "tan") {
return BuiltinFn::kTan;
}
@@ -1618,6 +1682,9 @@ BuiltinFn ParseBuiltinFn(std::string_view name) {
if (name == "unpack4xU8") {
return BuiltinFn::kUnpack4XU8;
}
if (name == "storageBarrier") {
return BuiltinFn::kStorageBarrier;
}
if (name == "workgroupBarrier") {
return BuiltinFn::kWorkgroupBarrier;
}
@@ -1705,6 +1772,12 @@ BuiltinFn ParseBuiltinFn(std::string_view name) {
if (name == "atomicCompareExchangeWeak") {
return BuiltinFn::kAtomicCompareExchangeWeak;
}
if (name == "atomicStoreMax") {
return BuiltinFn::kAtomicStoreMax;
}
if (name == "atomicStoreMin") {
return BuiltinFn::kAtomicStoreMin;
}
if (name == "subgroupBallot") {
return BuiltinFn::kSubgroupBallot;
}
@@ -1981,8 +2054,6 @@ const char* str(BuiltinFn i) {
return "sqrt";
case BuiltinFn::kStep:
return "step";
case BuiltinFn::kStorageBarrier:
return "storageBarrier";
case BuiltinFn::kTan:
return "tan";
case BuiltinFn::kTanh:
@@ -2005,6 +2076,8 @@ const char* str(BuiltinFn i) {
return "unpack4xI8";
case BuiltinFn::kUnpack4XU8:
return "unpack4xU8";
case BuiltinFn::kStorageBarrier:
return "storageBarrier";
case BuiltinFn::kWorkgroupBarrier:
return "workgroupBarrier";
case BuiltinFn::kTextureBarrier:
@@ -2063,6 +2136,10 @@ const char* str(BuiltinFn i) {
return "atomicExchange";
case BuiltinFn::kAtomicCompareExchangeWeak:
return "atomicCompareExchangeWeak";
case BuiltinFn::kAtomicStoreMax:
return "atomicStoreMax";
case BuiltinFn::kAtomicStoreMin:
return "atomicStoreMin";
case BuiltinFn::kSubgroupBallot:
return "subgroupBallot";
case BuiltinFn::kSubgroupElect:
3rdparty/dawn/src/tint/lang/core/enums.h
Vendored
+77 -3
View File
@@ -198,6 +198,66 @@ constexpr std::string_view kInterpolationTypeStrings[] = {
"perspective",
};
/// The sampler filtering.
enum class SamplerFiltering : uint8_t {
kUndefined,
kFiltering,
kNonFiltering,
};
/// @param value the enum value
/// @returns the string for the given enum value
std::string_view ToString(SamplerFiltering value);
/// @param out the stream to write to
/// @param value the SamplerFiltering
/// @returns @p out so calls can be chained
template <typename STREAM>
requires(traits::IsOStream<STREAM>)
auto& operator<<(STREAM& out, SamplerFiltering value) {
return out << ToString(value);
}
/// ParseSamplerFiltering parses a SamplerFiltering from a string.
/// @param str the string to parse
/// @returns the parsed enum, or SamplerFiltering::kUndefined if the string could not be parsed.
SamplerFiltering ParseSamplerFiltering(std::string_view str);
constexpr std::string_view kSamplerFilteringStrings[] = {
"filtering",
"non_filtering",
};
/// The texture filterable.
enum class TextureFilterable : uint8_t {
kUndefined,
kFilterable,
kUnfilterable,
};
/// @param value the enum value
/// @returns the string for the given enum value
std::string_view ToString(TextureFilterable value);
/// @param out the stream to write to
/// @param value the TextureFilterable
/// @returns @p out so calls can be chained
template <typename STREAM>
requires(traits::IsOStream<STREAM>)
auto& operator<<(STREAM& out, TextureFilterable value) {
return out << ToString(value);
}
/// ParseTextureFilterable parses a TextureFilterable from a string.
/// @param str the string to parse
/// @returns the parsed enum, or TextureFilterable::kUndefined if the string could not be parsed.
TextureFilterable ParseTextureFilterable(std::string_view str);
constexpr std::string_view kTextureFilterableStrings[] = {
"filterable",
"unfilterable",
};
/// Address space of a given pointer.
enum class SubgroupMatrixKind : uint8_t {
kUndefined,
@@ -426,6 +486,7 @@ enum class BuiltinType : uint8_t {
kTextureStorage2D,
kTextureStorage2DArray,
kTextureStorage3D,
kU16,
kU32,
kU8,
kVec2,
@@ -551,6 +612,7 @@ constexpr std::string_view kBuiltinTypeStrings[] = {
"texture_storage_2d",
"texture_storage_2d_array",
"texture_storage_3d",
"u16",
"u32",
"u8",
"vec2",
@@ -580,6 +642,7 @@ enum class BuiltinValue : uint8_t {
kFragDepth,
kFrontFacing,
kGlobalInvocationId,
kGlobalInvocationIndex,
kInstanceIndex,
kLocalInvocationId,
kLocalInvocationIndex,
@@ -594,6 +657,7 @@ enum class BuiltinValue : uint8_t {
kSubgroupSize,
kVertexIndex,
kWorkgroupId,
kWorkgroupIndex,
};
/// @param value the enum value
@@ -620,6 +684,7 @@ constexpr std::string_view kBuiltinValueStrings[] = {
"frag_depth",
"front_facing",
"global_invocation_id",
"global_invocation_index",
"instance_index",
"local_invocation_id",
"local_invocation_index",
@@ -634,6 +699,7 @@ constexpr std::string_view kBuiltinValueStrings[] = {
"subgroup_size",
"vertex_index",
"workgroup_id",
"workgroup_index",
};
/// Builtin depth mode defined with `@builtin(<name>, <depth_mode>)`.
@@ -739,6 +805,7 @@ enum class ParameterUsage : uint8_t {
kBase,
kBias,
kBits,
kColMajor,
kCompareValue,
kComponent,
kConstOffset,
@@ -780,6 +847,7 @@ enum class ParameterUsage : uint8_t {
kSamples,
kScope,
kSourceLaneIndex,
kStride,
kTexel,
kTexture,
kValue,
@@ -890,7 +958,6 @@ enum class BuiltinFn : uint8_t {
kSmoothstep,
kSqrt,
kStep,
kStorageBarrier,
kTan,
kTanh,
kTranspose,
@@ -902,6 +969,7 @@ enum class BuiltinFn : uint8_t {
kUnpack4X8Unorm,
kUnpack4XI8,
kUnpack4XU8,
kStorageBarrier,
kWorkgroupBarrier,
kTextureBarrier,
kTextureDimensions,
@@ -931,6 +999,8 @@ enum class BuiltinFn : uint8_t {
kAtomicXor,
kAtomicExchange,
kAtomicCompareExchangeWeak,
kAtomicStoreMax,
kAtomicStoreMin,
kSubgroupBallot,
kSubgroupElect,
kSubgroupBroadcast,
@@ -1070,7 +1140,6 @@ constexpr BuiltinFn kBuiltinFns[] = {
BuiltinFn::kSmoothstep,
BuiltinFn::kSqrt,
BuiltinFn::kStep,
BuiltinFn::kStorageBarrier,
BuiltinFn::kTan,
BuiltinFn::kTanh,
BuiltinFn::kTranspose,
@@ -1082,6 +1151,7 @@ constexpr BuiltinFn kBuiltinFns[] = {
BuiltinFn::kUnpack4X8Unorm,
BuiltinFn::kUnpack4XI8,
BuiltinFn::kUnpack4XU8,
BuiltinFn::kStorageBarrier,
BuiltinFn::kWorkgroupBarrier,
BuiltinFn::kTextureBarrier,
BuiltinFn::kTextureDimensions,
@@ -1111,6 +1181,8 @@ constexpr BuiltinFn kBuiltinFns[] = {
BuiltinFn::kAtomicXor,
BuiltinFn::kAtomicExchange,
BuiltinFn::kAtomicCompareExchangeWeak,
BuiltinFn::kAtomicStoreMax,
BuiltinFn::kAtomicStoreMin,
BuiltinFn::kSubgroupBallot,
BuiltinFn::kSubgroupElect,
BuiltinFn::kSubgroupBroadcast,
@@ -1231,7 +1303,6 @@ constexpr const char* kBuiltinFnStrings[] = {
"smoothstep",
"sqrt",
"step",
"storageBarrier",
"tan",
"tanh",
"transpose",
@@ -1243,6 +1314,7 @@ constexpr const char* kBuiltinFnStrings[] = {
"unpack4x8unorm",
"unpack4xI8",
"unpack4xU8",
"storageBarrier",
"workgroupBarrier",
"textureBarrier",
"textureDimensions",
@@ -1272,6 +1344,8 @@ constexpr const char* kBuiltinFnStrings[] = {
"atomicXor",
"atomicExchange",
"atomicCompareExchangeWeak",
"atomicStoreMax",
"atomicStoreMin",
"subgroupBallot",
"subgroupElect",
"subgroupBroadcast",
3rdparty/dawn/src/tint/lang/core/fluent_types.h
Vendored
+12 -11
View File
@@ -39,13 +39,14 @@ using f16 = tint::core::f16;
using f32 = tint::core::f32;
using i32 = tint::core::i32;
using i8 = tint::core::i8;
using u16 = tint::core::u16;
using u32 = tint::core::u32;
using u64 = tint::core::u64;
using u8 = tint::core::u8;
using AFloat = tint::core::AFloat;
using AInt = tint::core::AInt;
// A sentinel type used by some template arguments to signal that the a type should be inferred.
// A sentinel type used by some template arguments to signal that the type should be inferred.
struct Infer {};
/// A 'fluent' type helper used to construct an ast::Array or type::Array.
@@ -195,20 +196,20 @@ using vec4u = vec<4, u32>;
////////////////////////////////////////////////////////////////////////////////
// Address space aliases
////////////////////////////////////////////////////////////////////////////////
static constexpr core::AddressSpace function = core::AddressSpace::kFunction;
static constexpr core::AddressSpace handle = core::AddressSpace::kHandle;
static constexpr core::AddressSpace private_ = core::AddressSpace::kPrivate;
static constexpr core::AddressSpace immediate = core::AddressSpace::kImmediate;
static constexpr core::AddressSpace storage = core::AddressSpace::kStorage;
static constexpr core::AddressSpace uniform = core::AddressSpace::kUniform;
static constexpr core::AddressSpace workgroup = core::AddressSpace::kWorkgroup;
inline constexpr core::AddressSpace function = core::AddressSpace::kFunction;
inline constexpr core::AddressSpace handle = core::AddressSpace::kHandle;
inline constexpr core::AddressSpace private_ = core::AddressSpace::kPrivate;
inline constexpr core::AddressSpace immediate = core::AddressSpace::kImmediate;
inline constexpr core::AddressSpace storage = core::AddressSpace::kStorage;
inline constexpr core::AddressSpace uniform = core::AddressSpace::kUniform;
inline constexpr core::AddressSpace workgroup = core::AddressSpace::kWorkgroup;
////////////////////////////////////////////////////////////////////////////////
// Access control aliases
////////////////////////////////////////////////////////////////////////////////
static constexpr core::Access read = core::Access::kRead;
static constexpr core::Access read_write = core::Access::kReadWrite;
static constexpr core::Access write = core::Access::kWrite;
inline constexpr core::Access read = core::Access::kRead;
inline constexpr core::Access read_write = core::Access::kReadWrite;
inline constexpr core::Access write = core::Access::kWrite;
////////////////////////////////////////////////////////////////////////////////
// Traits
3rdparty/dawn/src/tint/lang/core/intrinsic/data.cc
Vendored
+13097 -5904
View File
File diff suppressed because it is too large Load Diff
3rdparty/dawn/src/tint/lang/core/intrinsic/table.cc
Vendored
+34 -34
View File
@@ -405,7 +405,7 @@ Candidate ScoreOverload(Context& context,
auto* matcher_indices = context.data[parameter.matcher_indices];
auto* ty =
context.Match(templates, overload, matcher_indices, earliest_eval_stage).Type(args[p]);
parameters.Emplace(ty, parameter.usage);
parameters.Emplace(ty, parameter.usage, parameter.is_const);
}
return Candidate{score, &overload, templates, parameters};
@@ -521,7 +521,7 @@ void PrintCandidate(StyledText& ss,
auto prev_style = ss.Style();
TINT_DEFER(ss << prev_style);
auto& overload = *candidate.overload;
const OverloadInfo& overload = *candidate.overload;
TemplateState templates = candidate.templates;
@@ -533,11 +533,11 @@ void PrintCandidate(StyledText& ss,
if (overload.num_explicit_templates > 0) {
ss << "<";
for (size_t i = 0; i < overload.num_explicit_templates; i++) {
const auto& tmpl = context.data[overload.templates + i];
const TemplateInfo& tmpl = context.data[overload.templates + i];
bool matched = false;
if (i < template_args.Length()) {
auto* matcher_indices = context.data[tmpl.matcher_indices];
const MatcherIndex* matcher_indices = context.data[tmpl.matcher_indices];
matched = (matcher_indices == nullptr) ||
(context.Match(templates, overload, matcher_indices, earliest_eval_stage)
.Type(template_args[i]) != nullptr);
@@ -559,8 +559,8 @@ void PrintCandidate(StyledText& ss,
bool all_params_match = true;
ss << "(";
for (size_t i = 0; i < overload.num_parameters; i++) {
const auto& parameter = context.data[overload.parameters + i];
auto* matcher_indices = context.data[parameter.matcher_indices];
const ParameterInfo& parameter = context.data[overload.parameters + i];
const MatcherIndex* matcher_indices = context.data[parameter.matcher_indices];
bool matched = false;
if (i < args.Length()) {
@@ -588,7 +588,7 @@ void PrintCandidate(StyledText& ss,
ss << ")";
if (overload.return_matcher_indices.IsValid()) {
ss << " -> ";
auto* matcher_indices = context.data[overload.return_matcher_indices];
const MatcherIndex* matcher_indices = context.data[overload.return_matcher_indices];
context.Match(templates, overload, matcher_indices, earliest_eval_stage).PrintType(ss);
}
@@ -615,35 +615,35 @@ void PrintCandidate(StyledText& ss,
}
for (size_t i = 0; i < overload.num_templates; i++) {
auto& tmpl = context.data[overload.templates + i];
if (auto* matcher_indices = context.data[tmpl.matcher_indices]) {
separator();
bool matched = false;
if (tmpl.kind == TemplateInfo::Kind::kType) {
if (auto* ty = templates.Type(i)) {
matched =
(context.Match(templates, overload, matcher_indices, earliest_eval_stage)
.Type(ty) != nullptr);
}
} else {
matched = context.Match(templates, overload, matcher_indices, earliest_eval_stage)
.Num(templates.Num(i))
.IsValid();
}
if (matched) {
ss << style::Match(" ✓ ") << style::Plain(" ");
} else {
ss << style::Mismatch(" ✗ ") << style::Plain(" ");
}
const TemplateInfo& tmpl = context.data[overload.templates + i];
const MatcherIndex* matcher_indices = context.data[tmpl.matcher_indices];
if (!matcher_indices) {
continue;
}
ss << style::Type(tmpl.name) << style::Plain(" is ");
if (tmpl.kind == TemplateInfo::Kind::kType) {
context.Match(templates, overload, matcher_indices, earliest_eval_stage)
.PrintType(ss);
} else {
context.Match(templates, overload, matcher_indices, earliest_eval_stage)
.PrintNum(ss);
separator();
bool matched = false;
if (tmpl.kind == TemplateInfo::Kind::kType) {
if (const core::type::Type* ty = templates.Type(i)) {
matched = (context.Match(templates, overload, matcher_indices, earliest_eval_stage)
.Type(ty) != nullptr);
}
} else {
matched = context.Match(templates, overload, matcher_indices, earliest_eval_stage)
.Num(templates.Num(i))
.IsValid();
}
if (matched) {
ss << style::Match(" ✓ ") << style::Plain(" ");
} else {
ss << style::Mismatch(" ✗ ") << style::Plain(" ");
}
ss << style::Type(tmpl.name) << style::Plain(" is ");
if (tmpl.kind == TemplateInfo::Kind::kType) {
context.Match(templates, overload, matcher_indices, earliest_eval_stage).PrintType(ss);
} else {
context.Match(templates, overload, matcher_indices, earliest_eval_stage).PrintNum(ss);
}
}
}
3rdparty/dawn/src/tint/lang/core/intrinsic/table.h
Vendored
+3 -6
View File
@@ -43,11 +43,6 @@
#include "src/tint/utils/text/string_stream.h"
#include "src/tint/utils/text/styled_text.h"
// Forward declarations
namespace tint::diag {
class List;
} // namespace tint::diag
namespace tint::core::intrinsic {
/// Overload describes a fully matched builtin function overload
@@ -60,12 +55,14 @@ struct Overload {
const core::type::Type* const type;
/// Parameter usage
core::ParameterUsage const usage = core::ParameterUsage::kNone;
/// True if the parameter is required to be const.
const bool is_const;
/// Equality operator
/// @param other the parameter to compare against
/// @returns true if this parameter and @p other are the same
bool operator==(const Parameter& other) const {
return type == other.type && usage == other.usage;
return type == other.type && usage == other.usage && is_const == other.is_const;
}
/// Inequality operator
3rdparty/dawn/src/tint/lang/core/intrinsic/table_data.h
Vendored
+6 -2
View File
@@ -29,6 +29,7 @@
#define SRC_TINT_LANG_CORE_INTRINSIC_TABLE_DATA_H_
#include <stdint.h>
#include <limits>
#include <string>
@@ -116,7 +117,7 @@ struct TableIndex {
};
/// Index type used to index TableData::template_types
using TemplateIndex = TableIndex<TableIndexNamespace::kTemplate, uint8_t>;
using TemplateIndex = TableIndex<TableIndexNamespace::kTemplate, uint16_t>;
/// Index type used to index TableData::type_matchers or TableData::number_matchers
using MatcherIndex = TableIndex<TableIndexNamespace::kMatcher, uint8_t>;
@@ -161,6 +162,9 @@ struct ParameterInfo {
/// The parameter usage (parameter name in definition file)
const ParameterUsage usage;
/// True if the parameter is required to be constant.
const bool is_const;
/// Index of the matcher indices that are used to match the parameter types.
/// These indices are consumed by the matchers themselves.
const MatcherIndicesIndex matcher_indices;
@@ -212,7 +216,7 @@ struct IntrinsicInfo {
};
/// A IntrinsicInfo with no overloads
static constexpr IntrinsicInfo kNoOverloads{0, OverloadIndex(OverloadIndex::kInvalid)};
inline constexpr IntrinsicInfo kNoOverloads{0, OverloadIndex(OverloadIndex::kInvalid)};
/// Number is an 32 bit unsigned integer, which can be in one of three states:
/// * Invalid - Number has not been assigned a value
3rdparty/dawn/src/tint/lang/core/intrinsic/type_matchers.h
Vendored
+80 -2
View File
@@ -55,6 +55,7 @@
#include "src/tint/lang/core/type/storage_texture.h"
#include "src/tint/lang/core/type/string.h"
#include "src/tint/lang/core/type/texture_dimension.h"
#include "src/tint/lang/core/type/u16.h"
#include "src/tint/lang/core/type/u32.h"
#include "src/tint/lang/core/type/u64.h"
#include "src/tint/lang/core/type/u8.h"
@@ -128,6 +129,14 @@ inline bool MatchU8(intrinsic::MatchState&, const type::Type* ty) {
return ty->IsAnyOf<intrinsic::Any, type::U8, type::AbstractInt>();
}
inline const type::U16* BuildU16(intrinsic::MatchState& state, const type::Type*) {
return state.types.u16();
}
inline bool MatchU16(intrinsic::MatchState&, const type::Type* ty) {
return ty->IsAnyOf<intrinsic::Any, type::U16, type::AbstractInt>();
}
inline bool MatchVec(intrinsic::MatchState&,
const type::Type* ty,
intrinsic::Number& N,
@@ -502,13 +511,38 @@ inline bool MatchSampler(intrinsic::MatchState&, const type::Type* ty) {
if (ty->Is<intrinsic::Any>()) {
return true;
}
return ty->Is([](const type::Sampler* s) { return s->Kind() == type::SamplerKind::kSampler; });
return ty->Is([](const type::Sampler* s) {
return s->Kind() == type::SamplerKind::kSampler &&
s->Filtering() == SamplerFiltering::kUndefined;
});
}
inline const type::Sampler* BuildSampler(intrinsic::MatchState& state, const type::Type*) {
return state.types.sampler();
}
inline bool MatchSamplerFiltered(intrinsic::MatchState&,
const type::Type* ty,
intrinsic::Number& F) {
if (ty->Is<intrinsic::Any>()) {
F = intrinsic::Number::any;
return true;
}
if (auto* s = ty->As<type::Sampler>()) {
if (s->Filtering() != SamplerFiltering::kUndefined) {
F = intrinsic::Number(static_cast<uint32_t>(s->Filtering()));
return true;
}
}
return false;
}
inline const type::Sampler* BuildSamplerFiltered(intrinsic::MatchState& state,
const type::Type*,
intrinsic::Number& F) {
return state.types.sampler(static_cast<core::SamplerFiltering>(F.Value()));
}
inline bool MatchSamplerComparison(intrinsic::MatchState&, const type::Type* ty) {
if (ty->Is<intrinsic::Any>()) {
return true;
@@ -531,7 +565,7 @@ inline bool MatchTexture(intrinsic::MatchState&,
return true;
}
if (auto* v = ty->As<type::SampledTexture>()) {
if (v->Dim() == dim) {
if (v->Dim() == dim && v->Filterable() == TextureFilterable::kUndefined) {
T = v->Type();
return true;
}
@@ -559,6 +593,50 @@ DECLARE_SAMPLED_TEXTURE(Cube, type::TextureDimension::kCube)
DECLARE_SAMPLED_TEXTURE(CubeArray, type::TextureDimension::kCubeArray)
#undef DECLARE_SAMPLED_TEXTURE
inline bool MatchTextureFiltered(intrinsic::MatchState&,
const type::Type* ty,
type::TextureDimension dim,
const type::Type*& T,
intrinsic::Number& F) {
if (ty->Is<intrinsic::Any>()) {
T = ty;
F = intrinsic::Number::any;
return true;
}
if (auto* v = ty->As<type::SampledTexture>()) {
if (v->Dim() == dim && v->Filterable() != TextureFilterable::kUndefined) {
T = v->Type();
F = intrinsic::Number(static_cast<uint32_t>(v->Filterable()));
return true;
}
}
return false;
}
#define JOIN2(a, b, c) a##b##c
#define DECLARE_SAMPLED_TEXTURE_FILTERED(suffix, dim) \
inline bool JOIN2(MatchTexture, suffix, Filtered)(intrinsic::MatchState & state, \
const type::Type* ty, const type::Type*& T, \
intrinsic::Number& F) { \
return MatchTextureFiltered(state, ty, dim, T, F); \
} \
inline const type::SampledTexture* JOIN2(BuildTexture, suffix, Filtered)( \
intrinsic::MatchState & state, const type::Type*, const type::Type* T, \
intrinsic::Number& F) { \
return state.types.sampled_texture(dim, T, \
static_cast<core::TextureFilterable>(F.Value())); \
}
DECLARE_SAMPLED_TEXTURE_FILTERED(1D, type::TextureDimension::k1d)
DECLARE_SAMPLED_TEXTURE_FILTERED(2D, type::TextureDimension::k2d)
DECLARE_SAMPLED_TEXTURE_FILTERED(2DArray, type::TextureDimension::k2dArray)
DECLARE_SAMPLED_TEXTURE_FILTERED(3D, type::TextureDimension::k3d)
DECLARE_SAMPLED_TEXTURE_FILTERED(Cube, type::TextureDimension::kCube)
DECLARE_SAMPLED_TEXTURE_FILTERED(CubeArray, type::TextureDimension::kCubeArray)
#undef DECLARE_SAMPLED_TEXTURE_FILTERED
#undef JOIN2
inline bool MatchTextureMultisampled(intrinsic::MatchState&,
const type::Type* ty,
type::TextureDimension dim,
3rdparty/dawn/src/tint/lang/core/ir/builder.h
Vendored
+27
View File
@@ -83,6 +83,7 @@
#include "src/tint/lang/core/type/memory_view.h"
#include "src/tint/lang/core/type/pointer.h" // IWYU pragma: export
#include "src/tint/lang/core/type/type.h" // IWYU pragma: export
#include "src/tint/lang/core/type/u16.h" // IWYU pragma: export
#include "src/tint/lang/core/type/u32.h" // IWYU pragma: export
#include "src/tint/lang/core/type/u64.h" // IWYU pragma: export
#include "src/tint/lang/core/type/u8.h" // IWYU pragma: export
@@ -380,6 +381,11 @@ class Builder {
/// @returns the new constant
ir::Constant* Constant(core::u8 v) { return Constant(ConstantValue(v)); }
/// Creates a ir::Constant for a u16 Scalar
/// @param v the value
/// @returns the new constant
ir::Constant* Constant(core::u16 v) { return Constant(ConstantValue(v)); }
/// Creates a ir::Constant for a f32 Scalar
/// @param v the value
/// @returns the new constant
@@ -433,6 +439,11 @@ class Builder {
/// @returns the new constant
const core::constant::Value* ConstantValue(core::u8 v) { return ir.constant_values.Get(v); }
/// Creates a core::constant::Value for a u16 Scalar
/// @param v the value
/// @returns the new constant
const core::constant::Value* ConstantValue(core::u16 v) { return ir.constant_values.Get(v); }
/// Creates a core::constant::Value for a f32 Scalar
/// @param v the value
/// @returns the new constant
@@ -1346,6 +1357,22 @@ class Builder {
return val->Type()->Equals(*to) ? val : Convert(to, val)->Result();
}
/// Adds a call to bitcast if destination type is different then the value's type
/// @param to the type converted to
/// @param val the value to be converted
/// @returns either result of the conversion or original value
ir::Value* InsertBitcastIfNeeded(const core::type::Type* to, ir::Value* val) {
return val->Type()->Equals(*to) ? val : Bitcast(to, val)->Result();
}
/// Adds a call to bitcast if destination type is different then the instruction's type
/// @param to the type converted to
/// @param inst the instruction to be converted
/// @returns either result of the conversion or original instruction
ir::Instruction* InsertBitcastIfNeeded(const core::type::Type* to, ir::Instruction* inst) {
return inst->Result()->Type()->Equals(*to) ? inst : Bitcast(to, inst);
}
/// Creates a value constructor instruction with an existing instruction result
/// @param result the instruction result to use
/// @param args the arguments to the constructor
3rdparty/dawn/src/tint/lang/core/ir/constant.cc
Vendored
+1
View File
@@ -26,6 +26,7 @@
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
#include "src/tint/lang/core/ir/constant.h"
#include "src/tint/utils/ice/ice.h"
TINT_INSTANTIATE_TYPEINFO(tint::core::ir::Constant);
3rdparty/dawn/src/tint/lang/core/ir/core_builtin_call.cc
Vendored
+2
View File
@@ -85,6 +85,8 @@ tint::core::ir::Instruction::Accesses CoreBuiltinCall::GetSideEffects() const {
case BuiltinFn::kAtomicXor:
case BuiltinFn::kAtomicExchange:
case BuiltinFn::kAtomicCompareExchangeWeak:
case BuiltinFn::kAtomicStoreMax:
case BuiltinFn::kAtomicStoreMin:
case BuiltinFn::kDpdx:
case BuiltinFn::kDpdxCoarse:
case BuiltinFn::kDpdxFine:
3rdparty/dawn/src/tint/lang/core/ir/disassembler.cc
Vendored
+3
View File
@@ -449,6 +449,9 @@ StyledText Disassembler::ValueToStyledText(const Value* val) {
[&](const core::constant::Scalar<u8>* scalar) {
text << StyleLiteral(u32(scalar->ValueAs<u8>().value), "u8");
},
[&](const core::constant::Scalar<u16>* scalar) {
text << StyleLiteral(u32(scalar->ValueAs<u16>().value), "u16");
},
[&](const core::constant::Scalar<f32>* scalar) {
text << StyleLiteral(scalar->ValueAs<f32>().value, "f");
},
3rdparty/dawn/src/tint/lang/core/ir/evaluator.cc
Vendored
+1
View File
@@ -26,6 +26,7 @@
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
#include "src/tint/lang/core/ir/evaluator.h"
#include "src/tint/lang/core/binary_op.h"
#include "src/tint/lang/core/ir/constant.h"
#include "src/tint/lang/core/ir/constexpr_if.h"
3rdparty/dawn/src/tint/lang/core/ir/function.cc
Vendored
+2
View File
@@ -27,6 +27,8 @@
#include "src/tint/lang/core/ir/function.h"
#include <utility>
#include "src/tint/lang/core/ir/clone_context.h"
#include "src/tint/lang/core/ir/module.h"
#include "src/tint/utils/containers/predicates.h"
3rdparty/dawn/src/tint/lang/core/ir/function.h
Vendored
-2
View File
@@ -30,7 +30,6 @@
#include <array>
#include <optional>
#include <utility>
#include "src/tint/lang/core/io_attributes.h"
#include "src/tint/lang/core/ir/constant.h"
@@ -42,7 +41,6 @@
// Forward declarations
namespace tint::core::ir {
class Block;
class FunctionTerminator;
} // namespace tint::core::ir
namespace tint::core::ir {
3rdparty/dawn/src/tint/lang/core/ir/if.h
Vendored
-5
View File
@@ -32,11 +32,6 @@
#include "src/tint/lang/core/ir/control_instruction.h"
// Forward declarations
namespace tint::core::ir {
class MultiInBlock;
} // namespace tint::core::ir
namespace tint::core::ir {
/// If instruction.
3rdparty/dawn/src/tint/lang/core/ir/module.cc
Vendored
+18 -10
View File
@@ -28,6 +28,7 @@
#include "src/tint/lang/core/ir/module.h"
#include <limits>
#include <type_traits>
#include <utility>
#include "src/tint/lang/core/ir/control_instruction.h"
@@ -39,19 +40,22 @@ namespace tint::core::ir {
namespace {
/// Helper to non-recursively sort a module's function in dependency order.
/// Helper to non-recursively sort a module's functions in dependency order.
template <typename FUNCTION>
struct FunctionSorter {
using BLOCK = std::conditional_t<std::is_const_v<FUNCTION>, const Block, Block>;
/// The dependency-ordered list of functions.
UniqueVector<Function*, 16> ordered_functions{};
UniqueVector<FUNCTION*, 16> ordered_functions{};
/// The functions that have been visited and checked for dependencies.
Hashset<Function*, 16> visited{};
Hashset<FUNCTION*, 16> visited{};
/// A stack of functions that need to processed and eventually added to the ordered list.
Vector<Function*, 16> function_stack{};
Vector<FUNCTION*, 16> function_stack{};
/// Visit a function and check for dependencies, and eventually add it to the ordered list.
/// @param func the function to visit
void Visit(Function* func) {
void Visit(FUNCTION* func) {
function_stack.Push(func);
while (!function_stack.IsEmpty()) {
// Visit the next function on the stack, if it hasn't already been visited.
@@ -76,15 +80,15 @@ struct FunctionSorter {
/// Visit a function body block and look for dependencies.
/// @param block the function body to visit
void Visit(Block* block) {
Vector<Block*, 64> block_stack;
void Visit(BLOCK* block) {
Vector<BLOCK*, 64> block_stack;
block_stack.Push(block);
while (!block_stack.IsEmpty()) {
auto* current_block = block_stack.Pop();
for (auto* inst : *current_block) {
if (auto* control = inst->template As<ControlInstruction>()) {
// Enqueue child blocks.
control->ForeachBlock([&](Block* b) { block_stack.Push(b); });
control->ForeachBlock([&](BLOCK* b) { block_stack.Push(b); });
} else if (auto* call = inst->template As<UserCall>()) {
// Enqueue the function that is being called.
if (!visited.Contains(call->Target())) {
@@ -98,7 +102,7 @@ struct FunctionSorter {
/// Sort the functions of a module.
/// @param mod the IR module
/// @returns the sorted function list
static Vector<Function*, 16> SortFunctions(Module& mod) {
static Vector<FUNCTION*, 16> SortFunctions(const Module& mod) {
FunctionSorter sorter;
for (auto& func : mod.functions) {
sorter.Visit(func);
@@ -175,7 +179,11 @@ Source Module::SourceOf(const Value* value) const {
}
Vector<Function*, 16> Module::DependencyOrderedFunctions() {
return FunctionSorter::SortFunctions(*this);
return FunctionSorter<Function>::SortFunctions(*this);
}
Vector<const Function*, 16> Module::DependencyOrderedFunctions() const {
return FunctionSorter<const Function>::SortFunctions(*this);
}
void Module::Destroy(Function* func) {
3rdparty/dawn/src/tint/lang/core/ir/reflection.h
Vendored
+1 -15
View File
@@ -30,25 +30,11 @@
#include <optional>
#include "src/tint/api/common/workgroup_info.h"
#include "src/tint/lang/core/ir/module.h"
#include "src/tint/utils/result.h"
namespace tint::core::ir {
/// Workgroup size information
struct WorkgroupInfo {
/// The x-component
uint32_t x = 0;
/// The y-component
uint32_t y = 0;
/// The z-component
uint32_t z = 0;
/// The needed workgroup storage size
size_t storage_size = 0;
/// The `@subgroup_size` attribute
std::optional<uint32_t> subgroup_size;
};
/// Generate WorkgroupInfo for an IR module input.
/// @param ir the workgroup info for the IR module
3rdparty/dawn/src/tint/lang/core/ir/switch.h
Vendored
-1
View File
@@ -36,7 +36,6 @@
// Forward declarations
namespace tint::core::ir {
class Constant;
class MultiInBlock;
} // namespace tint::core::ir
namespace tint::core::ir {
3rdparty/dawn/src/tint/lang/core/ir/terminator.h
Vendored
-5
View File
@@ -32,11 +32,6 @@
#include "src/tint/lang/core/ir/value.h"
#include "src/tint/utils/rtti/castable.h"
// Forward declarations
namespace tint::core::ir {
class Block;
} // namespace tint::core::ir
namespace tint::core::ir {
/// The base class of all instructions that terminate a block.
3rdparty/dawn/src/tint/lang/core/ir/transform/array_length_from_immediate.cc
Vendored
+2 -2
View File
@@ -350,8 +350,8 @@ Result<ArrayLengthFromImmediateResult> ArrayLengthFromImmediates(
const uint32_t buffer_sizes_offset,
const uint32_t buffer_sizes_array_elements_num,
const std::unordered_map<BindingPoint, uint32_t>& bindpoint_to_size_index) {
TINT_CHECK_RESULT(ValidateAndDumpIfNeeded(ir, "core.ArrayLengthFromImmediates",
kArrayLengthFromImmediateCapabilities));
TINT_CHECK_RESULT(ValidateBeforeIfNeeded(ir, kArrayLengthFromImmediateCapabilities,
"core.ArrayLengthFromImmediates"));
State state{ir, immediate_data_layout, buffer_sizes_offset, buffer_sizes_array_elements_num,
bindpoint_to_size_index};
3rdparty/dawn/src/tint/lang/core/ir/transform/array_length_from_uniform.cc
Vendored
+3 -3
View File
@@ -316,7 +316,7 @@ struct State {
const uint32_t size_index = bindpoint_to_size_index.at(info.binding_point);
const uint32_t array_index = size_index / 4;
const uint32_t vec_index = size_index % 4;
auto* vec_ptr = b.Access<ptr<uniform, vec4<u32>>>(BufferSizes(), u32(array_index));
auto* vec_ptr = b.Access<ptr<uniform, vec4u>>(BufferSizes(), u32(array_index));
auto* total_buffer_size = b.LoadVectorElement(vec_ptr, u32(vec_index))->Result();
// Calculate actual array length:
@@ -361,8 +361,8 @@ Result<ArrayLengthFromUniformResult> ArrayLengthFromUniform(
Module& ir,
BindingPoint ubo_binding,
const std::unordered_map<BindingPoint, uint32_t>& bindpoint_to_size_index) {
TINT_CHECK_RESULT(ValidateAndDumpIfNeeded(ir, "core.ArrayLengthFromUniform",
kArrayLengthFromUniformCapabilities));
TINT_CHECK_RESULT(ValidateBeforeIfNeeded(ir, kArrayLengthFromUniformCapabilities,
"core.ArrayLengthFromUniform"));
State state{ir, ubo_binding, bindpoint_to_size_index};
state.Process();
3rdparty/dawn/src/tint/lang/core/ir/transform/array_length_from_uniform.h
Vendored
+1 -1
View File
@@ -29,9 +29,9 @@
#define SRC_TINT_LANG_CORE_IR_TRANSFORM_ARRAY_LENGTH_FROM_UNIFORM_H_
#include <unordered_map>
#include "src/tint/lang/core/ir/validator.h"
#include "src/tint/api/common/binding_point.h"
#include "src/tint/lang/core/ir/validator.h"
#include "src/tint/utils/result.h"
// Forward declarations.
3rdparty/dawn/src/tint/lang/core/ir/transform/bgra8unorm_polyfill.cc
Vendored
+49 -143
View File
@@ -27,8 +27,6 @@
#include "src/tint/lang/core/ir/transform/bgra8unorm_polyfill.h"
#include <utility>
#include "src/tint/lang/core/ir/builder.h"
#include "src/tint/lang/core/ir/module.h"
#include "src/tint/lang/core/ir/validator.h"
@@ -55,167 +53,76 @@ struct State {
/// Process the module.
void Process() {
// Find module-scope variables that need to be replaced.
if (!ir.root_block->IsEmpty()) {
Vector<Instruction*, 4> to_remove;
for (auto inst : *ir.root_block) {
auto* var = inst->As<Var>();
if (!var) {
continue;
}
auto* ptr = var->Result()->Type()->As<core::type::Pointer>();
if (!ptr) {
continue;
}
auto* storage_texture = ptr->StoreType()->As<core::type::StorageTexture>();
if (storage_texture &&
storage_texture->TexelFormat() == core::TexelFormat::kBgra8Unorm) {
ReplaceVar(var, storage_texture);
to_remove.Push(var);
continue;
}
auto* texel_buffer = ptr->StoreType()->As<core::type::TexelBuffer>();
if (texel_buffer && texel_buffer->TexelFormat() == core::TexelFormat::kBgra8Unorm) {
ReplaceVar(var, texel_buffer);
to_remove.Push(var);
}
}
for (auto* remove : to_remove) {
remove->Destroy();
// Find module-scope variables that may need to be updated.
for (auto inst : *ir.root_block) {
if (auto* var = inst->As<Var>()) {
ChangeTypeIfNeeded(var->Result());
}
}
// Find function parameters that need to be replaced.
// Find function parameters that may need to be updated.
for (auto& func : ir.functions) {
for (uint32_t index = 0; index < func->Params().Length(); index++) {
auto* param = func->Params()[index];
auto* storage_texture = param->Type()->As<core::type::StorageTexture>();
if (storage_texture &&
storage_texture->TexelFormat() == core::TexelFormat::kBgra8Unorm) {
ReplaceParameter(func, param, index, storage_texture);
continue;
}
auto* texel_buffer = param->Type()->As<core::type::TexelBuffer>();
if (texel_buffer && texel_buffer->TexelFormat() == core::TexelFormat::kBgra8Unorm) {
ReplaceParameter(func, param, index, texel_buffer);
}
for (auto* param : func->Params()) {
ChangeTypeIfNeeded(param);
}
}
}
/// Replace a variable declaration with one that uses rgba8unorm instead of bgra8unorm.
/// @param old_var the variable declaration to replace
/// @param bgra8 the bgra8unorm texture type
void ReplaceVar(Var* old_var, const core::type::StorageTexture* bgra8) {
// Redeclare the variable with a rgba8unorm texel format.
auto* rgba8 =
ty.storage_texture(bgra8->Dim(), core::TexelFormat::kRgba8Unorm, bgra8->Access());
auto* new_var = b.Var(ty.ptr(handle, rgba8));
auto bp = old_var->BindingPoint();
new_var->SetBindingPoint(bp->group, bp->binding);
new_var->InsertBefore(old_var);
if (auto name = ir.NameOf(old_var)) {
ir.SetName(new_var, name.NameView());
/// Update a value's type to use rgba8unorm instead of bgra8unorm if needed.
/// @param value the value that may have a bgra8unorm type
void ChangeTypeIfNeeded(Value* value) {
// Check if the type has a bgra8unorm texel format and create the rgba8unorm equivalent.
auto* rgba8unorm_type = tint::Switch(
value->Type()->UnwrapPtr(),
[&](const core::type::StorageTexture* st) -> const core::type::Type* {
if (st->TexelFormat() == TexelFormat::kBgra8Unorm) {
return ty.storage_texture(st->Dim(), TexelFormat::kRgba8Unorm, st->Access());
}
return nullptr;
},
[&](const core::type::TexelBuffer* tb) -> const core::type::Type* {
if (tb->TexelFormat() == TexelFormat::kBgra8Unorm) {
return ty.texel_buffer(TexelFormat::kRgba8Unorm, tb->Access());
}
return nullptr;
});
if (!rgba8unorm_type) {
// No change needed.
return;
}
// Replace all uses of the old variable with the new one.
ReplaceUses(old_var->Result(), new_var->Result());
}
/// Replace a variable declaration with one that uses rgba8unorm instead of bgra8unorm.
/// @param old_var the variable declaration to replace
/// @param bgra8 the bgra8unorm texel buffer type
void ReplaceVar(Var* old_var, const core::type::TexelBuffer* bgra8) {
// Redeclare the variable with a rgba8unorm texel format.
auto* rgba8 = ty.texel_buffer(core::TexelFormat::kRgba8Unorm, bgra8->Access());
auto* new_var = b.Var(ty.ptr(handle, rgba8));
auto bp = old_var->BindingPoint();
new_var->SetBindingPoint(bp->group, bp->binding);
new_var->InsertBefore(old_var);
if (auto name = ir.NameOf(old_var)) {
ir.SetName(new_var, name.NameView());
// If the original type was a pointer, make the new type a pointer too.
if (value->Type()->Is<core::type::Pointer>()) {
rgba8unorm_type = ty.ptr(handle, rgba8unorm_type);
}
// Replace all uses of the old variable with the new one.
ReplaceUses(old_var->Result(), new_var->Result());
// Change the value to use the new rgba8unorm type.
value->SetType(rgba8unorm_type);
// Update all uses of the value.
UpdateUses(value);
}
/// Replace a function parameter with one that uses rgba8unorm instead of bgra8unorm.
/// @param func the function
/// @param old_param the function parameter to replace
/// @param index the index of the function parameter
/// @param bgra8 the bgra8unorm texture type
void ReplaceParameter(Function* func,
FunctionParam* old_param,
uint32_t index,
const core::type::StorageTexture* bgra8) {
// Redeclare the parameter with a rgba8unorm texel format.
auto* rgba8 =
ty.storage_texture(bgra8->Dim(), core::TexelFormat::kRgba8Unorm, bgra8->Access());
auto* new_param = b.FunctionParam(rgba8);
if (auto name = ir.NameOf(old_param)) {
ir.SetName(new_param, name.NameView());
}
Vector<FunctionParam*, 4> new_params = func->Params();
new_params[index] = new_param;
func->SetParams(std::move(new_params));
// Replace all uses of the old parameter with the new one.
ReplaceUses(old_param, new_param);
}
/// Replace a function parameter with one that uses rgba8unorm instead of bgra8unorm.
/// @param func the function
/// @param old_param the function parameter to replace
/// @param index the index of the function parameter
/// @param bgra8 the bgra8unorm texel buffer type
void ReplaceParameter(Function* func,
FunctionParam* old_param,
uint32_t index,
const core::type::TexelBuffer* bgra8) {
// Redeclare the parameter with a rgba8unorm texel format.
auto* rgba8 = ty.texel_buffer(core::TexelFormat::kRgba8Unorm, bgra8->Access());
auto* new_param = b.FunctionParam(rgba8);
if (auto name = ir.NameOf(old_param)) {
ir.SetName(new_param, name.NameView());
}
Vector<FunctionParam*, 4> new_params = func->Params();
new_params[index] = new_param;
func->SetParams(std::move(new_params));
// Replace all uses of the old parameter with the new one.
ReplaceUses(old_param, new_param);
}
/// Recursively replace the uses of @p value with @p new_value.
/// @param old_value the value whose usages should be replaced
/// @param new_value the value to use instead
void ReplaceUses(Value* old_value, Value* new_value) {
old_value->ForEachUseUnsorted([&](Usage use) {
/// Recursively update the uses of @p value.
/// @param value the value to update
void UpdateUses(Value* value) {
value->ForEachUseUnsorted([&](Usage use) {
tint::Switch(
use.instruction,
[&](Load* load) {
// Replace load instructions with new ones that have the updated type.
auto* new_load = b.Load(new_value);
new_load->InsertBefore(load);
ReplaceUses(load->Result(), new_load->Result());
load->Destroy();
// Update the result type of the load instruction and then update its uses too.
load->Result()->SetType(value->Type()->UnwrapPtr());
UpdateUses(load->Result());
},
[&](CoreBuiltinCall* call) {
// Replace arguments to builtin functions and add swizzles if necessary.
call->SetOperand(use.operand_index, new_value);
if (call->Func() == core::BuiltinFn::kTextureStore) {
// Swizzle the value argument of a `textureStore()` builtin.
uint32_t index = 2u;
if (auto* tex = old_value->Type()->As<core::type::StorageTexture>()) {
if (auto* tex = value->Type()->As<core::type::StorageTexture>()) {
index = core::type::IsTextureArray(tex->Dim()) ? 3u : 2u;
}
auto* value = call->Args()[index];
auto* swizzle = b.Swizzle(value->Type(), value, Vector{2u, 1u, 0u, 3u});
auto* texel = call->Args()[index];
auto* swizzle = b.Swizzle(texel->Type(), texel, Vector{2u, 1u, 0u, 3u});
swizzle->InsertBefore(call);
call->SetOperand(index, swizzle->Result());
} else if (call->Func() == core::BuiltinFn::kTextureLoad) {
@@ -227,9 +134,8 @@ struct State {
swizzle->SetOperand(Swizzle::kObjectOperandOffset, call->Result());
}
},
[&](UserCall* call) {
// Just replace arguments to user functions and then stop.
call->SetOperand(use.operand_index, new_value);
[&](UserCall*) {
// Nothing to do since the parameter will already have been updated.
},
TINT_ICE_ON_NO_MATCH);
});
@@ -240,7 +146,7 @@ struct State {
Result<SuccessType> Bgra8UnormPolyfill(Module& ir) {
TINT_CHECK_RESULT(
ValidateAndDumpIfNeeded(ir, "core.Bgra8UnormPolyfill", kBgra8UnormPolyfillCapabilities));
ValidateBeforeIfNeeded(ir, kBgra8UnormPolyfillCapabilities, "core.Bgra8UnormPolyfill"));
State{ir}.Process();
3rdparty/dawn/src/tint/lang/core/ir/transform/binary_polyfill.cc
Vendored
+1 -1
View File
@@ -200,7 +200,7 @@ struct State {
Result<SuccessType> BinaryPolyfill(Module& ir, const BinaryPolyfillConfig& config) {
TINT_CHECK_RESULT(
ValidateAndDumpIfNeeded(ir, "core.BinaryPolyfill", kBinaryPolyfillCapabilities));
ValidateBeforeIfNeeded(ir, kBinaryPolyfillCapabilities, "core.BinaryPolyfill"));
State{config, ir}.Process();
3rdparty/dawn/src/tint/lang/core/ir/transform/binding_remapper.cc
Vendored
+1 -1
View File
@@ -73,7 +73,7 @@ Result<SuccessType> BindingRemapper(
Module& ir,
const std::unordered_map<BindingPoint, BindingPoint>& binding_points) {
TINT_CHECK_RESULT(
ValidateAndDumpIfNeeded(ir, "core.BindingRemapper", kBindingRemapperCapabilities));
ValidateBeforeIfNeeded(ir, kBindingRemapperCapabilities, "core.BindingRemapper"));
Run(ir, binding_points);
3rdparty/dawn/src/tint/lang/core/ir/transform/block_decorated_structs.cc
Vendored
+2 -2
View File
@@ -118,8 +118,8 @@ void Run(Module& ir) {
} // namespace
Result<SuccessType> BlockDecoratedStructs(Module& ir) {
TINT_CHECK_RESULT(ValidateAndDumpIfNeeded(ir, "core.BlockDecoratedStructs",
kBlockDecoratedStructsCapabilities));
TINT_CHECK_RESULT(ValidateBeforeIfNeeded(ir, kBlockDecoratedStructsCapabilities,
"core.BlockDecoratedStructs"));
Run(ir);
3rdparty/dawn/src/tint/lang/core/ir/transform/builtin_polyfill.cc
Vendored
+14 -9
View File
@@ -622,9 +622,10 @@ struct State {
// result = extractBits(e, offset, count)
// With:
// let s = min(offset, 32u);
// let t = min(32u, (s + count));
// let clamped_count = min(count, 32u);
// let t = min(32u, (s + clamped_count));
// let shl = (32u - t);
// let shr = (shl + s
// let shr = (shl + s);
// let shl_result = select(i32(), (e << shl), (shl < 32u));
// result = select(((shl_result >> 31u) >> 1u), (shl_result >> shr), (shr < 32u));
// }
@@ -634,7 +635,8 @@ struct State {
auto V = [&](uint32_t u) { return b.MatchWidth(u32(u), result_ty); };
b.InsertBefore(call, [&] {
auto* s = b.Min(offset, 32_u);
auto* t = b.Min(32_u, b.Add(s, count));
auto* clamped_count = b.Min(count, 32_u);
auto* t = b.Min(32_u, b.Add(s, clamped_count));
auto* shl = b.Subtract(32_u, t);
auto* shr = b.Add(shl, s);
auto* f1 = b.Zero(result_ty);
@@ -832,7 +834,10 @@ struct State {
};
b.InsertBefore(call, [&] {
auto* oc = b.Add(offset, count);
auto* clamped_count = b.Min(count, 32_u);
auto* clamped_offset = b.Min(offset, 32_u);
auto* oc = b.Add(clamped_offset, clamped_count);
auto* t1 = b.ShiftLeft(1_u, offset);
auto* s1 = b.Call<u32>(core::BuiltinFn::kSelect, b.Zero<u32>(), t1,
b.LessThan(offset, 32_u));
@@ -952,10 +957,10 @@ struct State {
auto* sampler = call->Args()[1];
auto* coords = call->Args()[2];
b.InsertBefore(call, [&] {
auto* dims = b.Call<vec2<u32>>(core::BuiltinFn::kTextureDimensions, texture);
auto* fdims = b.Convert<vec2<f32>>(dims);
auto* half_texel = b.Divide(b.Splat<vec2<f32>>(0.5_f), fdims);
auto* one_minus_half_texel = b.Subtract(b.Splat<vec2<f32>>(1_f), half_texel);
auto* dims = b.Call<vec2u>(core::BuiltinFn::kTextureDimensions, texture);
auto* fdims = b.Convert<vec2f>(dims);
auto* half_texel = b.Divide(b.Splat<vec2f>(0.5_f), fdims);
auto* one_minus_half_texel = b.Subtract(b.Splat<vec2f>(1_f), half_texel);
auto* clamped = b.Clamp(coords, half_texel, one_minus_half_texel);
b.CallWithResult(call->DetachResult(), core::BuiltinFn::kTextureSampleLevel, texture,
sampler, clamped, 0_f);
@@ -1239,7 +1244,7 @@ struct State {
Result<SuccessType> BuiltinPolyfill(Module& ir, const BuiltinPolyfillConfig& config) {
TINT_CHECK_RESULT(
ValidateAndDumpIfNeeded(ir, "core.BuiltinPolyfill", kBuiltinPolyfillCapabilities));
ValidateBeforeIfNeeded(ir, kBuiltinPolyfillCapabilities, "core.BuiltinPolyfill"));
State{config, ir}.Process();
3rdparty/dawn/src/tint/lang/core/ir/transform/builtin_scalarize.cc
Vendored
+2 -1
View File
@@ -26,6 +26,7 @@
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
#include "src/tint/lang/core/ir/transform/builtin_scalarize.h"
#include <cstdint>
#include <utility>
@@ -135,7 +136,7 @@ struct State {
Result<SuccessType> BuiltinScalarize(Module& ir, const BuiltinScalarizeConfig& config) {
TINT_CHECK_RESULT(
ValidateAndDumpIfNeeded(ir, "core.BuiltinScalarize", kBuiltinScalarizeCapabilities));
ValidateBeforeIfNeeded(ir, kBuiltinScalarizeCapabilities, "core.BuiltinScalarize"));
State{config, ir}.Process();
3rdparty/dawn/src/tint/lang/core/ir/transform/change_immediate_to_uniform.cc
Vendored
+10 -9
View File
@@ -148,15 +148,16 @@ struct State {
Result<SuccessType> ChangeImmediateToUniform(core::ir::Module& ir,
const ChangeImmediateToUniformConfig& config) {
TINT_CHECK_RESULT(
ValidateAndDumpIfNeeded(ir, "core.ChangeImmediateToUniform",
core::ir::Capabilities{
core::ir::Capability::kAllow8BitIntegers,
core::ir::Capability::kAllowPointSizeBuiltin,
core::ir::Capability::kAllowClipDistancesOnF32ScalarAndVector,
core::ir::Capability::kAllowDuplicateBindings,
core::ir::Capability::kAllowNonCoreTypes,
core::ir::Capability::kMslAllowEntryPointInterface,
}));
ValidateBeforeIfNeeded(ir,
core::ir::Capabilities{
core::ir::Capability::kAllow8BitIntegers,
core::ir::Capability::kAllowPointSizeBuiltin,
core::ir::Capability::kAllowClipDistancesOnF32ScalarAndVector,
core::ir::Capability::kAllowDuplicateBindings,
core::ir::Capability::kAllowNonCoreTypes,
core::ir::Capability::kMslAllowEntryPointInterface,
},
"core.ChangeImmediateToUniform"));
State{config, ir}.Process();
3rdparty/dawn/src/tint/lang/core/ir/transform/combine_access_instructions.cc
Vendored
+7 -6
View File
@@ -80,12 +80,13 @@ struct State {
} // namespace
Result<SuccessType> CombineAccessInstructions(Module& ir) {
TINT_CHECK_RESULT(ValidateAndDumpIfNeeded(ir, "core.CombineAccessInstructions",
core::ir::Capabilities{
core::ir::Capability::kAllowDuplicateBindings,
core::ir::Capability::kAllowNonCoreTypes,
core::ir::Capability::kAllow8BitIntegers,
}));
TINT_CHECK_RESULT(ValidateBeforeIfNeeded(ir,
core::ir::Capabilities{
core::ir::Capability::kAllowDuplicateBindings,
core::ir::Capability::kAllowNonCoreTypes,
core::ir::Capability::kAllow8BitIntegers,
},
"core.CombineAccessInstructions"));
State{ir}.Process();
3rdparty/dawn/src/tint/lang/core/ir/transform/conversion_polyfill.cc
Vendored
+1 -2
View File
@@ -117,7 +117,6 @@ struct State {
ir::Constant* high_limit_f = nullptr;
} limits;
// Largest integers representable in the source floating point format.
if (src_el_ty->Is<type::F32>()) {
// These values are chosen specifically to enable f32 clamping.
@@ -167,7 +166,7 @@ struct State {
Result<SuccessType> ConversionPolyfill(Module& ir, const ConversionPolyfillConfig& config) {
TINT_CHECK_RESULT(
ValidateAndDumpIfNeeded(ir, "core.ConversionPolyfill", kConversionPolyfillCapabilities));
ValidateBeforeIfNeeded(ir, kConversionPolyfillCapabilities, "core.ConversionPolyfill"));
State{config, ir}.Process();
3rdparty/dawn/src/tint/lang/core/ir/transform/dead_code_elimination.cc
Vendored
+1 -1
View File
@@ -106,7 +106,7 @@ struct State {
Result<SuccessType> DeadCodeElimination(Module& ir) {
TINT_CHECK_RESULT(
ValidateAndDumpIfNeeded(ir, "core.DeadCodeElimination", kDeadCodeEliminationCapabilities));
ValidateBeforeIfNeeded(ir, kDeadCodeEliminationCapabilities, "core.DeadCodeElimination"));
State{ir}.Process();
3rdparty/dawn/src/tint/lang/core/ir/transform/decompose_access.cc
Vendored
+1350
View File
File diff suppressed because it is too large Load Diff
3rdparty/dawn/src/tint/lang/core/ir/transform/decompose_access.h
Vendored
+66
View File
@@ -0,0 +1,66 @@
// Copyright 2024 The Dawn & Tint Authors
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are met:
//
// 1. Redistributions of source code must retain the above copyright notice, this
// list of conditions and the following disclaimer.
//
// 2. Redistributions in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// 3. Neither the name of the copyright holder nor the names of its
// contributors may be used to endorse or promote products derived from
// this software without specific prior written permission.
//
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
#ifndef SRC_TINT_LANG_CORE_IR_TRANSFORM_DECOMPOSE_ACCESS_H_
#define SRC_TINT_LANG_CORE_IR_TRANSFORM_DECOMPOSE_ACCESS_H_
#include <cstdint>
#include "src/tint/utils/result.h"
// Forward declarations.
namespace tint::core::ir {
class Module;
} // namespace tint::core::ir
namespace tint::core::ir::transform {
struct DecomposeAccessOptions {
/// Specify whether all variables in the address space should be decomposed.
bool storage = false;
bool uniform = false;
bool workgroup = false;
bool immediate = false;
// The minimum size of the resulting array, in bytes.
uint32_t minimum_array_size = 0;
// TODO(b/477295042): should there be a uniform standard layout option? When enabled uniform
// could be treated like every other storage class.
};
/// DecomposeAccess is a transform used to replace uniform, storage, or workgroup variables
/// with an a uniformly typed array instead.
///
/// @param module the module to transform
/// @returns success or failure
Result<SuccessType> DecomposeAccess(core::ir::Module& module,
const DecomposeAccessOptions& options);
} // namespace tint::core::ir::transform
#endif // SRC_TINT_LANG_CORE_IR_TRANSFORM_DECOMPOSE_ACCESS_H_
3rdparty/dawn/src/tint/lang/core/ir/transform/decompose_uniform_access.cc
Vendored
-624
View File
@@ -1,624 +0,0 @@
// Copyright 2024 The Dawn & Tint Authors
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are met:
//
// 1. Redistributions of source code must retain the above copyright notice, this
// list of conditions and the following disclaimer.
//
// 2. Redistributions in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// 3. Neither the name of the copyright holder nor the names of its
// contributors may be used to endorse or promote products derived from
// this software without specific prior written permission.
//
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
#include "src/tint/lang/core/ir/transform/decompose_uniform_access.h"
#include <utility>
#include "src/tint/lang/core/ir/builder.h"
#include "src/tint/lang/core/ir/validator.h"
namespace tint::core::ir::transform {
namespace {
using namespace tint::core::fluent_types; // NOLINT
using namespace tint::core::number_suffixes; // NOLINT
/// PIMPL state for the transform.
///
/// Each uniform buffer variable is rewritten so its store type is an array of vec4u
/// elements. Accesses to original store type contents are rewritten to pick out the
/// right set of vec4u elements and then use bitcasts as needed to construct the originally
/// loaded value.
struct State {
/// The IR module.
core::ir::Module& ir;
/// The IR builder.
core::ir::Builder b{ir};
/// The type manager.
core::type::Manager& ty{ir.Types()};
using VarTypePair = std::pair<core::ir::Var*, const core::type::Type*>;
/// Maps a struct to the load function
Hashmap<VarTypePair, core::ir::Function*, 2> var_and_type_to_load_fn_{};
/// Process the module.
void Process() {
Vector<core::ir::Var*, 4> var_worklist;
for (auto* inst : *ir.root_block) {
// Allow this to run before or after PromoteInitializers by handling non-var root_block
// entries
auto* var = inst->As<core::ir::Var>();
if (!var) {
continue;
}
// DecomposeStorageAccess maybe have converte the var pointers into ByteAddressBuffer
// objects. Since they've been changed, then they're Storage buffers and we don't care
// about them here.
auto* var_ty = var->Result()->Type()->As<core::type::Pointer>();
if (!var_ty) {
continue;
}
// Only care about uniform address space variables.
if (var_ty->AddressSpace() != core::AddressSpace::kUniform) {
continue;
}
var_worklist.Push(var);
}
for (auto* var : var_worklist) {
auto* result = var->Result();
auto usage_worklist = result->UsagesSorted();
auto* var_ty = result->Type()->As<core::type::Pointer>();
while (!usage_worklist.IsEmpty()) {
auto usage = usage_worklist.Pop();
auto* inst = usage.instruction;
// Load instructions can be destroyed by the replacing access function
if (!inst->Alive()) {
continue;
}
OffsetData od{};
tint::Switch(
inst, //
[&](core::ir::LoadVectorElement* l) { LoadVectorElement(l, var, od); },
[&](core::ir::Load* l) { Load(l, var, od); },
[&](core::ir::Access* a) { Access(a, var, a->Object()->Type(), od); },
[&](core::ir::Let* let) {
// The `let` is, essentially, an alias for the `var` as it's assigned
// directly. Gather all the `let` usages into our worklist, and then replace
// the `let` with the `var` itself.
for (auto& use : let->Result()->UsagesSorted()) {
usage_worklist.Push(use);
}
let->Result()->ReplaceAllUsesWith(result);
let->Destroy();
},
TINT_ICE_ON_NO_MATCH);
}
// Swap the result type of the `var` to the new result type
auto array_length = (var_ty->StoreType()->Size() + 15) / 16;
result->SetType(ty.ptr(var_ty->AddressSpace(), ty.array(ty.vec4u(), array_length),
var_ty->Access()));
}
}
// OffsetData represents an unsigned integer expression.
// The value is the sum of a const part, held in `byte_offset`, and
// non-const parts in `byte_offset_expr`.
struct OffsetData {
uint32_t byte_offset = 0;
Vector<core::ir::Value*, 4> byte_offset_expr{};
};
// Note, must be called inside a builder insert block (Append, InsertBefore, etc)
core::ir::Value* OffsetToValue(OffsetData offset) {
core::ir::Value* val = nullptr;
// If the offset is zero, skip setting val. This way, we won't add `0 +` and create useless
// addition expressions, but if the offset is zero, and there are no expressions, make sure
// we return the 0 value.
if (offset.byte_offset != 0) {
val = b.Value(u32(offset.byte_offset));
} else if (offset.byte_offset_expr.IsEmpty()) {
return b.Value(0_u);
}
for (core::ir::Value* expr : offset.byte_offset_expr) {
if (!val) {
val = expr;
} else {
val = b.Add(val, expr)->Result();
}
}
return val;
}
// Converts a byte offset to an index into a vec4u array.
// After the transform runs the store type of the uniform buffer variable is an
// array of vec4u.
// Note, this must be called inside a builder insert block (Append, InsertBefore, etc)
core::ir::Value* OffsetValueToArrayIndex(core::ir::Value* val) {
if (auto* cnst = val->As<core::ir::Constant>()) {
auto v = cnst->Value()->ValueAs<uint32_t>();
return b.Value(u32(v / 16u));
}
return b.Divide(val, 16_u)->Result();
}
// Calculates the index of the vec4u element containing the byte at (byte_idx % 16).
// Assumes the upper bits of byte_idx have already been used to access the correct
// vec4u array element in the underlying variable.
core::ir::Value* CalculateVectorOffset(core::ir::Value* byte_idx) {
if (auto* byte_cnst = byte_idx->As<core::ir::Constant>()) {
return b.Value(u32((byte_cnst->Value()->ValueAs<uint32_t>() % 16u) / 4u));
}
// Note: Using bitwise-and and shift instead of modulo and divide here was necessary to
// avoid an FXC miscompile. See https://crbug.com/454366353.
return b.ShiftRight(b.And(byte_idx, 15_u), 2_u)->Result();
}
void Access(core::ir::Access* a,
core::ir::Var* var,
const core::type::Type* obj_ty,
OffsetData offset) {
// Note, because we recurse through the `access` helper, the object passed in isn't
// necessarily the originating `var` object, but maybe a partially resolved access chain
// object.
if (auto* view = obj_ty->As<core::type::MemoryView>()) {
obj_ty = view->StoreType();
}
auto update_offset = [&](core::ir::Value* idx_value, uint32_t size) {
tint::Switch(
idx_value,
[&](core::ir::Constant* cnst) {
uint32_t idx = cnst->Value()->ValueAs<uint32_t>();
offset.byte_offset += size * idx;
},
[&](core::ir::Value* val) {
b.InsertBefore(a, [&] {
offset.byte_offset_expr.Push(
b.Multiply(u32(size), b.InsertConvertIfNeeded(ty.u32(), val))
->Result());
});
});
};
for (auto* idx_value : a->Indices()) {
tint::Switch(
obj_ty,
[&](const core::type::Vector* v) {
update_offset(idx_value, v->Type()->Size());
obj_ty = v->Type();
},
[&](const core::type::Matrix* m) {
update_offset(idx_value, m->ColumnStride());
obj_ty = m->ColumnType();
},
[&](const core::type::Array* ary) {
update_offset(idx_value, ary->ImplicitStride());
obj_ty = ary->ElemType();
},
[&](const core::type::Struct* s) {
auto* cnst = idx_value->As<core::ir::Constant>();
// A struct index must be a constant
TINT_IR_ASSERT(ir, cnst);
uint32_t idx = cnst->Value()->ValueAs<uint32_t>();
auto* mem = s->Members()[idx];
obj_ty = mem->Type();
offset.byte_offset += mem->Offset();
},
TINT_ICE_ON_NO_MATCH);
}
auto usages = a->Result()->UsagesSorted();
while (!usages.IsEmpty()) {
auto usage = usages.Pop();
tint::Switch(
usage.instruction,
[&](core::ir::Let* let) {
// The `let` is essentially an alias to the `access`. So, add the `let`
// usages into the usage worklist, and replace the let with the access chain
// directly.
for (auto& u : let->Result()->UsagesSorted()) {
usages.Push(u);
}
let->Result()->ReplaceAllUsesWith(a->Result());
let->Destroy();
},
[&](core::ir::Access* sub_access) {
// Treat an access chain of the access chain as a continuation of the outer
// chain. Pass through the object we stopped at and the current byte_offset
// and then restart the access chain replacement for the new access chain.
Access(sub_access, var, obj_ty, offset);
},
[&](core::ir::Load* ld) {
a->Result()->RemoveUsage(usage);
Load(ld, var, offset);
},
[&](core::ir::LoadVectorElement* lve) {
a->Result()->RemoveUsage(usage);
LoadVectorElement(lve, var, offset);
},
TINT_ICE_ON_NO_MATCH);
}
a->Destroy();
}
void Load(core::ir::Load* ld, core::ir::Var* var, OffsetData offset) {
b.InsertBefore(ld, [&] {
auto* byte_idx = OffsetToValue(offset);
auto* result = MakeLoad(ld, var, ld->Result()->Type(), byte_idx);
ld->Result()->ReplaceAllUsesWith(result->Result());
});
ld->Destroy();
}
void LoadVectorElement(core::ir::LoadVectorElement* lve,
core::ir::Var* var,
OffsetData offset) {
b.InsertBefore(lve, [&] {
// Add the byte count from the start of the vector to the requested element to the
// current offset calculation
auto elem_byte_size = lve->Result()->Type()->DeepestElement()->Size();
if (auto* cnst = lve->Index()->As<core::ir::Constant>()) {
offset.byte_offset += (cnst->Value()->ValueAs<uint32_t>() * elem_byte_size);
} else {
offset.byte_offset_expr.Push(
b.Multiply(b.InsertConvertIfNeeded(ty.u32(), lve->Index()), u32(elem_byte_size))
->Result());
}
auto* byte_idx = OffsetToValue(offset);
auto* result = MakeLoad(lve, var, lve->Result()->Type(), byte_idx);
lve->Result()->ReplaceAllUsesWith(result->Result());
});
lve->Destroy();
}
// Creates the appropriate load instructions for the given result type.
core::ir::Instruction* MakeLoad(core::ir::Instruction* inst,
core::ir::Var* var,
const core::type::Type* result_ty,
core::ir::Value* byte_idx) {
if (result_ty->IsFloatScalar() || result_ty->IsIntegerScalar()) {
return MakeScalarLoad(var, result_ty, byte_idx);
}
return tint::Switch(
result_ty,
[&](const core::type::Struct* s) {
auto* fn = GetLoadFunctionFor(inst, var, s);
return b.Call(fn, byte_idx);
},
[&](const core::type::Matrix* m) {
auto* fn = GetLoadFunctionFor(inst, var, m);
return b.Call(fn, byte_idx);
},
[&](const core::type::Array* a) {
auto* fn = GetLoadFunctionFor(inst, var, a);
return b.Call(fn, byte_idx);
},
[&](const core::type::Vector* v) { return MakeVectorLoad(var, v, byte_idx); },
TINT_ICE_ON_NO_MATCH);
}
core::ir::Instruction* MakeScalarLoad(core::ir::Var* var,
const core::type::Type* result_ty,
core::ir::Value* byte_idx) {
auto* array_idx = OffsetValueToArrayIndex(byte_idx);
auto* access = b.Access(ty.ptr(uniform, ty.vec4u()), var, array_idx);
auto* vec_idx = CalculateVectorOffset(byte_idx);
core::ir::Instruction* load = b.LoadVectorElement(access, vec_idx);
if (result_ty->Is<core::type::F16>()) {
return ExtractScalarF16(load, byte_idx);
}
return b.Bitcast(result_ty, load);
}
core::ir::Access* ExtractScalarF16(core::ir::Instruction* load, core::ir::Value* byte_idx) {
// We will bitcast the u32 to a vec2h and then extract the element that we want.
core::ir::Value* element_index = nullptr;
if (auto* cnst = byte_idx->As<core::ir::Constant>()) {
if (cnst->Value()->ValueAs<uint32_t>() % 4 == 0) {
element_index = b.Constant(0_u);
} else {
element_index = b.Constant(1_u);
}
} else {
auto* false_ = b.Value(1_u);
auto* true_ = b.Value(0_u);
auto* cond = b.Equal(b.Modulo(byte_idx, 4_u), 0_u);
Vector<core::ir::Value*, 3> args{false_, true_, cond->Result()};
element_index = b.Call(ty.u32(), core::BuiltinFn::kSelect, args)->Result();
}
auto* bitcast = b.Bitcast<vec2<f16>>(load);
return b.Access<f16>(bitcast, element_index);
}
// When loading a vector we have to take the alignment into account to determine which part of
// the `uint4` to load. A `vec` of `u32`, `f32` or `i32` has an alignment requirement of
// a multiple of 8-bytes (`f16` is 4-bytes). So, this means we'll have memory like:
//
// Byte: | 0 | 1 | 2 | 3 | 4 | 5 | 6 | 7 | 8 | 9 | 10 | 11 | 12 | 13 | 14 | 15 | 16 |
// Value:| v1 | v2 | v3 | v4 |
// Scalar Index: 0 1 2 3
//
// Start with a byte address which is `offset + (column * columnStride)`, the array index is
// `byte_address / 16`. This gives us the `uint4` which contains our values. We can then
// calculate the vector offset as `(byte_address % 16) / 4`.
//
// * A 4-element row we access all 4-elements at the `array_idx`
// * A 3-element row we access `array_idx` at `.xyz` as it must be padded to a vec4.
// * A 2-element row, we have to decide if we want the `xy` or `zw` element. We have a minimum
// alignment of 8-bytes as per the WGSL spec. So if the `vector_idx != 2` is `0` then we
// access the `.xy` component, otherwise it is in the `.zw` component.
core::ir::Instruction* MakeVectorLoad(core::ir::Var* var,
const core::type::Vector* result_ty,
core::ir::Value* byte_idx) {
auto* array_idx = OffsetValueToArrayIndex(byte_idx);
auto* access = b.Access(ty.ptr(uniform, ty.vec4u()), var, array_idx);
if (result_ty->DeepestElement()->Is<core::type::F16>()) {
return MakeVectorLoadF16(access, result_ty, byte_idx);
}
TINT_IR_ASSERT(ir, result_ty->DeepestElement()->Size() == 4);
core::ir::Instruction* load = nullptr;
if (result_ty->Width() == 4) {
load = b.Load(access);
} else if (result_ty->Width() == 3) {
load = b.Swizzle(ty.vec3u(), b.Load(access), {0, 1, 2});
} else if (result_ty->Width() == 2) {
auto* vec_idx = CalculateVectorOffset(byte_idx);
if (auto* cnst = vec_idx->As<core::ir::Constant>()) {
if (cnst->Value()->ValueAs<uint32_t>() == 2u) {
load = b.Swizzle(ty.vec2u(), b.Load(access), {2, 3});
} else {
load = b.Swizzle(ty.vec2u(), b.Load(access), {0, 1});
}
} else {
auto* ubo = b.Load(access);
// if vec_idx == 2 -> zw
auto* sw_lhs = b.Swizzle(ty.vec2u(), ubo, {2, 3});
// else -> xy
auto* sw_rhs = b.Swizzle(ty.vec2u(), ubo, {0, 1});
auto* cond = b.Equal(vec_idx, 2_u);
Vector<core::ir::Value*, 3> args{sw_rhs->Result(), sw_lhs->Result(),
cond->Result()};
load = b.Call(ty.vec2u(), core::BuiltinFn::kSelect, args);
}
} else {
TINT_IR_UNREACHABLE(ir);
}
return b.Bitcast(result_ty, load);
}
// Returns the instruction for getting a vector-of-f16 value of type `result_ty`
// out of the pointer-to-vec4u value `access`. Get the components starting
// at byte offset (byte_idx % 16).
// A `vec4` or `vec3` of `f16` has 8-byte alignment, while a `vec2` of `f16` has 4-byte
// alignment. So, this means we'll have memory like:
//
// Byte: | 0 | 1 | 2 | 3 | 4 | 5 | 6 | 7 | 8 | 9 | 10 | 11 | 12 | 13 | 14 | 15 |
// Scalar Index: 0 | 1 | 2 | 3 |
// vec4<u32>:| x | y | z | w |
// vec4<f16>:| x | y | z | w | x | y | z | w |
// vec3<f16>:| x | y | z | | x | y | z | |
// vec2<f16>:| x | y | x | y | x | y | x | y |
core::ir::Instruction* MakeVectorLoadF16(core::ir::Access* access,
const core::type::Vector* result_ty,
core::ir::Value* byte_idx) {
// Vec4 ends up being the same as a bitcast of vec2<u32> to a vec4<f16>.
// A vec3 will be stored as a vec4, so we can bitcast as if we're a vec4
// and swizzle out the last element.
if (result_ty->Width() == 3 || result_ty->Width() == 4) {
core::ir::Instruction* load = nullptr;
auto* vec_idx = CalculateVectorOffset(byte_idx); // 0 or 2
if (auto* cnst = vec_idx->As<core::ir::Constant>()) {
if (cnst->Value()->ValueAs<uint32_t>() == 2u) {
load = b.Swizzle(ty.vec2u(), b.Load(access), {2, 3});
} else {
load = b.Swizzle(ty.vec2u(), b.Load(access), {0, 1});
}
} else {
auto* ubo = b.Load(access);
// if vec_idx == 2 -> zw
auto* sw_lhs = b.Swizzle(ty.vec2u(), ubo, {2, 3});
// else -> xy
auto* sw_rhs = b.Swizzle(ty.vec2u(), ubo, {0, 1});
auto* cond = b.Equal(vec_idx, 2_u);
auto args = Vector{sw_rhs->Result(), sw_lhs->Result(), cond->Result()};
load = b.Call(ty.vec2u(), core::BuiltinFn::kSelect, std::move(args));
}
if (result_ty->Width() == 3) {
auto* bc = b.Bitcast(ty.vec4(result_ty->Type()), load);
return b.Swizzle(result_ty, bc, {0, 1, 2});
}
return b.Bitcast(result_ty, load);
}
// Vec2 ends up being the same as a bitcast u32 to vec2<f16>
if (result_ty->Width() == 2) {
core::ir::Instruction* load = nullptr;
auto* vec_idx = CalculateVectorOffset(byte_idx); // 0, 1, 2, or 3
if (auto* cnst = vec_idx->As<core::ir::Constant>()) {
const auto vec_idx_val = cnst->Value()->ValueAs<uint32_t>();
load = b.Swizzle(ty.u32(), b.Load(access), {vec_idx_val});
} else {
load = b.Access(ty.u32(), b.Load(access), vec_idx);
}
return b.Bitcast(result_ty, load);
}
TINT_IR_UNREACHABLE(ir);
}
// Creates a load function for the given `var` and `matrix` combination. Essentially creates
// a function similar to:
//
// fn custom_load_M(offset: u32) {
// const uint scalar_offset = ((offset + 0u)) / 4;
// const uint scalar_offset_1 = ((offset + (1 * ColumnStride))) / 4;
// const uint scalar_offset_2 = ((offset + (2 * ColumnStride))) / 4;
// const uint scalar_offset_3 = ((offset + (3 * ColumnStride)))) / 4;
// return float4x4(
// asfloat(v[scalar_offset / 4]),
// asfloat(v[scalar_offset_1 / 4]),
// asfloat(v[scalar_offset_2 / 4]),
// asfloat(v[scalar_offset_3 / 4])
// );
// }
core::ir::Function* GetLoadFunctionFor(core::ir::Instruction* inst,
core::ir::Var* var,
const core::type::Matrix* mat) {
return var_and_type_to_load_fn_.GetOrAdd(VarTypePair{var, mat}, [&] {
auto* start_byte_offset = b.FunctionParam("start_byte_offset", ty.u32());
auto* fn = b.Function(mat);
fn->SetParams({start_byte_offset});
b.Append(fn->Block(), [&] {
Vector<core::ir::Value*, 4> values;
for (size_t i = 0; i < mat->Columns(); ++i) {
uint32_t stride = static_cast<uint32_t>(i * mat->ColumnStride());
OffsetData od{stride, {start_byte_offset}};
auto* byte_idx = OffsetToValue(od);
values.Push(MakeLoad(inst, var, mat->ColumnType(), byte_idx)->Result());
}
b.Return(fn, b.Construct(mat, values));
});
return fn;
});
}
// Creates a load function for the given `var` and `array` combination. Essentially creates
// a function similar to:
//
// fn custom_load_A(offset: u32) {
// A a = A();
// u32 i = 0;
// loop {
// if (i >= A length) {
// break;
// }
// offset = (offset + (i * A->ImplicitStride())) / 16
// a[i] = cast(v[offset].xyz)
// i = i + 1;
// }
// return a;
// }
core::ir::Function* GetLoadFunctionFor(core::ir::Instruction* inst,
core::ir::Var* var,
const core::type::Array* arr) {
return var_and_type_to_load_fn_.GetOrAdd(VarTypePair{var, arr}, [&] {
auto* start_byte_offset = b.FunctionParam("start_byte_offset", ty.u32());
auto* fn = b.Function(arr);
fn->SetParams({start_byte_offset});
b.Append(fn->Block(), [&] {
auto* result_arr = b.Var<function>("a", b.Zero(arr));
auto* count = arr->Count()->As<core::type::ConstantArrayCount>();
TINT_IR_ASSERT(ir, count);
b.LoopRange(0_u, u32(count->value), 1_u, [&](core::ir::Value* idx) {
auto* stride = b.Multiply(idx, u32(arr->ImplicitStride()))->Result();
OffsetData od{0, {start_byte_offset, stride}};
auto* byte_idx = OffsetToValue(od);
auto* access = b.Access(ty.ptr<function>(arr->ElemType()), result_arr, idx);
b.Store(access, MakeLoad(inst, var, arr->ElemType(), byte_idx));
});
b.Return(fn, b.Load(result_arr));
});
return fn;
});
}
// Creates a load function for the given `var` and `struct` combination. Essentially creates
// a function similar to:
//
// fn custom_load_S(start_offset: u32) {
// let a = load object at (start_offset + member_offset)
// let b = load object at (start_offset + member 1 offset);
// ...
// return S(a, b, ..., z);
// }
core::ir::Function* GetLoadFunctionFor(core::ir::Instruction* inst,
core::ir::Var* var,
const core::type::Struct* s) {
return var_and_type_to_load_fn_.GetOrAdd(VarTypePair{var, s}, [&] {
auto* start_byte_offset = b.FunctionParam("start_byte_offset", ty.u32());
auto* fn = b.Function(s);
fn->SetParams({start_byte_offset});
b.Append(fn->Block(), [&] {
Vector<core::ir::Value*, 4> values;
for (const auto* mem : s->Members()) {
uint32_t stride = static_cast<uint32_t>(mem->Offset());
OffsetData od{stride, {start_byte_offset}};
auto* byte_idx = OffsetToValue(od);
values.Push(MakeLoad(inst, var, mem->Type(), byte_idx)->Result());
}
b.Return(fn, b.Construct(s, values));
});
return fn;
});
}
};
} // namespace
Result<SuccessType> DecomposeUniformAccess(core::ir::Module& ir) {
TINT_CHECK_RESULT(
ValidateAndDumpIfNeeded(ir, "core.DecomposeUniformAccess",
core::ir::Capabilities{
core::ir::Capability::kAllow8BitIntegers,
core::ir::Capability::kAllowClipDistancesOnF32ScalarAndVector,
core::ir::Capability::kAllowDuplicateBindings,
core::ir::Capability::kAllowNonCoreTypes,
}));
State{ir}.Process();
return Success;
}
} // namespace tint::core::ir::transform
3rdparty/dawn/src/tint/lang/core/ir/transform/demote_to_helper.cc
Vendored
+1 -1
View File
@@ -215,7 +215,7 @@ struct State {
Result<SuccessType> DemoteToHelper(Module& ir) {
TINT_CHECK_RESULT(
ValidateAndDumpIfNeeded(ir, "core.DemoteToHelper", kDemoteToHelperCapabilities));
ValidateBeforeIfNeeded(ir, kDemoteToHelperCapabilities, "core.DemoteToHelper"));
State{ir}.Process();
3rdparty/dawn/src/tint/lang/core/ir/transform/direct_variable_access.cc
Vendored
+2 -2
View File
@@ -672,8 +672,8 @@ struct State {
} // namespace
Result<SuccessType> DirectVariableAccess(Module& ir, const DirectVariableAccessOptions& options) {
TINT_CHECK_RESULT(ValidateAndDumpIfNeeded(ir, "core.DirectVariableAccess",
kDirectVariableAccessCapabilities));
TINT_CHECK_RESULT(
ValidateBeforeIfNeeded(ir, kDirectVariableAccessCapabilities, "core.DirectVariableAccess"));
State{ir, options}.Process();
3rdparty/dawn/src/tint/lang/core/ir/transform/multiplanar_external_texture.cc
Vendored
+4 -4
View File
@@ -227,8 +227,8 @@ struct State {
if (call->Func() == core::BuiltinFn::kTextureDimensions) {
// Use params.apparentSize + vec2u(1, 1) instead of the textureDimensions.
b.InsertBefore(call, [&] {
auto* apparent_size = b.Access<vec2<u32>>(params, 12_u);
auto* vec2u_1_1 = b.Splat<vec2<u32>>(1_u);
auto* apparent_size = b.Access<vec2u>(params, 12_u);
auto* vec2u_1_1 = b.Splat<vec2u>(1_u);
auto* dimensions = b.Add(apparent_size, vec2u_1_1);
dimensions->SetResult(call->DetachResult());
});
@@ -603,8 +603,8 @@ struct State {
Result<SuccessType> MultiplanarExternalTexture(
Module& ir,
const tint::transform::multiplanar::BindingsMap& multiplanar_map) {
TINT_CHECK_RESULT(ValidateAndDumpIfNeeded(ir, "core.MultiplanarExternalTexture",
kMultiplanarExternalTextureCapabilities));
TINT_CHECK_RESULT(ValidateBeforeIfNeeded(ir, kMultiplanarExternalTextureCapabilities,
"core.MultiplanarExternalTexture"));
return State{multiplanar_map, ir}.Process();
}
3rdparty/dawn/src/tint/lang/core/ir/transform/prepare_immediate_data.cc
Vendored
+7 -6
View File
@@ -149,12 +149,13 @@ struct State {
Result<ImmediateDataLayout> PrepareImmediateData(Module& ir,
const PrepareImmediateDataConfig& config) {
TINT_CHECK_RESULT(ValidateAndDumpIfNeeded(ir, "core.PrepareImmediateData",
core::ir::Capabilities{
core::ir::Capability::kAllowDuplicateBindings,
core::ir::Capability::kAllow8BitIntegers,
core::ir::Capability::kAllowNonCoreTypes,
}));
TINT_CHECK_RESULT(ValidateBeforeIfNeeded(ir,
core::ir::Capabilities{
core::ir::Capability::kAllowDuplicateBindings,
core::ir::Capability::kAllow8BitIntegers,
core::ir::Capability::kAllowNonCoreTypes,
},
"core.PrepareImmediateData"));
return State{config, ir}.Run();
}
3rdparty/dawn/src/tint/lang/core/ir/transform/preserve_padding.cc
Vendored
+1 -1
View File
@@ -170,7 +170,7 @@ struct State {
Result<SuccessType> PreservePadding(Module& ir) {
TINT_CHECK_RESULT(
ValidateAndDumpIfNeeded(ir, "core.PreservePadding", kPreservePaddingCapabilities));
ValidateBeforeIfNeeded(ir, kPreservePaddingCapabilities, "core.PreservePadding"));
State{ir}.Process();
3rdparty/dawn/src/tint/lang/core/ir/transform/prevent_infinite_loops.cc
Vendored
+5 -5
View File
@@ -78,7 +78,7 @@ struct State {
// idx.y -= u32(idx.x == UINT32_MAX);
// Declare a new index variable at the top of the loop initializer.
auto* idx = b.Var<function>("tint_loop_idx", b.Splat<vec2<u32>>(u32::Highest()));
auto* idx = b.Var<function>("tint_loop_idx", b.Splat<vec2u>(u32::Highest()));
if (loop->Initializer()->IsEmpty()) {
loop->Initializer()->Append(b.NextIteration(loop));
}
@@ -86,8 +86,8 @@ struct State {
// Insert the new exit condition at the top of the loop body.
b.InsertBefore(loop->Body()->Front(), [&] {
auto* ifelse = b.If(b.Call<bool>(
BuiltinFn::kAll, b.Equal(b.Load(idx)->Result(), b.Splat<vec2<u32>>(0_u))));
auto* ifelse = b.If(
b.Call<bool>(BuiltinFn::kAll, b.Equal(b.Load(idx)->Result(), b.Splat<vec2u>(0_u))));
b.Append(ifelse->True(), [&] {
// If the loop produces result values, just use `undef` as this exit condition
// should never actually be hit.
@@ -116,8 +116,8 @@ struct State {
} // namespace
Result<SuccessType> PreventInfiniteLoops(Module& ir) {
TINT_CHECK_RESULT(ValidateAndDumpIfNeeded(ir, "core.PreventInfiniteLoops",
kPreventInfiniteLoopsCapabilities));
TINT_CHECK_RESULT(
ValidateBeforeIfNeeded(ir, kPreventInfiniteLoopsCapabilities, "core.PreventInfiniteLoops"));
State{ir}.Process();
3rdparty/dawn/src/tint/lang/core/ir/transform/remove_continue_in_switch.cc
Vendored
+9 -8
View File
@@ -121,14 +121,15 @@ struct State {
Result<SuccessType> RemoveContinueInSwitch(Module& ir) {
TINT_CHECK_RESULT(
ValidateAndDumpIfNeeded(ir, "core.RemoveContinueInSwitch",
core::ir::Capabilities{
core::ir::Capability::kAllowVectorElementPointer,
core::ir::Capability::kAllowHandleVarsWithoutBindings,
core::ir::Capability::kAllowClipDistancesOnF32ScalarAndVector,
core::ir::Capability::kAllowDuplicateBindings,
core::ir::Capability::kAllowNonCoreTypes,
}));
ValidateBeforeIfNeeded(ir,
core::ir::Capabilities{
core::ir::Capability::kAllowVectorElementPointer,
core::ir::Capability::kAllowHandleVarsWithoutBindings,
core::ir::Capability::kAllowClipDistancesOnF32ScalarAndVector,
core::ir::Capability::kAllowDuplicateBindings,
core::ir::Capability::kAllowNonCoreTypes,
},
"core.RemoveContinueInSwitch"));
State{ir}.Process();
3rdparty/dawn/src/tint/lang/core/ir/transform/remove_terminator_args.cc
Vendored
+2 -2
View File
@@ -163,8 +163,8 @@ struct State {
} // namespace
Result<SuccessType> RemoveTerminatorArgs(Module& ir) {
TINT_CHECK_RESULT(ValidateAndDumpIfNeeded(ir, "core.RemoveTerminatorArgs",
kRemoveTerminatorArgsCapabilities));
TINT_CHECK_RESULT(
ValidateBeforeIfNeeded(ir, kRemoveTerminatorArgsCapabilities, "core.RemoveTerminatorArgs"));
State{ir}.Process();
3rdparty/dawn/src/tint/lang/core/ir/transform/remove_uniform_vector_component_loads.cc
Vendored
+7 -6
View File
@@ -78,12 +78,13 @@ struct State {
} // namespace
Result<SuccessType> RemoveUniformVectorComponentLoads(core::ir::Module& ir) {
TINT_CHECK_RESULT(ValidateAndDumpIfNeeded(ir, "core.RemoveUniformVectorComponentLoads",
core::ir::Capabilities{
core::ir::Capability::kAllowDuplicateBindings,
core::ir::Capability::kAllowNonCoreTypes,
core::ir::Capability::kAllow8BitIntegers,
}));
TINT_CHECK_RESULT(ValidateBeforeIfNeeded(ir,
core::ir::Capabilities{
core::ir::Capability::kAllowDuplicateBindings,
core::ir::Capability::kAllowNonCoreTypes,
core::ir::Capability::kAllow8BitIntegers,
},
"core.RemoveUniformVectorComponentLoads"));
State{ir}.Process();
3rdparty/dawn/src/tint/lang/core/ir/transform/rename_conflicts.cc
Vendored
+2 -1
View File
@@ -27,6 +27,7 @@
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
#include "src/tint/lang/core/ir/transform/rename_conflicts.h"
#include "src/tint/lang/core/ir/construct.h"
#include "src/tint/lang/core/ir/control_instruction.h"
#include "src/tint/lang/core/ir/convert.h"
@@ -301,7 +302,7 @@ struct State {
Result<SuccessType> RenameConflicts(core::ir::Module& ir) {
TINT_CHECK_RESULT(
ValidateAndDumpIfNeeded(ir, "core.RenameConflicts", kRenameConflictsCapabilities));
ValidateBeforeIfNeeded(ir, kRenameConflictsCapabilities, "core.RenameConflicts"));
State{ir}.Process();
3rdparty/dawn/src/tint/lang/core/ir/transform/resource_table.cc
Vendored
+24 -5
View File
@@ -133,9 +133,28 @@ struct State {
b.Append(has_check->True(), [&] {
auto* type_val = b.Access(ty.ptr<storage, u32, read>(), storage_buffer, 1_u, idx);
auto* v = b.Load(type_val);
auto* eq = b.Equal(v, u32(static_cast<uint32_t>(core::type::TypeToResourceType(type))));
ResourceType resource_ty = core::type::TypeToResourceType(type);
core::ir::Value* eq = nullptr;
std::vector<ResourceType> conv = core::type::ConvertsFrom(type);
if (!conv.empty()) {
auto* conv_ty = ty.vec(ty.u32(), static_cast<uint32_t>(conv.size()) + 1);
auto* lhs = b.Construct(conv_ty, v);
Vector<Value*, 4> vals;
vals.Push(b.Value(u32(resource_ty)));
for (auto& r : conv) {
vals.Push(b.Value(u32(r)));
}
auto* rhs = b.Construct(conv_ty, vals);
auto* cmp = b.Equal(lhs, rhs);
eq = b.Call(ty.bool_(), core::BuiltinFn::kAny, cmp)->Result();
} else {
eq = b.Equal(v, u32(resource_ty))->Result();
}
b.ExitIf(has_check, eq);
});
@@ -181,14 +200,14 @@ struct State {
}
core::ir::Var* InjectStorageBuffer(const BindingPoint& bp) {
auto* str = ty.Struct(ir.symbols.New("tint_resource_table_buffer"),
auto* str = ty.Struct(ir.symbols.New("tint_resource_table_metadata_struct"),
Vector<core::type::Manager::StructMemberDesc, 2>{
{ir.symbols.New("array_length"), ty.u32()},
{ir.symbols.New("bindings"), ty.array<u32>()},
});
auto* sb_ty = ty.ptr(storage, str, read);
auto* sb = b.Var(sb_ty);
auto* sb = b.Var("tint_resource_table_metadata", sb_ty);
sb->SetBindingPoint(bp.group, bp.binding);
return sb;
@@ -202,7 +221,7 @@ ResourceTableHelper::~ResourceTableHelper() = default;
Result<SuccessType> ResourceTable(core::ir::Module& ir,
const ResourceTableConfig& config,
ResourceTableHelper* helper) {
TINT_CHECK_RESULT(ValidateAndDumpIfNeeded(ir, "core.ResourceTable"));
TINT_CHECK_RESULT(ValidateBeforeIfNeeded(ir, {}, "core.ResourceTable"));
State{config, ir, helper}.Process();
3rdparty/dawn/src/tint/lang/core/ir/transform/resource_table_helper.cc
Vendored
+2 -5
View File
@@ -32,17 +32,14 @@
#include <vector>
#include "src/tint/lang/core/ir/core_builtin_call.h"
#include "src/tint/lang/core/ir/load.h"
#include "src/tint/lang/core/ir/var.h"
#include "src/tint/lang/core/type/resource_type.h"
#include "src/tint/utils/rtti/switch.h"
namespace tint::core::ir::transform {
std::optional<ResourceTableConfig> GenerateResourceTableConfig(Module& mod) {
ResourceTableConfig cfg{
.resource_table_binding = BindingPoint{.group = 42, .binding = 41},
.storage_buffer_binding = BindingPoint{.group = 42, .binding = 42},
.resource_table_binding = BindingPoint{.group = 43, .binding = 51},
.storage_buffer_binding = BindingPoint{.group = 42, .binding = 52},
.default_binding_type_order = {},
};
3rdparty/dawn/src/tint/lang/core/ir/transform/robustness.cc
Vendored
+1 -1
View File
@@ -568,7 +568,7 @@ struct State {
} // namespace
Result<SuccessType> Robustness(Module& ir, const RobustnessConfig& config) {
TINT_CHECK_RESULT(ValidateAndDumpIfNeeded(ir, "core.Robustness", kRobustnessCapabilities));
TINT_CHECK_RESULT(ValidateBeforeIfNeeded(ir, kRobustnessCapabilities, "core.Robustness"));
State{config, ir}.Process();
3rdparty/dawn/src/tint/lang/core/ir/transform/robustness.h
Vendored
+1 -1
View File
@@ -29,9 +29,9 @@
#define SRC_TINT_LANG_CORE_IR_TRANSFORM_ROBUSTNESS_H_
#include <unordered_set>
#include "src/tint/lang/core/ir/validator.h"
#include "src/tint/api/common/binding_point.h"
#include "src/tint/lang/core/ir/validator.h"
#include "src/tint/utils/reflection.h"
#include "src/tint/utils/result.h"
3rdparty/dawn/src/tint/lang/core/ir/transform/shader_io.cc
Vendored
+60
View File
@@ -257,6 +257,66 @@ struct State {
} // namespace
core::ir::Value* ShaderIOBackendState::PolyfillWorkgroupIndex(Builder& builder,
uint32_t workgroup_id_index,
uint32_t num_workgroups_index) {
if (tint_workgroup_index != nullptr) {
return tint_workgroup_index;
}
// workgroup_index = workgroup_id.x +
// (workgroup_id.y * num_workgroups.x) +
// (workgroup_id.z * num_workgroups.x * num_workgroups.y)
auto* workgroup_id = GetInput(builder, workgroup_id_index);
auto* num_workgroups = GetInput(builder, num_workgroups_index);
auto* num_workgroups_x = builder.Access(ty.u32(), num_workgroups, 0_u);
auto* num_workgroups_y = builder.Access(ty.u32(), num_workgroups, 1_u);
auto* z_part = builder.Multiply(num_workgroups_x, num_workgroups_y)->Result();
z_part = builder.Multiply(builder.Access(ty.u32(), workgroup_id, 2_u), z_part)->Result();
auto* y_part =
builder.Multiply(builder.Access(ty.u32(), workgroup_id, 1_u), num_workgroups_x)->Result();
auto* init = builder.Add(builder.Access(ty.u32(), workgroup_id, 0_u), y_part)->Result();
init = builder.Add(init, z_part)->Result();
tint_workgroup_index = init;
return tint_workgroup_index;
}
core::ir::Value* ShaderIOBackendState::PolyfillGlobalInvocationIndex(
Builder& builder,
uint32_t global_invocation_id_index,
uint32_t num_workgroups_index) {
if (tint_global_invocation_index) {
return tint_global_invocation_index;
}
// global_invocation_index =
// global_invocation_id.x +
// (global_invocation_id.y * num_workgroups.x * workgroup_size.x) +
// (global_invocation_id.z * num_workgroups.x * workgroup_size.x * num_workgroups.y *
// workgroup_size.y)
auto* num_workgroups = GetInput(builder, num_workgroups_index);
auto* global_id = GetInput(builder, global_invocation_id_index);
auto* global_id_x = builder.Access(ty.u32(), global_id, 0_u);
auto* global_id_y = builder.Access(ty.u32(), global_id, 1_u);
auto* global_id_z = builder.Access(ty.u32(), global_id, 2_u);
auto* num_workgroups_x = builder.Access(ty.u32(), num_workgroups, 0_u);
auto* num_workgroups_y = builder.Access(ty.u32(), num_workgroups, 1_u);
auto* x_size = builder.Multiply(num_workgroups_x, u32(workgroup_size->at(0)));
auto* y_size = builder.Multiply(num_workgroups_y, u32(workgroup_size->at(1)));
auto* z_part = builder.Multiply(x_size, y_size);
z_part = builder.Multiply(global_id_z, z_part);
auto* y_part = builder.Multiply(global_id_y, x_size);
auto* value = builder.Add(global_id_x, y_part);
value = builder.Add(value, z_part);
tint_global_invocation_index = value->Result();
return tint_global_invocation_index;
}
void RunShaderIOBase(Module& module, std::function<MakeBackendStateFunc> make_backend_state) {
State{make_backend_state, module}.Process();
}
3rdparty/dawn/src/tint/lang/core/ir/transform/shader_io.h
Vendored
+28
View File
@@ -103,6 +103,31 @@ struct ShaderIOBackendState {
/// @returns true if a vertex point size builtin should be added
virtual bool NeedsVertexPointSize() const { return false; }
/// Creates the polyfilled workgroup_index builtin value.
/// Each backend is responsible for tracking the indices of workgroup_id and num_workgroups
/// builtin values (and ensuring they are available in the module).
/// @param builder The IR builder
/// @param workgroup_id_index The index for GetInputs of the workgroup_id builtin value
/// @param num_workgroups_index The index for GetInputs of the num_workgroups builtin value
core::ir::Value* PolyfillWorkgroupIndex(Builder& builder,
uint32_t workgroup_id_index,
uint32_t num_workgroups_index);
/// Creates the polyfilled global_invocation_index builtin value.
/// Each backend is responsible for tracking the indices of global_invocation_id and
/// num_workgroups builtin values (and ensuring they are available in the module).
/// @param builder The IR builder
/// @param global_invocation_id_index The index for GetInputs of the global_invocation_id
/// builtin value
/// @param num_workgroups_index The index for GetInputs of the num_workgroups builtin value
core::ir::Value* PolyfillGlobalInvocationIndex(Builder& builder,
uint32_t global_invocation_id_index,
uint32_t num_workgroups_index);
/// The workgroup size of the entry point.
/// Backends are responsible for caching this value.
std::optional<std::array<uint32_t, 3>> workgroup_size = std::nullopt;
protected:
/// The IR module.
Module& ir;
@@ -121,6 +146,9 @@ struct ShaderIOBackendState {
/// The list of shader outputs.
Vector<core::type::Manager::StructMemberDesc, 4> outputs;
core::ir::Value* tint_workgroup_index = nullptr;
core::ir::Value* tint_global_invocation_index = nullptr;
};
/// The signature for a function that creates a backend state object.
3rdparty/dawn/src/tint/lang/core/ir/transform/signed_integer_polyfill.cc
Vendored
+16 -15
View File
@@ -146,21 +146,22 @@ struct State {
Result<SuccessType> SignedIntegerPolyfill(core::ir::Module& ir,
const SignedIntegerPolyfillConfig& cfg) {
TINT_CHECK_RESULT(
ValidateAndDumpIfNeeded(ir, "ir.SignedIntegerPolyfill",
core::ir::Capabilities{
core::ir::Capability::kAllowDuplicateBindings,
core::ir::Capability::kAllow8BitIntegers,
core::ir::Capability::kAllow64BitIntegers,
core::ir::Capability::kAllowPointSizeBuiltin,
core::ir::Capability::kAllowVectorElementPointer,
core::ir::Capability::kAllowHandleVarsWithoutBindings,
core::ir::Capability::kAllowClipDistancesOnF32ScalarAndVector,
core::ir::Capability::kAllowAnyLetType,
core::ir::Capability::kMslAllowEntryPointInterface,
core::ir::Capability::kAllowModuleScopeLets,
core::ir::Capability::kAllowAnyInputAttachmentIndexType,
core::ir::Capability::kAllowNonCoreTypes,
}));
ValidateBeforeIfNeeded(ir,
core::ir::Capabilities{
core::ir::Capability::kAllowDuplicateBindings,
core::ir::Capability::kAllow8BitIntegers,
core::ir::Capability::kAllow64BitIntegers,
core::ir::Capability::kAllowPointSizeBuiltin,
core::ir::Capability::kAllowVectorElementPointer,
core::ir::Capability::kAllowHandleVarsWithoutBindings,
core::ir::Capability::kAllowClipDistancesOnF32ScalarAndVector,
core::ir::Capability::kAllowAnyLetType,
core::ir::Capability::kMslAllowEntryPointInterface,
core::ir::Capability::kAllowModuleScopeLets,
core::ir::Capability::kAllowAnyInputAttachmentIndexType,
core::ir::Capability::kAllowNonCoreTypes,
},
"ir.SignedIntegerPolyfill"));
State{ir, cfg}.Process();
3rdparty/dawn/src/tint/lang/core/ir/transform/single_entry_point.cc
Vendored
+1 -1
View File
@@ -106,7 +106,7 @@ Result<SuccessType> Run(ir::Module& ir, std::string_view entry_point_name) {
Result<SuccessType> SingleEntryPoint(Module& ir, std::string_view entry_point_name) {
TINT_CHECK_RESULT(
ValidateAndDumpIfNeeded(ir, "core.SingleEntryPoint", kSingleEntryPointCapabilities));
ValidateBeforeIfNeeded(ir, kSingleEntryPointCapabilities, "core.SingleEntryPoint"));
return Run(ir, entry_point_name);
}
3rdparty/dawn/src/tint/lang/core/ir/transform/std140.cc
Vendored
+1 -1
View File
@@ -448,7 +448,7 @@ struct State {
} // namespace
Result<SuccessType> Std140(Module& ir) {
TINT_CHECK_RESULT(ValidateAndDumpIfNeeded(ir, "core.Std140", kStd140Capabilities));
TINT_CHECK_RESULT(ValidateBeforeIfNeeded(ir, kStd140Capabilities, "core.Std140"));
State{ir}.Process();
3rdparty/dawn/src/tint/lang/core/ir/transform/substitute_overrides.cc
Vendored
+20 -17
View File
@@ -29,11 +29,9 @@
#include <cstdint>
#include <functional>
#include <limits>
#include <utility>
#include "src/tint/lang/core/binary_op.h"
#include "src/tint/lang/core/fluent_types.h"
#include "src/tint/lang/core/ir/binary.h"
#include "src/tint/lang/core/ir/builder.h"
#include "src/tint/lang/core/ir/const_param_validator.h"
#include "src/tint/lang/core/ir/constexpr_if.h"
@@ -49,9 +47,6 @@
#include "src/tint/lang/core/ir/value.h"
#include "src/utils/numeric.h"
using namespace tint::core::fluent_types; // NOLINT
using namespace tint::core::number_suffixes; // NOLINT
namespace tint::core::ir::transform {
namespace {
@@ -69,6 +64,13 @@ struct State {
/// The type manager.
core::type::Manager& ty{ir.Types()};
diag::Diagnostic MakeError(const Source& src) {
diag::Diagnostic error{};
error.severity = diag::Severity::Error;
error.source = src;
return error;
}
/// Process the module.
diag::Result<SuccessType> Process() {
Vector<Instruction*, 8> to_remove;
@@ -122,9 +124,7 @@ struct State {
TINT_ICE_ON_NO_MATCH);
if (!substitution_representation_valid) {
diag::Diagnostic error{};
error.severity = diag::Severity::Error;
error.source = ir.SourceOf(override);
diag::Diagnostic error = MakeError(ir.SourceOf(override));
error << "Pipeline overridable constant " << iter->first.value
<< " with value (" << iter->second
<< ") is not representable in type ("
@@ -138,9 +138,7 @@ struct State {
}
if (override->Initializer() == nullptr) {
diag::Diagnostic error{};
error.severity = diag::Severity::Error;
error.source = ir.SourceOf(override);
diag::Diagnostic error = MakeError(ir.SourceOf(override));
error << "Initializer not provided for override, and override not overridden.";
return diag::Failure(error);
}
@@ -203,17 +201,22 @@ struct State {
// not check constant evaluation access against zero size.
int64_t cnt_size_check = new_value->Value()->ValueAs<AInt>();
if (cnt_size_check < 1) {
diag::Diagnostic error{};
error.severity = diag::Severity::Error;
error.source = ir.SourceOf(cnt->value);
diag::Diagnostic error = MakeError(ir.SourceOf(cnt->value));
error << "array count (" << cnt_size_check << ") must be greater than 0";
return diag::Failure(error);
}
uint32_t num_elements = new_value->Value()->ValueAs<uint32_t>();
uint64_t new_ary_size = uint64_t{num_elements} * old_ty->ImplicitStride();
if (new_ary_size > std::numeric_limits<uint32_t>::max()) {
diag::Diagnostic error = MakeError(ir.SourceOf(cnt->value));
error << "array size (" << new_ary_size << ") is too large";
return diag::Failure(error);
}
auto* new_cnt = ty.Get<core::type::ConstantArrayCount>(num_elements);
auto* new_ty = ty.Get<core::type::Array>(old_ty->ElemType(), new_cnt,
num_elements * old_ty->ImplicitStride());
static_cast<uint32_t>(new_ary_size));
auto* new_ptr = ty.ptr(old_ptr->AddressSpace(), new_ty, old_ptr->Access());
var->Result()->SetType(new_ptr);
@@ -370,7 +373,7 @@ struct State {
Result<SuccessType> SubstituteOverrides(Module& ir, const SubstituteOverridesConfig& cfg) {
TINT_CHECK_RESULT(
ValidateAndDumpIfNeeded(ir, "core.SubstituteOverrides", kSubstituteOverridesCapabilities));
ValidateBeforeIfNeeded(ir, kSubstituteOverridesCapabilities, "core.SubstituteOverrides"));
{
auto result = State{ir, cfg}.Process();
if (result != Success) {
3rdparty/dawn/src/tint/lang/core/ir/transform/value_to_let.cc
Vendored
+1 -1
View File
@@ -230,7 +230,7 @@ struct State {
} // namespace
Result<SuccessType> ValueToLet(Module& ir, const ValueToLetConfig& cfg) {
TINT_CHECK_RESULT(ValidateAndDumpIfNeeded(ir, "core.ValueToLet", kValueToLetCapabilities));
TINT_CHECK_RESULT(ValidateBeforeIfNeeded(ir, kValueToLetCapabilities, "core.ValueToLet"));
State{ir, cfg}.Process();
3rdparty/dawn/src/tint/lang/core/ir/transform/vectorize_scalar_matrix_constructors.cc
Vendored
+2 -2
View File
@@ -94,8 +94,8 @@ struct State {
} // namespace
Result<SuccessType> VectorizeScalarMatrixConstructors(Module& ir) {
TINT_CHECK_RESULT(ValidateAndDumpIfNeeded(ir, "core.VectorizeScalarMatrixConstructors",
kVectorizeScalarMatrixConstructorsCapabilities));
TINT_CHECK_RESULT(ValidateBeforeIfNeeded(ir, kVectorizeScalarMatrixConstructorsCapabilities,
"core.VectorizeScalarMatrixConstructors"));
State{ir}.Process();
3rdparty/dawn/src/tint/lang/core/ir/transform/vertex_pulling.cc
Vendored
+48 -49
View File
@@ -363,7 +363,7 @@ struct State {
auto load_i32 = [&](uint32_t offset) { return b.Bitcast<i32>(load_u32(offset))->Result(); };
auto load_f32 = [&](uint32_t offset) { return b.Bitcast<f32>(load_u32(offset))->Result(); };
auto load_ivec = [&](uint32_t offset, uint32_t bits, const core::type::Vector* vec) {
// For a vec2<u32>, we read the `xxxx'yyyy` u32 word. We then splat to a vec2 and left
// For a vec2u, we read the `xxxx'yyyy` u32 word. We then splat to a vec2 and left
// shift so we have `(xxxx'yyyy, yyyy'xxxx)`. Finally, we right shift to produce
// `(0000'xxxx, 0000'yyyy)`
auto* uvec = ty.MatchWidth(ty.u32(), vec);
@@ -421,27 +421,27 @@ struct State {
case VertexFormat::kUint16x4: {
auto* xy = load_ivec(0, 16, ty.vec2u());
auto* zw = load_ivec(1, 16, ty.vec2u());
return b.Construct<vec4<u32>>(xy, zw)->Result();
return b.Construct<vec4u>(xy, zw)->Result();
}
case VertexFormat::kUint32:
return load_u32(0);
case VertexFormat::kUint32x2: {
auto* x = load_u32(0);
auto* y = load_u32(1);
return b.Construct<vec2<u32>>(x, y)->Result();
return b.Construct<vec2u>(x, y)->Result();
}
case VertexFormat::kUint32x3: {
auto* x = load_u32(0);
auto* y = load_u32(1);
auto* z = load_u32(2);
return b.Construct<vec3<u32>>(x, y, z)->Result();
return b.Construct<vec3u>(x, y, z)->Result();
}
case VertexFormat::kUint32x4: {
auto* x = load_u32(0);
auto* y = load_u32(1);
auto* z = load_u32(2);
auto* w = load_u32(3);
return b.Construct<vec4<u32>>(x, y, z, w)->Result();
return b.Construct<vec4u>(x, y, z, w)->Result();
}
// Formats that are always i32 in the shader (or vectors of i32).
@@ -467,27 +467,27 @@ struct State {
case VertexFormat::kSint16x4: {
auto* xy = load_ivec(0, 16, ty.vec2i());
auto* zw = load_ivec(1, 16, ty.vec2i());
return b.Construct<vec4<i32>>(xy, zw)->Result();
return b.Construct<vec4i>(xy, zw)->Result();
}
case VertexFormat::kSint32:
return load_i32(0);
case VertexFormat::kSint32x2: {
auto* x = load_i32(0);
auto* y = load_i32(1);
return b.Construct<vec2<i32>>(x, y)->Result();
return b.Construct<vec2i>(x, y)->Result();
}
case VertexFormat::kSint32x3: {
auto* x = load_i32(0);
auto* y = load_i32(1);
auto* z = load_i32(2);
return b.Construct<vec3<i32>>(x, y, z)->Result();
return b.Construct<vec3i>(x, y, z)->Result();
}
case VertexFormat::kSint32x4: {
auto* x = load_i32(0);
auto* y = load_i32(1);
auto* z = load_i32(2);
auto* w = load_i32(3);
return b.Construct<vec4<i32>>(x, y, z, w)->Result();
return b.Construct<vec4i>(x, y, z, w)->Result();
}
// Unsigned normalized formats.
@@ -496,7 +496,7 @@ struct State {
// ******xx
auto* word = load_u32(0);
// 000000xx, ********, ********, ********
auto* unpack = b.Call<vec4<f32>>(core::BuiltinFn::kUnpack4X8Unorm, word);
auto* unpack = b.Call<vec4f>(core::BuiltinFn::kUnpack4X8Unorm, word);
// 000000xx
return float_value(b.Access<f32>(unpack, 0_u)->Result());
}
@@ -504,30 +504,30 @@ struct State {
// ****yyxx
auto* word = load_u32(0);
// 000000xx, 000000yy, ********, ********
auto* unpack = b.Call<vec4<f32>>(core::BuiltinFn::kUnpack4X8Unorm, word);
auto* unpack = b.Call<vec4f>(core::BuiltinFn::kUnpack4X8Unorm, word);
// 000000xx, 000000yy
return float_value(b.Swizzle<vec2<f32>>(unpack, Vector{0u, 1u})->Result());
return float_value(b.Swizzle<vec2f>(unpack, Vector{0u, 1u})->Result());
}
case VertexFormat::kUnorm8x4: {
// wwzzyyxx
auto* word = load_u32(0);
// 000000xx, 000000yy, 000000zz, 000000ww
auto* unpack = b.Call<vec4<f32>>(core::BuiltinFn::kUnpack4X8Unorm, word);
auto* unpack = b.Call<vec4f>(core::BuiltinFn::kUnpack4X8Unorm, word);
return float_value(unpack->Result());
}
case VertexFormat::kUnorm8x4BGRA: {
// wwzzyyxx
auto* word = load_u32(0);
// 000000xx, 000000yy, 000000zz, 000000ww
auto* unpack = b.Call<vec4<f32>>(core::BuiltinFn::kUnpack4X8Unorm, word);
auto* unpack = b.Call<vec4f>(core::BuiltinFn::kUnpack4X8Unorm, word);
// 000000zz, 000000yy, 000000xx, 000000ww
return float_value(b.Swizzle<vec4<f32>>(unpack, Vector{2u, 1u, 0u, 3u})->Result());
return float_value(b.Swizzle<vec4f>(unpack, Vector{2u, 1u, 0u, 3u})->Result());
}
case VertexFormat::kUnorm16: {
// ****xxxx
auto* word = load_u32(0);
// 0000xxxx, ********
auto* unpack = b.Call<vec2<f32>>(core::BuiltinFn::kUnpack2X16Unorm, word);
auto* unpack = b.Call<vec2f>(core::BuiltinFn::kUnpack2X16Unorm, word);
// 0000xxxx
return float_value(b.Access<f32>(unpack, 0_u)->Result());
}
@@ -535,7 +535,7 @@ struct State {
// yyyyxxxx
auto* word = load_u32(0);
// 0000xxxx, 0000yyyy
auto* unpack = b.Call<vec2<f32>>(core::BuiltinFn::kUnpack2X16Unorm, word);
auto* unpack = b.Call<vec2f>(core::BuiltinFn::kUnpack2X16Unorm, word);
return float_value(unpack->Result());
}
case VertexFormat::kUnorm16x4: {
@@ -543,9 +543,9 @@ struct State {
auto* word0 = load_u32(0);
auto* word1 = load_u32(1);
// 0000xxxx, 0000yyyy, 0000zzzz, 0000wwww
auto* unpack0 = b.Call<vec2<f32>>(core::BuiltinFn::kUnpack2X16Unorm, word0);
auto* unpack1 = b.Call<vec2<f32>>(core::BuiltinFn::kUnpack2X16Unorm, word1);
return float_value(b.Construct<vec4<f32>>(unpack0, unpack1)->Result());
auto* unpack0 = b.Call<vec2f>(core::BuiltinFn::kUnpack2X16Unorm, word0);
auto* unpack1 = b.Call<vec2f>(core::BuiltinFn::kUnpack2X16Unorm, word1);
return float_value(b.Construct<vec4f>(unpack0, unpack1)->Result());
}
// Signed normalized formats.
@@ -554,7 +554,7 @@ struct State {
// ******xx
auto* word = load_u32(0);
// 000000xx, ********, ********, ********
auto* unpack = b.Call<vec4<f32>>(core::BuiltinFn::kUnpack4X8Snorm, word);
auto* unpack = b.Call<vec4f>(core::BuiltinFn::kUnpack4X8Snorm, word);
// 000000xx
return float_value(b.Access<f32>(unpack, 0_u)->Result());
}
@@ -562,22 +562,22 @@ struct State {
// ****yyxx
auto* word = load_u32(0);
// 000000xx, 000000yy, ********, ********
auto* unpack = b.Call<vec4<f32>>(core::BuiltinFn::kUnpack4X8Snorm, word);
auto* unpack = b.Call<vec4f>(core::BuiltinFn::kUnpack4X8Snorm, word);
// 000000xx, 000000yy
return float_value(b.Swizzle<vec2<f32>>(unpack, Vector{0u, 1u})->Result());
return float_value(b.Swizzle<vec2f>(unpack, Vector{0u, 1u})->Result());
}
case VertexFormat::kSnorm8x4: {
// wwzzyyxx
auto* word = load_u32(0);
// 000000xx, 000000yy, 000000zz, 000000ww
auto* unpack = b.Call<vec4<f32>>(core::BuiltinFn::kUnpack4X8Snorm, word);
auto* unpack = b.Call<vec4f>(core::BuiltinFn::kUnpack4X8Snorm, word);
return float_value(unpack->Result());
}
case VertexFormat::kSnorm16: {
// ****xxxx
auto* word = load_u32(0);
// 0000xxxx, ********
auto* unpack = b.Call<vec2<f32>>(core::BuiltinFn::kUnpack2X16Snorm, word);
auto* unpack = b.Call<vec2f>(core::BuiltinFn::kUnpack2X16Snorm, word);
// 0000xxxx
return float_value(b.Access<f32>(unpack, 0_u)->Result());
}
@@ -585,7 +585,7 @@ struct State {
// yyyyxxxx
auto* word = load_u32(0);
// 0000xxxx, 0000yyyy
auto* unpack = b.Call<vec2<f32>>(core::BuiltinFn::kUnpack2X16Snorm, word);
auto* unpack = b.Call<vec2f>(core::BuiltinFn::kUnpack2X16Snorm, word);
return float_value(unpack->Result());
}
case VertexFormat::kSnorm16x4: {
@@ -593,9 +593,9 @@ struct State {
auto* word0 = load_u32(0);
auto* word1 = load_u32(1);
// 0000xxxx, 0000yyyy, 0000zzzz, 0000wwww
auto* unpack0 = b.Call<vec2<f32>>(core::BuiltinFn::kUnpack2X16Snorm, word0);
auto* unpack1 = b.Call<vec2<f32>>(core::BuiltinFn::kUnpack2X16Snorm, word1);
return float_value(b.Construct<vec4<f32>>(unpack0, unpack1)->Result());
auto* unpack0 = b.Call<vec2f>(core::BuiltinFn::kUnpack2X16Snorm, word0);
auto* unpack1 = b.Call<vec2f>(core::BuiltinFn::kUnpack2X16Snorm, word1);
return float_value(b.Construct<vec4f>(unpack0, unpack1)->Result());
}
// F16 formats that can either be f16 or f32 in the shader.
@@ -605,12 +605,12 @@ struct State {
auto* word = load_u32(0);
if (shader_element_type->Is<core::type::F16>()) {
// xxxx, ****
auto* bitcast = b.Bitcast<vec2<f16>>(word);
auto* bitcast = b.Bitcast<vec2h>(word);
// xxxx
return b.Access<f16>(bitcast, 0_u)->Result();
} else {
// 0000xxxx, ********
auto* unpack = b.Call<vec2<f32>>(core::BuiltinFn::kUnpack2X16Float, word);
auto* unpack = b.Call<vec2f>(core::BuiltinFn::kUnpack2X16Float, word);
// 0000xxxx
return b.Access<f32>(unpack, 0_u)->Result();
}
@@ -620,10 +620,10 @@ struct State {
auto* word = load_u32(0);
if (shader_element_type->Is<core::type::F16>()) {
// xxxx, yyyy
return b.Bitcast<vec2<f16>>(word)->Result();
return b.Bitcast<vec2h>(word)->Result();
} else {
// 0000xxxx, 0000yyyy
auto* unpack = b.Call<vec2<f32>>(core::BuiltinFn::kUnpack2X16Float, word);
auto* unpack = b.Call<vec2f>(core::BuiltinFn::kUnpack2X16Float, word);
return unpack->Result();
}
}
@@ -633,14 +633,14 @@ struct State {
auto* word1 = load_u32(1);
if (shader_element_type->Is<core::type::F16>()) {
// xxxx, yyyy, zzzz, wwww
auto* bitcast0 = b.Bitcast<vec2<f16>>(word0);
auto* bitcast1 = b.Bitcast<vec2<f16>>(word1);
return b.Construct<vec4<f16>>(bitcast0, bitcast1)->Result();
auto* bitcast0 = b.Bitcast<vec2h>(word0);
auto* bitcast1 = b.Bitcast<vec2h>(word1);
return b.Construct<vec4h>(bitcast0, bitcast1)->Result();
} else {
// 0000xxxx, 0000yyyy, 0000zzzz, 0000wwww
auto* unpack0 = b.Call<vec2<f32>>(core::BuiltinFn::kUnpack2X16Float, word0);
auto* unpack1 = b.Call<vec2<f32>>(core::BuiltinFn::kUnpack2X16Float, word1);
return b.Construct<vec4<f32>>(unpack0, unpack1)->Result();
auto* unpack0 = b.Call<vec2f>(core::BuiltinFn::kUnpack2X16Float, word0);
auto* unpack1 = b.Call<vec2f>(core::BuiltinFn::kUnpack2X16Float, word1);
return b.Construct<vec4f>(unpack0, unpack1)->Result();
}
}
@@ -651,32 +651,32 @@ struct State {
case VertexFormat::kFloat32x2: {
auto* x = load_f32(0);
auto* y = load_f32(1);
return float_value(b.Construct<vec2<f32>>(x, y)->Result());
return float_value(b.Construct<vec2f>(x, y)->Result());
}
case VertexFormat::kFloat32x3: {
auto* x = load_f32(0);
auto* y = load_f32(1);
auto* z = load_f32(2);
return float_value(b.Construct<vec3<f32>>(x, y, z)->Result());
return float_value(b.Construct<vec3f>(x, y, z)->Result());
}
case VertexFormat::kFloat32x4: {
auto* x = load_f32(0);
auto* y = load_f32(1);
auto* z = load_f32(2);
auto* w = load_f32(3);
return float_value(b.Construct<vec4<f32>>(x, y, z, w)->Result());
return float_value(b.Construct<vec4f>(x, y, z, w)->Result());
}
// Miscellaneous other formats that need custom handling.
case VertexFormat::kUnorm10_10_10_2: {
auto* u32s = b.Construct<vec4<u32>>(load_u32(0));
auto* u32s = b.Construct<vec4u>(load_u32(0));
// shr = u32s >> vec4u(0, 10, 20, 30);
auto* shr = b.ShiftRight(u32s, b.Composite<vec4<u32>>(0_u, 10_u, 20_u, 30_u));
auto* shr = b.ShiftRight(u32s, b.Composite<vec4u>(0_u, 10_u, 20_u, 30_u));
// mask = shr & vec4u(0x3FF, 0x3FF, 0x3FF, 0x3);
auto* mask = b.And(shr, b.Composite<vec4<u32>>(0x3FF_u, 0x3FF_u, 0x3FF_u, 0x3_u));
auto* mask = b.And(shr, b.Composite<vec4u>(0x3FF_u, 0x3FF_u, 0x3FF_u, 0x3_u));
// vec4f(mask) / vec4f(1023, 1023, 1023, 3);
auto* div = b.Composite<vec4<f32>>(1023_f, 1023_f, 1023_f, 3_f);
return float_value(b.Divide(b.Convert<vec4<f32>>(mask), div)->Result());
auto* div = b.Composite<vec4f>(1023_f, 1023_f, 1023_f, 3_f);
return float_value(b.Divide(b.Convert<vec4f>(mask), div)->Result());
}
}
TINT_IR_UNREACHABLE(ir);
@@ -686,8 +686,7 @@ struct State {
} // namespace
Result<SuccessType> VertexPulling(core::ir::Module& ir, const VertexPullingConfig& config) {
TINT_CHECK_RESULT(
ValidateAndDumpIfNeeded(ir, "core.VertexPulling", kVertexPullingCapabilities));
TINT_CHECK_RESULT(ValidateBeforeIfNeeded(ir, kVertexPullingCapabilities, "core.VertexPulling"));
State{config, ir}.Process();
3rdparty/dawn/src/tint/lang/core/ir/transform/zero_init_workgroup_memory.cc
Vendored
+2 -2
View File
@@ -306,8 +306,8 @@ struct State {
} // namespace
Result<SuccessType> ZeroInitWorkgroupMemory(Module& ir) {
TINT_CHECK_RESULT(ValidateAndDumpIfNeeded(ir, "core.ZeroInitWorkgroupMemory",
kZeroInitWorkgroupMemoryCapabilities));
TINT_CHECK_RESULT(ValidateBeforeIfNeeded(ir, kZeroInitWorkgroupMemoryCapabilities,
"core.ZeroInitWorkgroupMemory"));
State{ir}.Process();
3rdparty/dawn/src/tint/lang/core/ir/validator.cc
Vendored
+174 -71
View File
@@ -96,6 +96,7 @@
#include "src/tint/lang/core/type/sampled_texture.h"
#include "src/tint/lang/core/type/storage_texture.h"
#include "src/tint/lang/core/type/type.h"
#include "src/tint/lang/core/type/u16.h"
#include "src/tint/lang/core/type/u32.h"
#include "src/tint/lang/core/type/u64.h"
#include "src/tint/lang/core/type/u8.h"
@@ -119,6 +120,7 @@
#define TINT_DUMP_IR_WHEN_VALIDATING 0
#if TINT_DUMP_IR_WHEN_VALIDATING
#include <iostream>
#include "src/tint/utils/text/styled_text_printer.h"
#endif
@@ -133,6 +135,19 @@ struct ValidatedType {
namespace {
/// Prints out the current IR state, iff TINT_DUMP_IR_WHEN_VALIDATING is set.
void DumpIRIfEnabled([[maybe_unused]] const Module& ir,
[[maybe_unused]] const char* msg = "",
[[maybe_unused]] std::string_view timing = "") {
#if TINT_DUMP_IR_WHEN_VALIDATING
auto printer = StyledTextPrinter::Create(stdout);
std::cout << "=========================================================\n";
std::cout << "== IR dump " << timing << " " << msg << ":\n";
std::cout << "=========================================================\n";
printer->Print(Disassembler(ir).Text());
#endif
}
using SupportedStages = tint::EnumSet<Function::PipelineStage>;
/// @returns a human-readable string of all the entries in a EnumSet
@@ -243,6 +258,19 @@ void WalkTypeAndMembers(CTX& ctx,
[&](const core::type::Array* a) { WalkArrayElements(ctx, a, impl); });
}
/// @returns true if the type or any contained types are atomic
/// @param ty root of the types to walks
bool ContainsAtomic(const core::type::Type* ty) {
bool found = false;
WalkTypeAndMembers(found, ty, IOAttributes{},
[&](bool& ctx, const core::type::Type* t, const IOAttributes&) {
if (t != nullptr && t->Is<core::type::Atomic>()) {
ctx = true;
}
});
return found;
}
/// The IO direction of an operation.
enum class IODirection : uint8_t { kInput, kOutput, kResource };
@@ -654,6 +682,8 @@ const BuiltInChecker& BuiltinCheckerFor(BuiltinValue builtin, const Capabilities
case BuiltinValue::kLocalInvocationId:
return kLocalInvocationIdChecker;
case BuiltinValue::kLocalInvocationIndex:
case BuiltinValue::kGlobalInvocationIndex:
case BuiltinValue::kWorkgroupIndex:
return kLocalInvocationIndexChecker;
case BuiltinValue::kNumSubgroups:
return kNumSubgroupsChecker;
@@ -1188,7 +1218,8 @@ class Validator {
/// @param ignore_caps a set of capabilities to ignore for this check
void CheckType(const core::type::Type* type,
std::function<diag::Diagnostic&()> diag,
Capabilities ignore_caps = {});
Capabilities ignore_caps = {},
Capabilities allow_caps = {});
/// Validates the root block
/// @param blk the block
@@ -1320,8 +1351,8 @@ class Validator {
void CheckInterpolation(const CastableBase* anchor,
const core::type::Type* ty,
const IOAttributes& attr,
const Function::PipelineStage stage,
const IODirection dir);
Function::PipelineStage stage,
IODirection dir);
/// Validates binding_point attributes on entry point IO.
/// @param anchor where to attach error messages to.
@@ -2106,7 +2137,8 @@ bool Validator::CheckResultsAndOperands(const ir::Instruction* inst,
void Validator::CheckType(const core::type::Type* root,
std::function<diag::Diagnostic&()> diag,
Capabilities ignore_caps) {
Capabilities ignore_caps,
Capabilities allow_caps) {
if (root == nullptr) {
return;
}
@@ -2301,7 +2333,8 @@ void Validator::CheckType(const core::type::Type* root,
},
[&](const core::type::U64*) {
// u64 types are guarded by the Allow64BitIntegers capability.
if (!capabilities_.Contains(Capability::kAllow64BitIntegers)) {
if (!capabilities_.Contains(Capability::kAllow64BitIntegers) &&
!allow_caps.Contains(Capability::kAllow64BitIntegers)) {
diag() << "64-bit integer types are not permitted";
return false;
}
@@ -2323,6 +2356,14 @@ void Validator::CheckType(const core::type::Type* root,
}
return true;
},
[&](const core::type::U16*) {
// u16 types are guarded by the Allow16BitIntegers capability.
if (!capabilities_.Contains(Capability::kAllow16BitIntegers)) {
diag() << "16-bit integer types are not permitted";
return false;
}
return true;
},
[&](const core::type::Array* arr) {
if (!arr->ElemType()->HasCreationFixedFootprint()) {
diag() << "array elements, " << NameOf(type)
@@ -2357,8 +2398,17 @@ void Validator::CheckType(const core::type::Type* root,
return true;
},
[&](const core::type::Atomic* a) {
if (!a->Type()->IsAnyOf<core::type::I32, core::type::U32>()) {
diag() << "atomic subtype must be i32 or u32";
// Prior to lowering we allow for atomic operations on vec2u to support the
// AtomicVec2UMinMax feature.
if (auto* vec = a->Type()->As<core::type::Vector>()) {
if (vec->Width() == 2 && vec->Type()->Is<core::type::U32>()) {
return true;
}
}
if (!a->Type()->IsAnyOf<core::type::I32, core::type::U32, core::type::U64>()) {
diag() << "atomic subtype must be i32, u32 or u64 type is "
<< NameOf(a->Type());
return false;
}
return true;
@@ -2468,15 +2518,21 @@ void Validator::CheckType(const core::type::Type* root,
continue;
}
// Visit the elements of a composite type.
auto type_count = ty->Elements();
if (type_count.type && seen.Add(type_count.type)) {
stack.Push(type_count.type);
continue;
}
for (uint32_t i = 0; i < type_count.count; i++) {
if (auto* subtype = ty->Element(i); subtype && seen.Add(subtype)) {
stack.Push(subtype);
if (type_count.type) {
// Every element has the same type (e.g. array, vector, matrix, ...), so validate that
// type once if it has not been seen before.
if (seen.Add(type_count.type)) {
stack.Push(type_count.type);
}
} else {
// Different elements have different types (e.g. a struct), so we need to validate each
// of them if they have not been seen before.
for (uint32_t i = 0; i < type_count.count; i++) {
if (auto* subtype = ty->Element(i); subtype && seen.Add(subtype)) {
stack.Push(subtype);
}
}
}
}
@@ -3335,13 +3391,21 @@ void Validator::CheckInstruction(const Instruction* inst) {
return;
}
Capabilities allowed_types{};
if (inst->Is<core::ir::Bitcast>()) {
allowed_types.Add(Capability::kAllow64BitIntegers);
}
auto results = inst->Results();
for (size_t i = 0; i < results.Length(); ++i) {
auto* res = results[i];
if (!res) {
continue;
}
CheckType(res->Type(), [&]() -> diag::Diagnostic& { return AddResultError(inst, i); });
CheckType(
res->Type(), [&]() -> diag::Diagnostic& { return AddResultError(inst, i); }, {},
allowed_types);
}
auto ops = inst->Operands();
@@ -3351,7 +3415,9 @@ void Validator::CheckInstruction(const Instruction* inst) {
continue;
}
CheckType(op->Type(), [&]() -> diag::Diagnostic& { return AddError(inst, i); });
CheckType(
op->Type(), [&]() -> diag::Diagnostic& { return AddError(inst, i); }, {},
allowed_types);
}
tint::Switch(
@@ -3436,6 +3502,21 @@ void Validator::CheckVar(const Var* var) {
}
}
if (ContainsAtomic(mv->StoreType())) {
if (mv->AddressSpace() == AddressSpace::kStorage) {
if (mv->Access() != core::Access::kReadWrite) {
AddError(var)
<< "atomic variables in 'storage' address space must have 'read_write' "
"access mode";
return;
}
} else if (mv->AddressSpace() != AddressSpace::kWorkgroup) {
AddError(var)
<< "atomic variables must be in the 'workgroup' or 'storage' address space";
return;
}
}
// Check that initializer and result type match
if (var->Initializer()) {
if (mv->AddressSpace() != AddressSpace::kFunction &&
@@ -3995,16 +4076,17 @@ void Validator::CheckBitcastTypes(const Bitcast* bitcast) {
return;
}
// S -> vec2<f16>
// S -> vec2<f16 | u16>
if (auto* vec_type = result_type->As<core::type::Vector>()) {
auto elements = vec_type->Elements();
if (elements.count == 2 && Is<core::type::F16>(elements.type)) {
if (elements.count == 2 && IsAnyOf<core::type::F16, core::type::U16>(elements.type)) {
return;
}
}
} else if (val_type->Is<core::type::F16>()) {
// f16 -> f16, identity
if (result_type->Is<core::type::F16>()) {
} else if (val_type->IsAnyOf<core::type::F16, core::type::U16>()) {
// S, where S is f16 or u16
// S -> S, identity and reinterpretation
if (result_type->IsAnyOf<core::type::F16, core::type::U16>()) {
return;
}
} else if (val_type->Is<core::type::Vector>()) {
@@ -4025,6 +4107,12 @@ void Validator::CheckBitcastTypes(const Bitcast* bitcast) {
}
}
// atomic vec2u support vec2u to u64 bitcast required
if (val_elements.count == 2 && val_elements.type->Is<core::type::U32>() &&
result_type->Is<core::type::U64>()) {
return; // assume backend supported
}
if (val_elements.type->IsAnyOf<core::type::U32, core::type::I32, core::type::F32>()) {
// vecN<S>, where S is i32, u32, or f32
// vecN<S> -> vecN<s>, identity and reinterpretation cases
@@ -4034,22 +4122,22 @@ void Validator::CheckBitcastTypes(const Bitcast* bitcast) {
return;
}
// vec2<S> -> vec4<f16>
// vec2<S> -> vec4<f16|u16>
if (val_elements.count == 2) {
if (result_elements.has_value() && result_elements->count == 4 &&
result_elements->type->Is<core::type::F16>()) {
result_elements->type->IsAnyOf<core::type::F16, core::type::U16>()) {
return;
}
}
} else if (val_elements.type->Is<core::type::F16>()) {
} else if (val_elements.type->IsAnyOf<core::type::F16, core::type::U16>()) {
if (result_elements.has_value()) {
if (result_elements->type->Is<core::type::F16>()) {
// vecN<f16> -> vecN<f16>, identity
if (result_elements->type->IsAnyOf<core::type::F16, core::type::U16>()) {
// vecN<f16|u16> -> vecN<f16|u16>, identity
if (val_elements.count == result_elements->count) {
return;
}
} else {
// vec4<f16> -> vec2<S>, where S is i32, u32, or f32
// vec4<f16|u16> -> vec2<S>, where S is i32, u32, or f32
if (val_elements.count == 4 && result_elements->count == 2 &&
result_elements->type
->IsAnyOf<core::type::U32, core::type::I32, core::type::F32>()) {
@@ -4057,7 +4145,7 @@ void Validator::CheckBitcastTypes(const Bitcast* bitcast) {
}
}
} else {
// vec2<f16> -> i32, u32, or f32
// vec2<f16|u16> -> i32, u32, or f32
if (val_elements.count == 2 &&
result_type->IsAnyOf<core::type::U32, core::type::I32, core::type::F32>()) {
return;
@@ -4103,6 +4191,17 @@ void Validator::CheckBuiltinCall(const BuiltinCall* call) {
return;
}
// Check evaluation stage of parameters that are required to be const-expressions.
for (uint32_t i = 0; i < builtin->parameters.Length(); i++) {
const auto& p = builtin->parameters[i];
const auto* arg = call->Args()[i];
if (p.is_const && !arg->Is<Constant>()) {
AddError(call, BuiltinCall::kArgsOperandOffset + i)
<< "the " << style::Variable(p.usage) << " argument must be a constant";
return;
}
}
if (auto* bc = call->As<CoreBuiltinCall>()) {
CheckCoreBuiltinCall(bc, builtin.Get());
}
@@ -4124,17 +4223,6 @@ void Validator::CheckBuiltinCall(const BuiltinCall* call) {
void Validator::CheckCoreBuiltinCall(const CoreBuiltinCall* call,
const core::intrinsic::Overload& overload) {
if (call->Func() == core::BuiltinFn::kQuadBroadcast ||
call->Func() == core::BuiltinFn::kSubgroupBroadcast) {
TINT_ASSERT(call->Args().Length() == 2);
constexpr uint32_t kIdArg = 1;
auto* id = call->Args()[kIdArg];
if (!id->Is<core::ir::Constant>()) {
AddError(call, kIdArg) << "non-constant ID provided";
}
return;
}
auto idx_for_usage = [&](core::ParameterUsage usage) -> std::optional<uint32_t> {
for (uint32_t i = 0; i < overload.parameters.Length(); ++i) {
auto& p = overload.parameters[i];
@@ -4154,20 +4242,13 @@ void Validator::CheckCoreBuiltinCall(const CoreBuiltinCall* call,
TINT_ASSERT(idx < call->Args().Length());
auto* val = call->Args()[idx];
if (auto* const_val = val->As<ir::Constant>()) {
auto* cnst = const_val->Value();
auto* const_val = val->As<ir::Constant>();
TINT_ASSERT(const_val);
auto* cnst = const_val->Value();
if (val->Type()->Is<core::type::Vector>()) {
for (size_t i = 0; i < cnst->NumElements(); i++) {
auto value = cnst->Index(i)->ValueAs<int32_t>();
if (value < min || value > max) {
AddError(call, idx)
<< value << " outside range of [" << min << ", " << max << "]";
return;
}
}
} else {
auto value = cnst->ValueAs<int32_t>();
if (val->Type()->Is<core::type::Vector>()) {
for (size_t i = 0; i < cnst->NumElements(); i++) {
auto value = cnst->Index(i)->ValueAs<int32_t>();
if (value < min || value > max) {
AddError(call, idx)
<< value << " outside range of [" << min << ", " << max << "]";
@@ -4175,13 +4256,18 @@ void Validator::CheckCoreBuiltinCall(const CoreBuiltinCall* call,
}
}
} else {
AddError(call, idx) << "expected a constant value";
return;
auto value = cnst->ValueAs<int32_t>();
if (value < min || value > max) {
AddError(call, idx) << value << " outside range of [" << min << ", " << max << "]";
return;
}
}
};
check_arg_in_range(core::ParameterUsage::kComponent, 0, 3);
check_arg_in_range(core::ParameterUsage::kOffset, -8, 7);
if (core::IsTexture(call->Func())) {
check_arg_in_range(core::ParameterUsage::kComponent, 0, 3);
check_arg_in_range(core::ParameterUsage::kOffset, -8, 7);
}
}
void Validator::CheckMemberBuiltinCall(const MemberBuiltinCall* call) {
@@ -5245,31 +5331,48 @@ const core::type::Type* Validator::GetVectorPtrElementType(const Instruction* in
} // namespace
Result<SuccessType> Validate(const Module& mod, Capabilities capabilities) {
Result<SuccessType> Validate(const Module& mod,
Capabilities capabilities,
[[maybe_unused]] const char* msg,
[[maybe_unused]] std::string_view timing) {
DumpIRIfEnabled(mod, msg, timing);
Validator v(mod, capabilities);
TINT_CHECK_RESULT(v.Run());
return Success;
}
Result<SuccessType> ValidateAndDumpIfNeeded([[maybe_unused]] const Module& ir,
[[maybe_unused]] const char* msg,
[[maybe_unused]] Capabilities capabilities,
[[maybe_unused]] std::string_view timing) {
#if TINT_DUMP_IR_WHEN_VALIDATING
auto printer = StyledTextPrinter::Create(stdout);
std::cout << "=========================================================\n";
std::cout << "== IR dump " << timing << " " << msg << ":\n";
std::cout << "=========================================================\n";
printer->Print(Disassembler(ir).Text());
#endif
Result<SuccessType> ValidateBefore(const Module& mod, Capabilities capabilities, const char* msg) {
return Validate(mod, capabilities, msg, "before");
}
Result<SuccessType> ValidateAfter(const Module& mod, Capabilities capabilities, const char* msg) {
return Validate(mod, capabilities, msg, "after");
}
Result<SuccessType> ValidateIfNeeded([[maybe_unused]] const Module& mod,
[[maybe_unused]] Capabilities capabilities,
[[maybe_unused]] const char* msg,
[[maybe_unused]] std::string_view timing) {
DumpIRIfEnabled(mod, msg, timing);
#if TINT_ENABLE_IR_VALIDATION
TINT_CHECK_RESULT(Validate(ir, capabilities));
Validator v(mod, capabilities);
TINT_CHECK_RESULT(v.Run());
#endif
return Success;
}
Result<SuccessType> ValidateBeforeIfNeeded(const Module& mod,
Capabilities capabilities,
const char* msg) {
return ValidateIfNeeded(mod, capabilities, msg, "before");
}
Result<SuccessType> ValidateAfterIfNeeded(const Module& mod,
Capabilities capabilities,
const char* msg) {
return ValidateIfNeeded(mod, capabilities, msg, "after");
}
} // namespace tint::core::ir
namespace std {
3rdparty/dawn/src/tint/lang/core/ir/validator.h
Vendored
+57 -12
View File
@@ -29,13 +29,13 @@
#define SRC_TINT_LANG_CORE_IR_VALIDATOR_H_
#include <vector>
#include "src/tint/utils/containers/enum_set.h"
#include "src/tint/utils/result.h"
// Forward declarations
namespace tint::core::ir {
class Module;
class Function;
} // namespace tint::core::ir
namespace tint::core::ir {
@@ -44,6 +44,8 @@ namespace tint::core::ir {
enum class Capability : uint8_t {
/// Allows 8-bit integer types.
kAllow8BitIntegers,
/// Allows 16-bit integer types.
kAllow16BitIntegers,
/// Allows 64-bit integer types.
kAllow64BitIntegers,
/// Allows ClipDistances on f32 and vecN<f32> parameters
@@ -97,22 +99,65 @@ enum class Capability : uint8_t {
/// Capabilities is a set of Capability
using Capabilities = EnumSet<Capability>;
/// Validates that a given IR module is correctly formed
/// Validates the module @p ir is correctly formed
/// @param mod the module to validate
/// @param capabilities the optional capabilities that are allowed
/// @returns success or failure
Result<SuccessType> Validate(const Module& mod, Capabilities capabilities = {});
/// Validates the module @p ir and dumps its contents if required by the build configuration.
/// @param ir the module to transform
/// @param msg the msg to accompany the output
/// @param capabilities the optional capabilities that are allowed
/// @param timing when the validation is run.
/// @returns success or failure
Result<SuccessType> ValidateAndDumpIfNeeded(const Module& ir,
const char* msg,
Capabilities capabilities = {},
std::string_view timing = "before");
Result<SuccessType> Validate(const Module& mod,
Capabilities capabilities = {},
const char* msg = "",
std::string_view timing = "");
/// Validates the module @p ir is correctly formed before an operation
/// @param mod the module to validate
/// @param capabilities the optional capabilities that are allowed
/// @param msg the msg to accompany the output to indicate the operation
/// @returns success or failure
Result<SuccessType> ValidateBefore(const Module& mod,
Capabilities capabilities = {},
const char* msg = "");
/// Validates the module @p ir is correctly formed after an operation
/// @param mod the module to validate
/// @param capabilities the optional capabilities that are allowed
/// @param msg the msg to accompany the output to indicate the step
/// @returns success or failure
Result<SuccessType> ValidateAfter(const Module& mod,
Capabilities capabilities = {},
const char* msg = "");
/// Validates the module @p ir is correctly formed, iff required by the build configuration.
/// @param mod the module to transform
/// @param capabilities the optional capabilities that are allowed
/// @param msg the msg to accompany the output
/// @param timing when the validation is run.
/// @returns success or failure
Result<SuccessType> ValidateIfNeeded(const Module& mod,
Capabilities capabilities = {},
const char* msg = "",
std::string_view timing = "");
/// Validates the module @p ir is correctly formed before an operation, iff required by the build
/// configuration.
/// @param mod the module to validate
/// @param capabilities the optional capabilities that are allowed
/// @param msg the msg to accompany the output to indicate the operation
/// @returns success or failure
Result<SuccessType> ValidateBeforeIfNeeded(const Module& mod,
Capabilities capabilities = {},
const char* msg = "");
/// Validates the module @p ir is correctly formed after an operation, iff required by the build
/// configuration.
/// @param mod the module to validate
/// @param capabilities the optional capabilities that are allowed
/// @param msg the msg to accompany the output to indicate the operation
/// @returns success or failure
Result<SuccessType> ValidateAfterIfNeeded(const Module& mod,
Capabilities capabilities = {},
const char* msg = "");
} // namespace tint::core::ir
3rdparty/dawn/src/tint/lang/core/number.h
Vendored
+5
View File
@@ -29,6 +29,7 @@
#define SRC_TINT_LANG_CORE_NUMBER_H_
#include <stdint.h>
#include <cmath>
#include <functional>
#include <limits>
@@ -291,6 +292,8 @@ using i8 = Number<int8_t>;
using i32 = Number<int32_t>;
/// `u8` is a type alias to `Number<uint8_t>`.
using u8 = Number<uint8_t>;
/// `u16` is a type alias to `Number<uint16_t>`.
using u16 = Number<uint16_t>;
/// `u32` is a type alias to `Number<uint32_t>`.
using u32 = Number<uint32_t>;
/// `u64` is a type alias to `Number<uint64_t>`.
@@ -322,6 +325,8 @@ const char* FriendlyName() {
return "abstract-float";
} else if constexpr (std::is_same_v<T, i32>) {
return "i32";
} else if constexpr (std::is_same_v<T, u16>) {
return "u16";
} else if constexpr (std::is_same_v<T, u32>) {
return "u32";
} else if constexpr (std::is_same_v<T, f32>) {
3rdparty/dawn/src/tint/lang/core/type/array.h
Vendored
+1
View File
@@ -29,6 +29,7 @@
#define SRC_TINT_LANG_CORE_TYPE_ARRAY_H_
#include <stdint.h>
#include <optional>
#include <string>
#include <variant>
3rdparty/dawn/src/tint/lang/core/type/array_count.h
Vendored
+1 -1
View File
@@ -83,7 +83,7 @@ class ConstantArrayCount final : public Castable<ConstantArrayCount, ArrayCount>
uint32_t value;
};
/// The variant of an ArrayCount when the array is is runtime-sized.
/// The variant of an ArrayCount when the array is runtime-sized.
/// Example:
/// ```
/// type arr = array<i32>
3rdparty/dawn/src/tint/lang/core/type/binding_array.h
Vendored
+1
View File
@@ -29,6 +29,7 @@
#define SRC_TINT_LANG_CORE_TYPE_BINDING_ARRAY_H_
#include <stdint.h>
#include <string>
#include <variant>
3rdparty/dawn/src/tint/lang/core/type/manager.cc
Vendored
+11
View File
@@ -46,6 +46,7 @@
#include "src/tint/lang/core/type/sampled_texture.h"
#include "src/tint/lang/core/type/storage_texture.h"
#include "src/tint/lang/core/type/type.h"
#include "src/tint/lang/core/type/u16.h"
#include "src/tint/lang/core/type/u32.h"
#include "src/tint/lang/core/type/u64.h"
#include "src/tint/lang/core/type/u8.h"
@@ -149,6 +150,10 @@ const core::type::U8* Manager::u8() {
return Get<core::type::U8>();
}
const core::type::U16* Manager::u16() {
return Get<core::type::U16>();
}
const core::type::U32* Manager::u32() {
return Get<core::type::U32>();
}
@@ -257,6 +262,12 @@ const core::type::SampledTexture* Manager::sampled_texture(TextureDimension dim,
return Get<core::type::SampledTexture>(dim, type);
}
const core::type::SampledTexture* Manager::sampled_texture(TextureDimension dim,
const core::type::Type* type,
TextureFilterable filterable) {
return Get<core::type::SampledTexture>(dim, type, filterable);
}
const core::type::MultisampledTexture* Manager::multisampled_texture(TextureDimension dim,
const core::type::Type* type) {
return Get<core::type::MultisampledTexture>(dim, type);
3rdparty/dawn/src/tint/lang/core/type/manager.h
Vendored
+20
View File
@@ -69,6 +69,7 @@ class Reference;
class SampledTexture;
class StorageTexture;
class U8;
class U16;
class U32;
class U64;
class Vector;
@@ -156,6 +157,8 @@ class Manager final {
return Get<core::type::I32>(std::forward<ARGS>(args)...);
} else if constexpr (std::is_same_v<T, tint::core::u8>) {
return Get<core::type::U8>(std::forward<ARGS>(args)...);
} else if constexpr (std::is_same_v<T, tint::core::u16>) {
return Get<core::type::U16>(std::forward<ARGS>(args)...);
} else if constexpr (std::is_same_v<T, tint::core::u32>) {
return Get<core::type::U32>(std::forward<ARGS>(args)...);
} else if constexpr (std::is_same_v<T, tint::core::u64>) {
@@ -221,6 +224,9 @@ class Manager final {
/// @returns a u8 type
const core::type::U8* u8();
/// @returns a u16 type
const core::type::U16* u16();
/// @returns a u32 type
const core::type::U32* u32();
@@ -314,6 +320,14 @@ class Manager final {
const core::type::SampledTexture* sampled_texture(TextureDimension dim,
const core::type::Type* type);
/// @param dim the dimensionality of the texture
/// @param type the data type of the sampled texture
/// @param filterable the filterablity
/// @returns a sampled texture type with the provided params
const core::type::SampledTexture* sampled_texture(TextureDimension dim,
const core::type::Type* type,
TextureFilterable filterable);
/// @param dim the dimensionality of the texture
/// @param type the data type of the sampled texture
/// @returns a multisampled texture type with the provided params
@@ -664,6 +678,12 @@ class Manager final {
return Get<core::type::Sampler>(core::type::SamplerKind::kSampler);
}
/// @param filtering the sampler filtering parameter
/// @returns the sampler type
const core::type::Sampler* sampler(SamplerFiltering filtering) {
return Get<core::type::Sampler>(core::type::SamplerKind::kSampler, filtering);
}
/// @returns the comparison sampler type
const core::type::Sampler* comparison_sampler() {
return Get<core::type::Sampler>(core::type::SamplerKind::kComparisonSampler);
3rdparty/dawn/src/tint/lang/spirv/type/resource_table.cc → 3rdparty/dawn/src/tint/lang/core/type/resource_table.cc
Vendored
+5 -5
View File
@@ -25,15 +25,15 @@
// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
#include "src/tint/lang/spirv/type/resource_table.h"
#include "src/tint/lang/core/type/resource_table.h"
#include <sstream>
#include "src/tint/lang/core/type/manager.h"
TINT_INSTANTIATE_TYPEINFO(tint::spirv::type::ResourceTable);
TINT_INSTANTIATE_TYPEINFO(tint::core::type::ResourceTable);
namespace tint::spirv::type {
namespace tint::core::type {
ResourceTable::ResourceTable(const core::type::Type* binding_type)
: Base(static_cast<size_t>(Hash(tint::TypeCode::Of<ResourceTable>().bits, binding_type)),
@@ -49,7 +49,7 @@ bool ResourceTable::Equals(const UniqueNode& other) const {
std::string ResourceTable::FriendlyName() const {
std::stringstream str;
str << "spirv.resource_table<" << binding_type_->FriendlyName() << ">";
str << "resource_table<" << binding_type_->FriendlyName() << ">";
return str.str();
}
@@ -68,4 +68,4 @@ ResourceTable* ResourceTable::Clone(core::type::CloneContext& ctx) const {
return ctx.dst.mgr->Get<ResourceTable>(binding_type);
}
} // namespace tint::spirv::type
} // namespace tint::core::type
3rdparty/dawn/src/tint/lang/spirv/type/resource_table.h → 3rdparty/dawn/src/tint/lang/core/type/resource_table.h
Vendored
+5 -5
View File
@@ -25,15 +25,15 @@
// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
#ifndef SRC_TINT_LANG_SPIRV_TYPE_RESOURCE_TABLE_H_
#define SRC_TINT_LANG_SPIRV_TYPE_RESOURCE_TABLE_H_
#ifndef SRC_TINT_LANG_CORE_TYPE_RESOURCE_TABLE_H_
#define SRC_TINT_LANG_CORE_TYPE_RESOURCE_TABLE_H_
#include <string>
#include "src/tint/lang/core/enums.h"
#include "src/tint/lang/core/type/type.h"
namespace tint::spirv::type {
namespace tint::core::type {
/// ResourceTable represents an OpTypeRuntimeArray of resources
class ResourceTable final : public Castable<ResourceTable, core::type::Type> {
@@ -68,6 +68,6 @@ class ResourceTable final : public Castable<ResourceTable, core::type::Type> {
const core::type::Type* binding_type_;
};
} // namespace tint::spirv::type
} // namespace tint::core::type
#endif // SRC_TINT_LANG_SPIRV_TYPE_RESOURCE_TABLE_H_
#endif // SRC_TINT_LANG_CORE_TYPE_RESOURCE_TABLE_H_
3rdparty/dawn/src/tint/lang/core/type/resource_type.cc
Vendored
+151 -18
View File
@@ -41,38 +41,62 @@ namespace tint::core::type {
const core::type::Type* ResourceTypeToType(core::type::Manager& ty, ResourceType type) {
switch (type) {
case ResourceType::kTexture1d_f32:
return ty.sampled_texture(core::type::TextureDimension::k1d, ty.f32());
case ResourceType::kTexture1d_f32_filterable:
return ty.sampled_texture(core::type::TextureDimension::k1d, ty.f32(),
core::TextureFilterable::kFilterable);
case ResourceType::kTexture1d_f32_unfilterable:
return ty.sampled_texture(core::type::TextureDimension::k1d, ty.f32(),
core::TextureFilterable::kUnfilterable);
case ResourceType::kTexture1d_i32:
return ty.sampled_texture(core::type::TextureDimension::k1d, ty.i32());
case ResourceType::kTexture1d_u32:
return ty.sampled_texture(core::type::TextureDimension::k1d, ty.u32());
case ResourceType::kTexture2d_f32:
return ty.sampled_texture(core::type::TextureDimension::k2d, ty.f32());
case ResourceType::kTexture2d_f32_filterable:
return ty.sampled_texture(core::type::TextureDimension::k2d, ty.f32(),
core::TextureFilterable::kFilterable);
case ResourceType::kTexture2d_f32_unfilterable:
return ty.sampled_texture(core::type::TextureDimension::k2d, ty.f32(),
core::TextureFilterable::kUnfilterable);
case ResourceType::kTexture2d_i32:
return ty.sampled_texture(core::type::TextureDimension::k2d, ty.i32());
case ResourceType::kTexture2d_u32:
return ty.sampled_texture(core::type::TextureDimension::k2d, ty.u32());
case ResourceType::kTexture2dArray_f32:
return ty.sampled_texture(core::type::TextureDimension::k2dArray, ty.f32());
case ResourceType::kTexture2dArray_f32_filterable:
return ty.sampled_texture(core::type::TextureDimension::k2dArray, ty.f32(),
core::TextureFilterable::kFilterable);
case ResourceType::kTexture2dArray_f32_unfilterable:
return ty.sampled_texture(core::type::TextureDimension::k2dArray, ty.f32(),
core::TextureFilterable::kUnfilterable);
case ResourceType::kTexture2dArray_i32:
return ty.sampled_texture(core::type::TextureDimension::k2dArray, ty.i32());
case ResourceType::kTexture2dArray_u32:
return ty.sampled_texture(core::type::TextureDimension::k2dArray, ty.u32());
case ResourceType::kTexture3d_f32:
return ty.sampled_texture(core::type::TextureDimension::k3d, ty.f32());
case ResourceType::kTexture3d_f32_filterable:
return ty.sampled_texture(core::type::TextureDimension::k3d, ty.f32(),
core::TextureFilterable::kFilterable);
case ResourceType::kTexture3d_f32_unfilterable:
return ty.sampled_texture(core::type::TextureDimension::k3d, ty.f32(),
core::TextureFilterable::kUnfilterable);
case ResourceType::kTexture3d_i32:
return ty.sampled_texture(core::type::TextureDimension::k3d, ty.i32());
case ResourceType::kTexture3d_u32:
return ty.sampled_texture(core::type::TextureDimension::k3d, ty.u32());
case ResourceType::kTextureCube_f32:
return ty.sampled_texture(core::type::TextureDimension::kCube, ty.f32());
case ResourceType::kTextureCube_f32_filterable:
return ty.sampled_texture(core::type::TextureDimension::kCube, ty.f32(),
core::TextureFilterable::kFilterable);
case ResourceType::kTextureCube_f32_unfilterable:
return ty.sampled_texture(core::type::TextureDimension::kCube, ty.f32(),
core::TextureFilterable::kUnfilterable);
case ResourceType::kTextureCube_i32:
return ty.sampled_texture(core::type::TextureDimension::kCube, ty.i32());
case ResourceType::kTextureCube_u32:
return ty.sampled_texture(core::type::TextureDimension::kCube, ty.u32());
case ResourceType::kTextureCubeArray_f32:
return ty.sampled_texture(core::type::TextureDimension::kCubeArray, ty.f32());
case ResourceType::kTextureCubeArray_f32_filterable:
return ty.sampled_texture(core::type::TextureDimension::kCubeArray, ty.f32(),
core::TextureFilterable::kFilterable);
case ResourceType::kTextureCubeArray_f32_unfilterable:
return ty.sampled_texture(core::type::TextureDimension::kCubeArray, ty.f32(),
core::TextureFilterable::kUnfilterable);
case ResourceType::kTextureCubeArray_i32:
return ty.sampled_texture(core::type::TextureDimension::kCubeArray, ty.i32());
case ResourceType::kTextureCubeArray_u32:
@@ -96,6 +120,12 @@ const core::type::Type* ResourceTypeToType(core::type::Manager& ty, ResourceType
case ResourceType::kTextureDepthMultisampled2d:
return ty.depth_multisampled_texture(core::type::TextureDimension::k2d);
case ResourceType::kSampler_filtering:
return ty.sampler(core::SamplerFiltering::kFiltering);
case ResourceType::kSampler_non_filtering:
return ty.sampler(core::SamplerFiltering::kNonFiltering);
case ResourceType::kSampler_comparison:
return ty.comparison_sampler();
default:
TINT_UNREACHABLE();
}
@@ -109,42 +139,96 @@ ResourceType TypeToResourceType(const core::type::Type* in_type) {
case core::type::TextureDimension::k1d:
return tint::Switch(
sa->Type(),
[&](const core::type::F32*) { return ResourceType::kTexture1d_f32; },
[&](const core::type::F32*) {
switch (sa->Filterable()) {
case core::TextureFilterable::kFilterable:
return ResourceType::kTexture1d_f32_filterable;
case core::TextureFilterable::kUnfilterable:
return ResourceType::kTexture1d_f32_unfilterable;
case core::TextureFilterable::kUndefined:
TINT_UNREACHABLE();
}
},
[&](const core::type::I32*) { return ResourceType::kTexture1d_i32; },
[&](const core::type::U32*) { return ResourceType::kTexture1d_u32; },
TINT_ICE_ON_NO_MATCH);
case core::type::TextureDimension::k2d:
return tint::Switch(
sa->Type(),
[&](const core::type::F32*) { return ResourceType::kTexture2d_f32; },
[&](const core::type::F32*) {
switch (sa->Filterable()) {
case core::TextureFilterable::kFilterable:
return ResourceType::kTexture2d_f32_filterable;
case core::TextureFilterable::kUnfilterable:
return ResourceType::kTexture2d_f32_unfilterable;
case core::TextureFilterable::kUndefined:
TINT_UNREACHABLE();
}
},
[&](const core::type::I32*) { return ResourceType::kTexture2d_i32; },
[&](const core::type::U32*) { return ResourceType::kTexture2d_u32; },
TINT_ICE_ON_NO_MATCH);
case core::type::TextureDimension::k2dArray:
return tint::Switch(
sa->Type(),
[&](const core::type::F32*) { return ResourceType::kTexture2dArray_f32; },
[&](const core::type::F32*) {
switch (sa->Filterable()) {
case core::TextureFilterable::kFilterable:
return ResourceType::kTexture2dArray_f32_filterable;
case core::TextureFilterable::kUnfilterable:
return ResourceType::kTexture2dArray_f32_unfilterable;
case core::TextureFilterable::kUndefined:
TINT_UNREACHABLE();
}
},
[&](const core::type::I32*) { return ResourceType::kTexture2dArray_i32; },
[&](const core::type::U32*) { return ResourceType::kTexture2dArray_u32; },
TINT_ICE_ON_NO_MATCH);
case core::type::TextureDimension::k3d:
return tint::Switch(
sa->Type(),
[&](const core::type::F32*) { return ResourceType::kTexture3d_f32; },
[&](const core::type::F32*) {
switch (sa->Filterable()) {
case core::TextureFilterable::kFilterable:
return ResourceType::kTexture3d_f32_filterable;
case core::TextureFilterable::kUnfilterable:
return ResourceType::kTexture3d_f32_unfilterable;
case core::TextureFilterable::kUndefined:
TINT_UNREACHABLE();
}
},
[&](const core::type::I32*) { return ResourceType::kTexture3d_i32; },
[&](const core::type::U32*) { return ResourceType::kTexture3d_u32; },
TINT_ICE_ON_NO_MATCH);
case core::type::TextureDimension::kCube:
return tint::Switch(
sa->Type(),
[&](const core::type::F32*) { return ResourceType::kTextureCube_f32; },
[&](const core::type::F32*) {
switch (sa->Filterable()) {
case core::TextureFilterable::kFilterable:
return ResourceType::kTextureCube_f32_filterable;
case core::TextureFilterable::kUnfilterable:
return ResourceType::kTextureCube_f32_unfilterable;
case core::TextureFilterable::kUndefined:
TINT_UNREACHABLE();
}
},
[&](const core::type::I32*) { return ResourceType::kTextureCube_i32; },
[&](const core::type::U32*) { return ResourceType::kTextureCube_u32; },
TINT_ICE_ON_NO_MATCH);
case core::type::TextureDimension::kCubeArray:
return tint::Switch(
sa->Type(),
[&](const core::type::F32*) { return ResourceType::kTextureCubeArray_f32; },
[&](const core::type::F32*) {
switch (sa->Filterable()) {
case core::TextureFilterable::kFilterable:
return ResourceType::kTextureCubeArray_f32_filterable;
case core::TextureFilterable::kUnfilterable:
return ResourceType::kTextureCubeArray_f32_unfilterable;
case core::TextureFilterable::kUndefined:
TINT_UNREACHABLE();
}
},
[&](const core::type::I32*) { return ResourceType::kTextureCubeArray_i32; },
[&](const core::type::U32*) { return ResourceType::kTextureCubeArray_u32; },
TINT_ICE_ON_NO_MATCH);
@@ -177,7 +261,56 @@ ResourceType TypeToResourceType(const core::type::Type* in_type) {
TINT_UNREACHABLE();
}
},
[&](const core::type::Sampler* s) {
if (s->IsComparison()) {
return ResourceType::kSampler_comparison;
}
TINT_ASSERT(s->Filtering() != core::SamplerFiltering::kUndefined);
if (s->Filtering() == core::SamplerFiltering::kFiltering) {
return ResourceType::kSampler_filtering;
}
return ResourceType::kSampler_non_filtering;
},
TINT_ICE_ON_NO_MATCH);
}
std::vector<ResourceType> ConvertsFrom(const core::type::Type* in_type) {
return tint::Switch(
in_type,
[&](const core::type::SampledTexture* sa) -> std::vector<ResourceType> {
if (sa->Filterable() != core::TextureFilterable::kUnfilterable) {
return {};
}
TINT_ASSERT((sa->Type()->Is<core::type::F32>()));
switch (sa->Dim()) {
case core::type::TextureDimension::k1d:
return {ResourceType::kTexture1d_f32_filterable};
case core::type::TextureDimension::k2d:
return {ResourceType::kTexture2d_f32_filterable, ResourceType::kTextureDepth2d};
case core::type::TextureDimension::k2dArray:
return {ResourceType::kTexture2dArray_f32_filterable,
ResourceType::kTextureDepth2dArray};
case core::type::TextureDimension::k3d:
return {ResourceType::kTexture3d_f32_filterable};
case core::type::TextureDimension::kCube:
return {ResourceType::kTextureCube_f32_filterable,
ResourceType::kTextureDepthCube};
case core::type::TextureDimension::kCubeArray:
return {ResourceType::kTextureCubeArray_f32_unfilterable,
ResourceType::kTextureDepthCubeArray};
case core::type::TextureDimension::kNone:
TINT_UNREACHABLE();
}
},
[&](const core::type::Sampler* s) -> std::vector<ResourceType> {
if (s->Filtering() == core::SamplerFiltering::kNonFiltering) {
return {ResourceType::kSampler_filtering};
}
return {};
},
[](Default) -> std::vector<ResourceType> { return {}; });
}
} // namespace tint::core::type
3rdparty/dawn/src/tint/lang/core/type/resource_type.h
Vendored
+5
View File
@@ -28,6 +28,8 @@
#ifndef SRC_TINT_LANG_CORE_TYPE_RESOURCE_TYPE_H_
#define SRC_TINT_LANG_CORE_TYPE_RESOURCE_TYPE_H_
#include <vector>
#include "src/tint/api/common/resource_type.h"
#include "src/tint/lang/core/type/type.h"
@@ -39,6 +41,9 @@ ResourceType TypeToResourceType(const core::type::Type* in_type);
// Converts a `ResourceType` back into a `core::type::Type`
const core::type::Type* ResourceTypeToType(core::type::Manager& ty, ResourceType type);
// For `in_type`, if it can be convert to from another type, returns the type it converts from.
std::vector<ResourceType> ConvertsFrom(const core::type::Type* in_type);
} // namespace tint::core::type
#endif // SRC_TINT_LANG_CORE_TYPE_RESOURCE_TYPE_H_
3rdparty/dawn/src/tint/lang/core/type/sampled_texture.cc
Vendored
+14 -6
View File
@@ -27,9 +27,9 @@
#include "src/tint/lang/core/type/sampled_texture.h"
#include "src/tint/lang/core/enums.h"
#include "src/tint/lang/core/type/manager.h"
#include "src/tint/lang/core/type/texture_dimension.h"
#include "src/tint/utils/diagnostic/diagnostic.h"
#include "src/tint/utils/ice/ice.h"
#include "src/tint/utils/math/hash.h"
#include "src/tint/utils/text/string_stream.h"
@@ -38,8 +38,12 @@ TINT_INSTANTIATE_TYPEINFO(tint::core::type::SampledTexture);
namespace tint::core::type {
SampledTexture::SampledTexture(TextureDimension dim, const type::Type* type)
: Base(Hash(TypeCode::Of<SampledTexture>().bits, dim, type), dim), type_(type) {
SampledTexture::SampledTexture(TextureDimension dim,
const type::Type* type,
TextureFilterable filterable)
: Base(Hash(TypeCode::Of<SampledTexture>().bits, dim, type, filterable), dim),
type_(type),
filterable_(filterable) {
TINT_ASSERT(type_);
}
@@ -47,20 +51,24 @@ SampledTexture::~SampledTexture() = default;
bool SampledTexture::Equals(const UniqueNode& other) const {
if (auto* o = other.As<SampledTexture>()) {
return o->Dim() == Dim() && o->type_ == type_;
return o->Dim() == Dim() && o->type_ == type_ && o->filterable_ == filterable_;
}
return false;
}
std::string SampledTexture::FriendlyName() const {
StringStream out;
out << "texture_" << Dim() << "<" << type_->FriendlyName() << ">";
out << "texture_" << Dim() << "<" << type_->FriendlyName();
if (filterable_ != TextureFilterable::kUndefined) {
out << ", " << filterable_;
}
out << ">";
return out.str();
}
SampledTexture* SampledTexture::Clone(CloneContext& ctx) const {
auto* ty = type_->Clone(ctx);
return ctx.dst.mgr->Get<SampledTexture>(Dim(), ty);
return ctx.dst.mgr->Get<SampledTexture>(Dim(), ty, filterable_);
}
} // namespace tint::core::type
3rdparty/dawn/src/tint/lang/core/type/sampled_texture.h
Vendored
+9 -1
View File
@@ -30,6 +30,7 @@
#include <string>
#include "src/tint/lang/core/enums.h"
#include "src/tint/lang/core/type/texture.h"
#include "src/tint/lang/core/type/texture_dimension.h"
@@ -41,7 +42,10 @@ class SampledTexture final : public Castable<SampledTexture, Texture> {
/// Constructor
/// @param dim the dimensionality of the texture
/// @param type the data type of the sampled texture
SampledTexture(TextureDimension dim, const type::Type* type);
/// @param filterable the texture filterability
SampledTexture(TextureDimension dim,
const type::Type* type,
TextureFilterable filterable = TextureFilterable::kUndefined);
/// Destructor
~SampledTexture() override;
@@ -53,6 +57,9 @@ class SampledTexture final : public Castable<SampledTexture, Texture> {
/// @returns the subtype of the sampled texture
const type::Type* Type() const { return type_; }
/// @returns the filterability
TextureFilterable Filterable() const { return filterable_; }
/// @returns the name for this type that closely resembles how it would be
/// declared in WGSL.
std::string FriendlyName() const override;
@@ -63,6 +70,7 @@ class SampledTexture final : public Castable<SampledTexture, Texture> {
private:
const type::Type* type_;
TextureFilterable filterable_;
};
} // namespace tint::core::type
3rdparty/dawn/src/tint/lang/core/type/sampler.cc
Vendored
+13 -5
View File
@@ -34,14 +34,16 @@ TINT_INSTANTIATE_TYPEINFO(tint::core::type::Sampler);
namespace tint::core::type {
Sampler::Sampler(SamplerKind kind)
: Base(Hash(tint::TypeCode::Of<Sampler>().bits, kind), core::type::Flags{}), kind_(kind) {}
Sampler::Sampler(SamplerKind kind, SamplerFiltering filtering)
: Base(Hash(tint::TypeCode::Of<Sampler>().bits, kind, filtering), core::type::Flags{}),
kind_(kind),
filtering_(filtering) {}
Sampler::~Sampler() = default;
bool Sampler::Equals(const UniqueNode& other) const {
if (auto* o = other.As<Sampler>()) {
return o->kind_ == kind_;
return o->kind_ == kind_ && o->filtering_ == filtering_;
}
return false;
}
@@ -51,11 +53,17 @@ uint32_t Sampler::Align() const {
}
std::string Sampler::FriendlyName() const {
return kind_ == SamplerKind::kSampler ? "sampler" : "sampler_comparison";
if (kind_ == SamplerKind::kComparisonSampler) {
return "sampler_comparison";
}
if (filtering_ == SamplerFiltering::kUndefined) {
return "sampler";
}
return "sampler<" + std::string(ToString(filtering_)) + ">";
}
Sampler* Sampler::Clone(CloneContext& ctx) const {
return ctx.dst.mgr->Get<Sampler>(kind_);
return ctx.dst.mgr->Get<Sampler>(kind_, filtering_);
}
} // namespace tint::core::type
3rdparty/dawn/src/tint/lang/core/type/sampler.h
Vendored
+7 -1
View File
@@ -30,6 +30,7 @@
#include <string>
#include "src/tint/lang/core/enums.h"
#include "src/tint/lang/core/type/sampler_kind.h"
#include "src/tint/lang/core/type/type.h"
@@ -40,7 +41,8 @@ class Sampler final : public Castable<Sampler, Type> {
public:
/// Constructor
/// @param kind the kind of sampler
explicit Sampler(SamplerKind kind);
/// @param filtering the filtering flag
explicit Sampler(SamplerKind kind, SamplerFiltering filtering = SamplerFiltering::kUndefined);
/// Destructor
~Sampler() override;
@@ -52,6 +54,9 @@ class Sampler final : public Castable<Sampler, Type> {
/// @returns the sampler type
SamplerKind Kind() const { return kind_; }
/// @returns the sampler Filtering
SamplerFiltering Filtering() const { return filtering_; }
/// @returns true if this is a comparison sampler
bool IsComparison() const { return kind_ == SamplerKind::kComparisonSampler; }
@@ -68,6 +73,7 @@ class Sampler final : public Castable<Sampler, Type> {
private:
SamplerKind const kind_;
SamplerFiltering filtering_;
};
} // namespace tint::core::type
3rdparty/dawn/src/tint/lang/core/type/storage_texture.h
Vendored
-5
View File
@@ -34,11 +34,6 @@
#include "src/tint/lang/core/type/texture.h"
#include "src/tint/lang/core/type/texture_dimension.h"
// Forward declarations
namespace tint::core::type {
class Manager;
} // namespace tint::core::type
namespace tint::core::type {
/// A storage texture type.
3rdparty/dawn/src/tint/lang/core/type/swizzle_view.cc
Vendored
+79
View File
@@ -0,0 +1,79 @@
// Copyright 2026 The Dawn & Tint Authors
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are met:
//
// 1. Redistributions of source code must retain the above copyright notice, this
// list of conditions and the following disclaimer.
//
// 2. Redistributions in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// 3. Neither the name of the copyright holder nor the names of its
// contributors may be used to endorse or promote products derived from
// this software without specific prior written permission.
//
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
#include "src/tint/lang/core/type/swizzle_view.h"
#include "src/tint/lang/core/type/manager.h"
#include "src/tint/lang/core/type/vector.h"
#include "src/tint/utils/math/hash.h"
#include "src/tint/utils/text/string_stream.h"
TINT_INSTANTIATE_TYPEINFO(tint::core::type::SwizzleView);
namespace tint::core::type {
SwizzleView::SwizzleView(core::AddressSpace address_space,
const type::Type* store_type,
core::Access access,
uint32_t from,
uint32_t to)
: Base(Hash(tint::TypeCode::Of<SwizzleView>().bits), address_space, store_type, access),
from_(from),
to_(to) {
TINT_ASSERT(1 <= from_ && from_ < 5);
TINT_ASSERT(1 <= to_ && to_ < 5);
}
bool SwizzleView::Equals(const UniqueNode& other) const {
if (auto* o = other.As<SwizzleView>()) {
return o->AddressSpace() == AddressSpace() && o->StoreType() == StoreType() &&
o->Access() == Access() && o->FromSize() == FromSize() && o->ToSize() == ToSize();
}
return false;
}
std::string SwizzleView::FriendlyName() const {
StringStream out;
out << "swizzle<";
if (AddressSpace() != core::AddressSpace::kUndefined) {
out << AddressSpace() << ", ";
}
out << StoreType()->FriendlyName() << ", " << Access();
out << ", " << from_;
out << ", " << to_;
out << ">";
return out.str();
}
SwizzleView::~SwizzleView() = default;
SwizzleView* SwizzleView::Clone(CloneContext& ctx) const {
auto* ty = StoreType()->Clone(ctx);
return ctx.dst.mgr->Get<SwizzleView>(AddressSpace(), ty, Access(), from_, to_);
}
} // namespace tint::core::type
3rdparty/dawn/src/tint/lang/core/type/swizzle_view.h
Vendored
+80
View File
@@ -0,0 +1,80 @@
// Copyright 2026 The Dawn & Tint Authors
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are met:
//
// 1. Redistributions of source code must retain the above copyright notice, this
// list of conditions and the following disclaimer.
//
// 2. Redistributions in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// 3. Neither the name of the copyright holder nor the names of its
// contributors may be used to endorse or promote products derived from
// this software without specific prior written permission.
//
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
#ifndef SRC_TINT_LANG_CORE_TYPE_SWIZZLE_VIEW_H_
#define SRC_TINT_LANG_CORE_TYPE_SWIZZLE_VIEW_H_
#include <string>
#include "src/tint/lang/core/type/memory_view.h"
namespace tint::core::type {
/// A swizzle view type.
class SwizzleView final : public Castable<SwizzleView, MemoryView> {
public:
/// Constructor
/// @param address_space the address space of the object reference
/// @param store_type the swizzle's result type
/// @param access the resolved access control of the object reference
/// @param from the size of the object vector
/// @param to the size of the resulting vector produced by the swizzle operation
SwizzleView(core::AddressSpace address_space,
const Type* store_type,
core::Access access,
uint32_t from,
uint32_t to);
/// Destructor
~SwizzleView() override;
/// @returns the number of elements in the object vector
uint32_t FromSize() const { return from_; }
/// @returns the number of elements in the swizzle result
uint32_t ToSize() const { return to_; }
/// @param other the other node to compare against
/// @returns true if the this type is equal to @p other
bool Equals(const UniqueNode& other) const override;
/// @returns the name for this type that closely resembles how it would be
/// declared in WGSL.
std::string FriendlyName() const override;
/// @param ctx the clone context
/// @returns a clone of this type
SwizzleView* Clone(CloneContext& ctx) const override;
private:
const uint32_t from_;
const uint32_t to_;
};
} // namespace tint::core::type
#endif // SRC_TINT_LANG_CORE_TYPE_SWIZZLE_VIEW_H_
3rdparty/dawn/src/tint/lang/core/type/texel_buffer.h
Vendored
-5
View File
@@ -33,11 +33,6 @@
#include "src/tint/lang/core/enums.h"
#include "src/tint/lang/core/type/texture.h"
// Forward declarations
namespace tint::core::type {
class Manager;
} // namespace tint::core::type
namespace tint::core::type {
/// A texel buffer type.
3rdparty/dawn/src/tint/lang/core/type/type.cc
Vendored
+4 -3
View File
@@ -43,6 +43,7 @@
#include "src/tint/lang/core/type/struct.h"
#include "src/tint/lang/core/type/texel_buffer.h"
#include "src/tint/lang/core/type/texture.h"
#include "src/tint/lang/core/type/u16.h"
#include "src/tint/lang/core/type/u32.h"
#include "src/tint/lang/core/type/u64.h"
#include "src/tint/lang/core/type/u8.h"
@@ -63,7 +64,7 @@ Type::~Type() = default;
const Type* Type::UnwrapPtr() const {
auto* type = this;
while (auto* ptr = type->As<Pointer>()) {
if (auto* ptr = type->As<Pointer>()) {
type = ptr->StoreType();
}
return type;
@@ -114,7 +115,7 @@ bool Type::IsFloatScalarOrVector() const {
}
bool Type::IsIntegerScalar() const {
return IsAnyOf<U32, I32, U64, U8, I8>();
return IsAnyOf<U32, I32, U64, U8, I8, U16>();
}
bool Type::IsIntegerVector() const {
@@ -126,7 +127,7 @@ bool Type::IsSignedIntegerScalar() const {
}
bool Type::IsUnsignedIntegerScalar() const {
return IsAnyOf<U32, U64, U8>();
return IsAnyOf<U32, U64, U8, U16>();
}
bool Type::IsSignedIntegerVector() const {
3rdparty/dawn/src/tint/lang/core/type/type.h
Vendored
-4
View File
@@ -37,10 +37,6 @@
#include "src/tint/utils/containers/vector.h"
// Forward declarations
namespace tint {
class ProgramBuilder;
class SymbolTable;
} // namespace tint
namespace tint::core::type {
class Type;
} // namespace tint::core::type
3rdparty/dawn/src/tint/lang/core/type/u16.cc
Vendored
+63
View File
@@ -0,0 +1,63 @@
// Copyright 2026 The Dawn & Tint Authors
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are met:
//
// 1. Redistributions of source code must retain the above copyright notice, this
// list of conditions and the following disclaimer.
//
// 2. Redistributions in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// 3. Neither the name of the copyright holder nor the names of its
// contributors may be used to endorse or promote products derived from
// this software without specific prior written permission.
//
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
#include "src/tint/lang/core/type/u16.h"
#include "src/tint/lang/core/type/manager.h"
TINT_INSTANTIATE_TYPEINFO(tint::core::type::U16);
namespace tint::core::type {
U16::U16()
: Base(static_cast<size_t>(tint::TypeCode::Of<U16>().bits),
core::type::Flags{
Flag::kConstructable,
Flag::kCreationFixedFootprint,
Flag::kFixedFootprint,
Flag::kHostShareable,
}) {}
U16::~U16() = default;
std::string U16::FriendlyName() const {
return "u16";
}
uint32_t U16::Size() const {
return 2;
}
uint32_t U16::Align() const {
return 2;
}
U16* U16::Clone(CloneContext& ctx) const {
return ctx.dst.mgr->Get<U16>();
}
} // namespace tint::core::type
3rdparty/dawn/src/tint/lang/core/type/u16.h
Vendored
+63
View File
@@ -0,0 +1,63 @@
// Copyright 2026 The Dawn & Tint Authors
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are met:
//
// 1. Redistributions of source code must retain the above copyright notice, this
// list of conditions and the following disclaimer.
//
// 2. Redistributions in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// 3. Neither the name of the copyright holder nor the names of its
// contributors may be used to endorse or promote products derived from
// this software without specific prior written permission.
//
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
#ifndef SRC_TINT_LANG_CORE_TYPE_U16_H_
#define SRC_TINT_LANG_CORE_TYPE_U16_H_
#include <string>
#include "src/tint/lang/core/type/numeric_scalar.h"
namespace tint::core::type {
/// A unsigned int 16 type.
class U16 final : public Castable<U16, NumericScalar> {
public:
/// Constructor
U16();
/// Destructor
~U16() override;
/// @returns the name for this type that closely resembles how it would be
/// declared in WGSL.
std::string FriendlyName() const override;
/// @returns the size in bytes of the type.
uint32_t Size() const override;
/// @returns the alignment in bytes of the type.
uint32_t Align() const override;
/// @param ctx the clone context
/// @returns a clone of this type
U16* Clone(CloneContext& ctx) const override;
};
} // namespace tint::core::type
#endif // SRC_TINT_LANG_CORE_TYPE_U16_H_
3rdparty/dawn/src/tint/lang/spirv/intrinsic/data.cc
Vendored
+6440 -5244
View File
File diff suppressed because it is too large Load Diff
3rdparty/dawn/src/tint/lang/spirv/reader/common/options.h
Vendored
-6
View File
@@ -37,12 +37,6 @@ namespace tint::spirv::reader {
/// Options that control how the SPIR-V parser should behave.
struct Options {
/// Set to `true` to allow calls to derivative builtins in non-uniform control flow.
bool allow_non_uniform_derivatives = false;
// TODO(jrprice): Remove this when SPIR-V -> IR and IR -> WGSL are separate steps.
/// The extensions and language features that are allowed to be used in the generated WGSL.
wgsl::AllowedFeatures allowed_features = {};
/// Mapping from a SPIR-V Sampler binding point to a WGSL sampler binding
/// point. This allows remapping samplers which are split out of the
/// combined texture/sampler pairs in SPIR-V.
3rdparty/dawn/src/tint/lang/spirv/reader/lower/atomics.cc
Vendored
+8 -7
View File
@@ -554,13 +554,14 @@ struct State {
} // namespace
Result<SuccessType> Atomics(core::ir::Module& ir) {
TINT_CHECK_RESULT(ValidateAndDumpIfNeeded(ir, "spirv.Atomics",
core::ir::Capabilities{
core::ir::Capability::kAllowMultipleEntryPoints,
core::ir::Capability::kAllowOverrides,
core::ir::Capability::kAllowNonCoreTypes,
core::ir::Capability::kAllowPointerToHandle,
}));
TINT_CHECK_RESULT(ValidateBeforeIfNeeded(ir,
core::ir::Capabilities{
core::ir::Capability::kAllowMultipleEntryPoints,
core::ir::Capability::kAllowOverrides,
core::ir::Capability::kAllowNonCoreTypes,
core::ir::Capability::kAllowPointerToHandle,
},
"spirv.Atomics"));
State{ir}.Process();
3rdparty/dawn/src/tint/lang/spirv/reader/lower/builtins.cc
Vendored
+9 -8
View File
@@ -1310,14 +1310,15 @@ struct State {
Result<SuccessType> Builtins(core::ir::Module& ir) {
TINT_CHECK_RESULT(
ValidateAndDumpIfNeeded(ir, "spirv.Builtins",
core::ir::Capabilities{
core::ir::Capability::kAllowMultipleEntryPoints,
core::ir::Capability::kAllowOverrides,
core::ir::Capability::kAllowNonCoreTypes,
core::ir::Capability::kAllowStructMatrixDecorations,
core::ir::Capability::kAllowPointerToHandle,
}));
ValidateBeforeIfNeeded(ir,
core::ir::Capabilities{
core::ir::Capability::kAllowMultipleEntryPoints,
core::ir::Capability::kAllowOverrides,
core::ir::Capability::kAllowNonCoreTypes,
core::ir::Capability::kAllowStructMatrixDecorations,
core::ir::Capability::kAllowPointerToHandle,
},
"spirv.Builtins"));
State{ir}.Process();
3rdparty/dawn/src/tint/lang/spirv/reader/lower/decompose_strided_array.cc
Vendored
+9 -7
View File
@@ -251,13 +251,15 @@ struct State {
} // namespace
Result<SuccessType> DecomposeStridedArray(core::ir::Module& ir) {
TINT_CHECK_RESULT(ValidateAndDumpIfNeeded(ir, "spirv.DecomposeStridedArray",
core::ir::Capabilities{
core::ir::Capability::kAllowMultipleEntryPoints,
core::ir::Capability::kAllowNonCoreTypes,
core::ir::Capability::kAllowOverrides,
core::ir::Capability::kAllowPointerToHandle,
}));
TINT_CHECK_RESULT(
core::ir::ValidateBeforeIfNeeded(ir,
core::ir::Capabilities{
core::ir::Capability::kAllowMultipleEntryPoints,
core::ir::Capability::kAllowNonCoreTypes,
core::ir::Capability::kAllowOverrides,
core::ir::Capability::kAllowPointerToHandle,
},
"spirv.DecomposeStridedArray"));
State{ir}.Process();

Some files were not shown because too many files have changed in this diff Show More