[SYCL][NATIVECPU] Support specialization constants on Native CPU #14446
Conversation
| @@ -932,6 +934,7 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, | |||
| unsigned Size = M.getDataLayout().getTypeStoreSize(SCTy); | |||
| uint64_t Align = M.getDataLayout().getABITypeAlign(SCTy).value(); | |||
|
|
|||
|
|
|||
| @@ -946,8 +949,14 @@ PreservedAnalyses SpecConstantsPass::run(Module &M, | |||
| updatePaddingInLastMDNode(Ctx, SCMetadata, Padding); | |||
| } | |||
|
|
|||
| SCMetadata[SymID] = generateSpecConstantMetadata( | |||
| M, SymID, SCTy, NextID, /* is native spec constant */ false); | |||
| if (sycl::utils::isSYCLNativeCPU(M) && isa<StructType>(DefaultValue->getType())) { | |||
There was a problem hiding this comment.
If using default value is a reliable enough method, then I suggest to simply use it for all targets, simplifying this code.
There was a problem hiding this comment.
I've changed this to always pass the default value type to generateSpecConstantMetadata, unfortunately we can't use the default type for everything because replaceAllUsesWith will fail since due to the ABI the type returned by the spec const builtin is different from the default type.
| SCMetadata[SymID] = generateSpecConstantMetadata( | ||
| M, SymID, SCTy, NextID, /* is native spec constant */ false); | ||
| if (sycl::utils::isSYCLNativeCPU(M) && isa<StructType>(DefaultValue->getType())) { | ||
| auto STy = cast<StructType>(DefaultValue->getType()); |
There was a problem hiding this comment.
Why is this cast needed? generateSpecConstantMetadata accepts Type *
There was a problem hiding this comment.
You are right, the cast is removed now since we always use the default type, thank you
| @@ -12,6 +12,7 @@ | |||
| // FIXME: ACC devices use emulation path, which is not yet supported | |||
| // UNSUPPORTED: accelerator | |||
| // UNSUPPORTED: hip | |||
| // UNSUPPORTED: native_cpu | |||
There was a problem hiding this comment.
Why is this test being marked as unsupported on native CPU if this PR is intended to add support for this feature? Does it mean that the test passed and this patch introduces a regression?
There was a problem hiding this comment.
We run Native CPU e2e tests downstream, we weren't running the tests in test-e2e/SpecConstants before this PR. With this PR all the tests in test-e2e/SpecConstants are passing, besides this one, which is also marked as UNSUPPORTED for other targets that use emulated SC, so this is not a regression.
PietroGhg
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
PietroGhg
had a problem deploying
to
WindowsCILock
GitHub Actions
Failure
PietroGhg
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
|
Hi @AlexeySachkov, do you have any more thoughts about this PR? |
omarahmed1111
force-pushed
the
pietro/native_cpu_specconstants
branch
from
29be54d to
a77b395
Compare
omarahmed1111
force-pushed
the
pietro/native_cpu_specconstants
branch
from
a77b395 to
d01810c
Compare
omarahmed1111
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
omarahmed1111
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
|
Hi @intel/llvm-gatekeepers, this looks ready to merge, thank you |
Adds support to specialization constants by scheduling the
SpecConstantsPassin the Native CPU pass pipeline. A small change is needed in that pass since, due to ABI, the return type of the spec const builtins can't be used to determine the type of the spec constant, so for Native CPU we use the default value instead.UR PR: oneapi-src/unified-runtime#1821