| Age | Commit message (Collapse) | Author | Files | Lines |
|---|
|
(#119585)
A prior PR added a portion of the global address space modifications
required for declare target to, this PR seeks to add
a small amount more leftover from that PR.
The intent is to allow for more correct IR that the backends (in
particular AMDGPU) can treat more aptly for optimisations and code
correctness
1/3 required PRs to enable declare target to mapping, should look at PR
3/3 to check for full green passes (this one will fail a number due to
some dependencies).
Co-authored-by: Raghu Maddhipatla raghu.maddhipatla@amd.com
|
|
Currently we assume that `0` is the default AS, which is usually true,
but it isn't for `SPIR-V`.
Pass down the AS from `clang` and use it to create types.
After this change, we finally generate fully valid SPIR-V for a basic
OpenMP Offloading example.
Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>
|
|
MC Static Samplers Representation currently depends on Object
structures. This PR removes that dependency and in order to facilitate
removing to_underlying usage in follow-up PRs.
|
|
This is a follow-up to #150124. This PR makes it so that the
`-fopenmp-host-ir-file-path` respects VFS overlays, like any other input
file.
|
|
MC Descriptor Range Representation currently depend on Object
structures. This PR removes that dependency and in order to facilitate
removing to_underlying usage in follow-up PRs.
|
|
This pr implements the following validations:
1. Check that descriptor tables don't mix Sample and non-Sampler
resources
2. Ensure that descriptor ranges don't append onto an unbounded range
3. Ensure that descriptor ranges don't overflow
4. Adds a missing validation to ensure that only a single `RootFlags`
parameter is provided
Resolves: https://github.com/llvm/llvm-project/issues/153868.
|
|
We should set the correct target-specific AS for the SrcLocStr global
created in OMPIRBuilder.
We also may have to insert a constexpr addrspacecast because the struct
field type may be different than the value used to initialize it.
I actually want the cast to be from AS 1 to AS 4, but getting the type
to be AS4 relies on a PR currently in-review, so leave the cast target
to AS 0 for now.
---------
Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>
|
|
In the code that fix ups the debug information, we handles both the
debug intrinsics and debug records. The debug intrinsics are being
phased out and I recently changed mlir translation to not generate them.
This means that we should not get debug intrinsics anymore and code can
be simplified by removing their handling.
|
|
|
|
Calling `BasicBlock::splice` in `spliceBB` when both `Old` and `New` are
empty is a `nop` currently but it can cause a crash once debug records
are used instead of debug intrinsics. This PR makes the call conditional
on at least one of `Old` or `New` being non-empty.
Consider the following mlir:
```
omp.target map_entries() {
llvm.intr.dbg.declare ...
llvm.intr.dbg.declare ...
omp.teams ...
...
}
```
Current code would translate llvm.intr Ops to llvm intrinsics. Old is
the BasicBlock where they were get inserted and it will have 2 llvm
debug intrinsics by the time the implementation of `omp.teams` starts.
This implementation creates many BasicBlocks by calling `splitBB`. The
`New` is the just created BasicBlock which is empty.
In the new scheme (using debug records), there will be no instruction in
the `Old` BB after llvm.intr Ops get translated but just 2 trailing
debug records. So both `Old` and `New` are empty. When control reaches
`BasicBlock::splice`, it calls `spliceDebugInfoEmptyBlock`. This
function expects that in this case (`Src` is empty but has trailing
debug records), the `ToIt` is valid and it can call `adoptDbgRecords` on
it. This assumption is not true in this case as `New` is empty and
`ToIt` is pointing to end(). The fix is to only call
`BasicBlock::splice` when at least of `Old` or `New` is not empty.
|
|
DXC checks if registers are correctly bound to root signature
descriptors. This implements the same check.
closes: #[126645](https://github.com/llvm/llvm-project/issues/126645)
---------
Co-authored-by: joaosaffran <joao.saffran@microsoft.com>
Co-authored-by: Joao Saffran <jderezende@microsoft.com>
|
|
The check of whether a construct can privatize is a general utility, not
dependent on language-specific definitions.
|
|
The emitted code tests whether the current tile should executing the
remainder iterations by checking the logical iteration number is the one
after the floor iterations that execute the non-remainder iterations.
There are two counts of how many iterations there are: Those of
non-remainder iterations (simply rounded-down division of tripcount and
tile size), and those including an additional floor iteration for the
remainder iterations. The code was used the wrong one that caused the
condition to never match.
|
|
Parameter Header (#154249)
This patch is refactoring Root Parameter Header in DX Container backend
to remove the usage of `to_underlying`. This requires some changes:
first, MC Root Signature should not depend on Object/DXContainer.h;
Second, we need to assume data to be valid in scenarios where it was
originally not expected, this made some tests be removed.
|
|
We just replaced SmallSet<T *, N> with SmallPtrSet<T *, N>, bypassing
the redirection found in SmallSet.h. With that, we no longer need to
include SmallSet.h in many files.
|
|
Currently, Flang can generate no-loop kernels for all OpenMP target
kernels in the program if the flags
-fopenmp-assume-teams-oversubscription or
-fopenmp-assume-threads-oversubscription are set.
If we add an additional parameter, we can choose
in the future which OpenMP kernels should be generated as no-loop
kernels.
This PR doesn't modify current behavior of oversubscription flags.
RFC for no-loop kernels:
https://discourse.llvm.org/t/rfc-no-loop-mode-for-openmp-gpu-kernels/87517
|
|
|
|
Remove:
* DescriptorType enum - this almost exactly shadowed the ResourceClass
enum
* ClauseType aliased ResourceClass
Although these were introduced to make the HLSL root signature handling
code a bit cleaner, they were ultimately causing confusion as they
appeared to be unique enums that needed to be converted between each
other.
Closes #153890
|
|
Co-authored-by: Michael Klemm <michael.klemm@amd.com>
|
|
This patch replaces SmallSet<T *, N> with SmallPtrSet<T *, N>. Note
that SmallSet.h "redirects" SmallSet to SmallPtrSet for pointer
element types:
template <typename PointeeType, unsigned N>
class SmallSet<PointeeType*, N> : public SmallPtrSet<PointeeType*, N>
{};
We only have 140 instances that rely on this "redirection", with the
vast majority of them under llvm/. Since relying on the redirection
doesn't improve readability, this patch replaces SmallSet with
SmallPtrSet for pointer element types.
|
|
This patch refactors DXILABI to remove the dependency on scope printer.
Closes: #153827
---------
Co-authored-by: Joao Saffran <{ID}+{username}@users.noreply.github.com>
|
|
As part of the Root Signature Spec, we need to validate if Root
Signatures are not defining overlapping ranges.
Closes: https://github.com/llvm/llvm-project/issues/126645
---------
Co-authored-by: joaosaffran <joao.saffran@microsoft.com>
Co-authored-by: Joao Saffran <{ID}+{username}@users.noreply.github.com>
Co-authored-by: Joao Saffran <jderezende@microsoft.com>
|
|
Fixes #153043.
This is another case of debug location not getting updated when the
insert point is changed by the `restoreIP`. Fixed by using the wrapper
function that updates the debug location.
|
|
A function added in pr#151306 was under NDEBUG macro which caused the
build to fail in certain cases. It has been moved out of the #ifdef
check to ensure it is always compiled.
|
|
(#152253)
The resource binding analysis was incorrectly reducing the size of the
`Bindings` vector by one element after sorting and de-duplication. This
led to an inaccurate setting of the `HasOverlappingBinding` flag in the
`DXILResourceBindingInfo` analysis, as the truncated vector no longer
reflected the true binding state.
This update corrects the shrink logic and introduces an `assert` in the
`DXILPostOptimizationValidation` pass. The assertion will trigger if
`HasOverlappingBinding` is set but no corresponding error is detected,
helping catch future inconsistencies.
The bug surfaced when the `srv_metadata.hlsl` and `uav_metadata.hlsl`
tests were updated to include unbounded resource arrays as part of
https://github.com/llvm/llvm-project/issues/145422. These updated test
files are included in this PR, as they would cause the new assertion to
fire if the original issue remained unresolved.
Depends on #152250
|
|
This fixes #147063.
I tried to fix this issue in more general way in
https://github.com/llvm/llvm-project/pull/147091 but the reviewer
suggested to fix the locations which are causing this issue. So this is
a more targeted approach.
The `restoreIP` is frequently used in the `OMPIRBuilder` to change the
insert position. This function eventually calls
`SetInsertPoint(BasicBlock *TheBB, BasicBlock::iterator IP)`. This
function updates the insert point and the debug location. But if the
`IP` is pointing to the end of the `TheBB`, then the debug location is
not updated and we could have a mismatch between insert point and the
debug location.
The problem can occur in 2 different code patterns.
This code below shows the first scenario.
```
1. auto curPos = builder.saveIP();
2. builder.restoreIP(/* some new pos */);
3. // generate some code
4. builder.restoreIP(curPos);
```
If `curPos` points to the end of basic block, we could have a problem.
But it is easy one to handle as we have the location before hand and can
save the correct debug location before 2 and then restore it after 3.
This can be done either manually or using the `llvm::InsertPointGuard`
as shown below.
```
// manual approach
auto curPos = builder.saveIP();
llvm::DebugLoc DbgLoc = builder.getCurrentDebugLocation();
builder.restoreIP(/* some new pos */);
// generate some code
builder.SetCurrentDebugLocation(DbgLoc);
builder.restoreIP(curPos);
{
// using InsertPointGuard
llvm::InsertPointGuard IPG(builder);
builder.restoreIP(/* some new pos */);
// generate some code
}
```
This PR fixes one problematic case using the manual approach.
For the 2nd scenario, look at the code below.
```
1. void fn(InsertPointTy allocIP, InsertPointTy codegenIP) {
2. builder.setInsertPoint(allocIP);
3. // generate some alloca
4. builder.setInsertPoint(codegenIP);
5. }
```
The `fn` can be called from anywhere and we can't assume the debug
location of the builder is valid at the start of the function. So if 4
does not update the debug location because the `codegenIP` points at the
end of the block, the rest of the code can end up using the debug
location of the `allocaIP`. Unlike the first case, we don't have a debug
location that we can save before hand and restore afterwards.
The solution here is to use the location of the last instruction in that
block. I have added a wrapper function over `restoreIP` that could be
called for such cases. This PR uses it to fix one problematic case.
|
|
Static analysis flagged the second part of this range check as always
true. RegisterSpace is uint32_t therefore the max value is 0xFFFFFFFF
and so the first check is sufficient.
|
|
During the split of the various `Frontend/HLSL` libraries, there was an
oversight to duplicate the `ResourceClassNames` definitions. This commit
simply consolidates the definitions into `DXContainer.h` as
`getResourceClasses`
|
|
Scan reductions are supported in OpenMP with the help of scan directive.
Reduction clause of the for loop/simd directive can take an `inscan`
modifier along with the body of the directive specifying a `scan`
directive. This PR implements the lowering logic for scan reductions in
workshare loops of OpenMP.
The body of the for loop is split into two loops (Input phase loop and
Scan Phase loop) and a scan reduction loop is added in the middle. The
Input phase loop populates a temporary buffer with initial values that
are to be reduced. The buffer is used by the reduction loop to perform
scan reduction. Scan phase loop copies the values of the buffer to the
reduction variable before executing the scan phase. Below is a high
level view of the code generated.
```
<declare pointer to buffer> ptr
omp parallel {
size num_iters = <num_iters>
// temp buffer allocation
omp masked {
buff = malloc(num_iters*scanvarstype)
*ptr = buff
}
barrier;
// input phase loop
for (i: 0..<num_iters>) {
<input phase>;
buffer = *ptr;
buffer[i] = red;
}
// scan reduction
omp masked
{
for (int k = 0; k != ceil(log2(num_iters)); ++k) {
i=pow(2,k)
for (size cnt = last_iter; cnt >= i; --cnt) {
buffer = *ptr;
buffer[cnt] op= buffer[cnt-i];
}
}
}
barrier;
// scan phase loop
for (0..<num_iters>) {
buffer = *ptr;
red = buffer[i] ;
<scan phase>;
}
// temp buffer deletion
omp masked {
free(*ptr)
}
barrier;
}
```
The temporary buffer needs to be shared between all threads performing
reduction since it is read/written in Input and Scan workshare Loops.
This is achieved by declaring a pointer to the buffer in the shared
region and dynamically allocating the buffer by the master thread.
This is the reason why allocation, deallocation and scan reduction are
performed within `masked`. The code is verified to produce correct
results for Fortran programs with the code changes in the PR
https://github.com/llvm/llvm-project/pull/133149
|
|
options (#151579)
Adding mlir to llvm support for atomic control options.
Atomic Control Options are used to specify architectural characteristics
to help lowering of atomic operations. The options used are:
`-f[no-]atomic-remote-memory`, `-f[no-]atomic-fine-grained-memory`,
`-f[no-]atomic-ignore-denormal-mode`.
Legacy option `-m[no-]unsafe-fp-atomics` is aliased to
`-f[no-]ignore-denormal-mode`.
More details can be found in
https://github.com/llvm/llvm-project/pull/102569. This PR implements the
MLIR to LLVM lowering support of atomic control attributes specified
with OpenMP `atomicUpdateOp`.
Initial support can be found in PR:
https://github.com/llvm/llvm-project/pull/150860
|
|
Introduce the `enumToStringRef` enum into `ScopedPrinter.h` that
replicates `enumToString` behaviour, expect that instead of returning a
hex value string, it just returns an empty string. This allows us to
return a StringRef and easily check if an invalid enum was provided
based on the StringRef size
This then uses `enumToStringRef` to remove the redundant
`getResourceName` and `printEnum` functions.
Resolves: https://github.com/llvm/llvm-project/issues/151200.
|
|
Clean up some duplicated logic. We had two ways to do the same thing
here, and BindingInfoBuilder is more flexible.
|
|
|
|
This patch fixes:
llvm/lib/Frontend/Offloading/PropertySet.cpp:95:12: error: unused
variable '[It, Inserted]' [-Werror,-Wunused-variable]
|
|
This change fixes one of the failures in
https://github.com/llvm/llvm-project/pull/147321
Following code snippet:
`
for (const auto &[CategoryName, PropSet] : PSRegistry) {
J.attributeObject(CategoryName, [&] {
for (const auto &[PropName, PropVal] : PropSet) {
`
causes a build warning that is emitted as an error.
error: reference to local binding 'PropSet' declared in enclosing lambda
expression
This is resolved by capturing PropSet in a local variable.
Thanks
Signed-off-by: Arvind Sudarsanam <arvind.sudarsanam@intel.com>
|
|
This PR adds the `PropertySet` type, along with a pair of functions used
to serialize and deserialize into a JSON representation. A property set
is a key-value map, with values being one of 2 types - uint32 or byte
array. A property set registry is a collection of property sets, indexed
by a "category" name.
In SYCL offloading, property sets will be used to communicate metadata
about device images needed by the SYCL runtime. For example, there is a
property set which has a byte array containing the numeric ID, offset,
and size of each SYCL2020 spec constant. Another example is a property
set describing the optional kernel features used in the module: does it
use fp64? fp16? atomic64?
This metadata will be computed by `clang-sycl-linker` and the JSON
representation will be inserted in the string table of each
output `OffloadBinary`. This JSON will be consumed the SYCL offload
wrapper and will be lowered to the binary form SYCL runtime expects.
For example, consider this SYCL program that calls a kernel that uses
fp64:
```c++
#include <sycl/sycl.hpp>
using namespace sycl;
class MyKernel;
int main() {
queue q;
auto *p = malloc_shared<double>(1, q);
*p = .1;
q.single_task<MyKernel>([=]{ *p *= 2; }).wait();
std::cout << *p << "\n";
free(p, q);
}
```
The device code for this program would have the kernel marked with
`!sycl_used_aspects`:
```
define spir_kernel void @_ZTS8MyKernel([...]) !sycl_used_aspects !n { [...] }
!n = {i32 6}
```
`clang-sycl-linker` would recognize this metadata and then would output
the following JSON in the `OffloadBinary`'s key-value map:
```
{
"SYCL/device requirements": {
// aspects contains a list of sycl::aspect values used
// by the module; in this case just the value 6 encoded
// as a 4-byte little-endian integer
"aspects": "BjAwMA=="
}
}
```
The SYCL offload wrapper would lower those property sets to something
like this:
```c++
struct _sycl_device_binary_property_set_struct {
char *CategoryName;
_sycl_device_binary_property *PropertiesBegin;
_sycl_device_binary_property *PropertiesEnd;
};
struct _sycl_device_binary_property_struct {
char *PropertyName;
void *ValAddr;
uint64_t ValSize;
};
//
_sycl_device_binary_property_struct device_requirements[] = {
/* PropertyName */ "aspects",
/* ValAddr */ [pointer to the bytes 0x06 0x00 0x00 0x00],
/* ValSize */ 4,
};
_sycl_device_binary_property_set_struct properties[] = {
/* CategoryName */ "SYCL/device requirements",
/* PropertiesBegin */ device_requirements,
/* PropertiesEnd */ std::end(device_requirments),
}
```
---------
Co-authored-by: Arvind Sudarsanam <arvind.sudarsanam@intel.com>
|
|
This PR addresses
https://github.com/llvm/llvm-project/pull/144465#issuecomment-3063422828.
Using `joinErrors` and `llvm:Error` instead of boolean values.
---------
Co-authored-by: joaosaffran <joao.saffran@microsoft.com>
Co-authored-by: Joao Saffran <{ID}+{username}@users.noreply.github.com>
|
|
We extract the binding logic out of the DXILResource analysis passes into the
FrontendHLSL library. This will allow us to use this logic for resource and
root signature bindings in both the DirectX backend and the HLSL frontend.
|
|
`LocationDescription` contains both the insertion point and the debug
location. When `LocationDescription` is available, it is better to use
`updateToLocation` which will update both. This PR replaces
`restoreIP(Loc.IP)` with `updateToLocation(Loc)` as former may not
update debug location in all cases.
I am not checking the return value of `updateToLocation` because that is
checked just a few lines above in all cases and we would have returned
early if it failed.
|
|
(#150832)
Reverts llvm/llvm-project#147950
I noticed some fails in the reduction tests with clang after this
change. I need to understand the failures better. Reverting this for
now.
|
|
Metadata lib (#149221)
This PR, moves the existing Root Signature Metadata Parsing logic used
in `DXILRootSignature` to the common library used by both frontend and
backend. Closes:
[#145942](https://github.com/llvm/llvm-project/issues/145942)
---------
Co-authored-by: joaosaffran <joao.saffran@microsoft.com>
|
|
Reduction support: https://github.com/llvm/llvm-project/pull/146671
If Support is fixed in this PR
The problem for the IF clause in composite constructs was that wsloop
and simd both operate on the same CanonicalLoopInfo structure: with the
SIMD processed first, followed by the wsloop. Previously the IF clause
generated code like
```
if (cond) {
while (...) {
simd_loop_body;
}
} else {
while (...) {
nonsimd_loop_body;
}
}
```
The problem with this is that this invalidates the CanonicalLoopInfo
structure to be processed by the wsloop later. To avoid this, in this
patch I preserve the original loop, moving the IF clause inside of the
loop:
```
while (...) {
if (cond) {
simd_loop_body;
} else {
non_simd_loop_body;
}
}
```
On simple examples I tried LLVM was able to hoist the if condition
outside of the loop at -O3.
The disadvantage of this is that we cannot add the
llvm.loop.vectorize.enable attribute on either the SIMD or non-SIMD
loops because they both share a loop back edge. There's no way of
solving this without keeping the old design of having two different
loops: which cannot be represented using only one CanonicalLoopInfo
structure. I don't think the presence or absence of this attribute makes
much difference. In my testing it is the llvm.loop.parallel_access
metadata which makes the difference to vectorization. LLVM will
vectorize if legal whether or not this attribute is there in the TRUE
branch. In the FALSE branch this means the loop might be vectorized even
when the condition is false: but I think this is still standards
compliant: OpenMP 6.0 says that when the if clause is false that should
be treated like the SIMDLEN clause is one. The SIMDLEN clause is defined
as a "hint". For the same reason, SIMDLEN and SAFELEN clauses are
silently ignored when SIMD IF is used.
I think it is better to implement SIMD IF and ignore SIMDLEN and SAFELEN
and some vectorization encouragement metadata when combined with IF than
to ignore IF because IF could have correctness consequences whereas the
rest are optimiztion hints. For example, the user might use the IF
clause to disable SIMD programatically when it is known not safe to
vectorize the loop. In this case it is not at all safe to add the
parallel access or SAFELEN metadata.
|
|
(#148728)
There is a sanitizer fail in CI after this which I need to investigate.
Reverting for now.
Reverts llvm/llvm-project#148284
|
|
This is similar to https://github.com/llvm/llvm-project/pull/147950 but
for task proxy function.
|
|
`SemaHLSL` diagnostics (#147115)
At the moment, when we report diagnostics from `SemaHLSL` we only
provide the source location of the root signature attr. This allows for
significantly less helpful diagnostics (for eg. reporting resource range
overlaps).
This pr implements a way to retain the source location of a root element
when it is parsed, so that we can output the `SourceLocation` of each
root element that causes the overlap in the diagnostics during semantic
analysis.
This pr defines a wrapper struct `clang::hlsl::RootSignatureElement` in
`SemaHLSL` that will contain the underlying `RootElement` and can hold
any additional diagnostic information. This struct will be what is used
in `HLSLRootSignatureParser` and in `SemaHLSL`. Then the diagnostic
information will be stripped and the underlying element will be stored
in the `RootSignatureDecl`.
For the reporting of diagnostics, we can now use the retained
`SourceLocation` of each `RootElement` when reporting the range overlap,
and we can add a `note` diagnostic to highlight the other root element
as well.
- Defines `RootSignatureElement` in the `hlsl` namespace in `SemaHLSL`
(defined in `SemaHLSL` because `Parse` has a dependency on `Sema`)
- Updates parsing logic to construct `RootSignatureElement`s and retain
the source loction in `ParseHLSLRootSignature`
- Updates `SemaHLSL` when it constructs the `RootSignatureDecl` to take
the new `RootSignatureElement` and store the underlying `RootElement`
- Updates the current tests to ensure the new `note` diagnostic is
produced and that the `SourceLocation` is seen
- Slight update to the `RootSignatureValidations` api to ensure the
caller sorts and owns the memory of the passed in `RangeInfo`
- Adds a test to demonstrate the `SourceLocation` of both elements being
correctly pointed out
Resolves: https://github.com/llvm/llvm-project/issues/145819
|
|
We have this pattern of code in OMPIRBuilder for many functions that are
used in reduction operations.
```
Function *LtGRFunc = Function::Create
BasicBlock *EntryBlock = BasicBlock::Create(Ctx, "entry", LtGRFunc);
Builder.SetInsertPoint(EntryBlock);
```
The insertion point is moved to the new function but the debug location is not updated. This means that reduction function will use the debug location that points to another function. This problem gets hidden because these functions gets inlined but the potential for failure exists.
This patch resets the debug location when insertion point is moved to new function. Some `InsertPointGuard` have been added to make sure why restore the debug location correctly when we are done with the reduction function.
|
|
(#147111)
This pr resolves some discrepancies in verification during `validate` in
`DXILRootSignature.cpp`.
Note: we don't add a backend test for version 1.0 flag values because it
treats the struct as though there is no flags value. However, this will
be used when we use the verifications in the frontend.
- Updates `verifyDescriptorFlag` to check for valid flags based on
version, as reflected [here](https://github.com/llvm/wg-hlsl/pull/297)
- Add test to demonstrate updated flag verifications
- Adds `verifyNumDescriptors` to the validation of `DescriptorRange`s
- Add a test to demonstrate `numDescriptors` verification
- Updates a number of tests that mistakenly had an invalid
`numDescriptors` specified
Resolves: https://github.com/llvm/llvm-project/issues/147107
|
|
`RootSignatureValidations` (#147117)
This pr abstracts out the logic of detecting resource range overlap from
`SemaHLSL` into the `RootSignatureValidations` library.
For more context see linked issue.
- Moves the validation logic from `SemaHLSL` to
`RootSignatureValidations`
- Updates `SemaHLSL` to use the new interface for the validations
Resolves: https://github.com/llvm/llvm-project/issues/146393
|
|
`RootSignatureValidations` library (#147110)
Simple code movement of the verification logic in `validate` of the
`DXILRootSignature` pass.
Moving this code to the `RootSignatureValidations` library allows for
the common verifications to be used in the frontend.
- Moves all the `static` verification functions used in
`DXILRootSignature` to the `RootSignatureValidations` library
- Updates `DXILRootSignature` to use the moved functions
Resolves: https://github.com/llvm/llvm-project/issues/145940
|
|
This pr breaks-up `HLSLRootSignatureUtils` into separate orthogonal and
meaningful libraries. This prevents it ending up as a dumping grounds of
many different parts.
- Creates a library `RootSignatureMetadata` to contain helper functions
for interacting the root signatures in their metadata representation
- Create a library `RootSignatureValidations` to contain helper
functions that will validate various values of root signatures
- Move the serialization of root signature elements to
`HLSLRootSignature`
Resolves: https://github.com/llvm/llvm-project/issues/145946
|
