[NATIVECPU] Emit Native CPU properties #19429
Conversation
uwedolinsky
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
clang/lib/CodeGen/BackendUtil.cpp
Outdated
| if (LangOpts.SYCLIsNativeCPU) | ||
| llvm::sycl::utils::addSYCLNativeCPUEarlyPasses(MPM); |
There was a problem hiding this comment.
As discussed internally, ideally we'd get rid of both LangOpts.SYCLIsNativeCPU and SYCLNativeCPUBackend and check what target we're building for, but this is blocked on #19344, without that we don't have enough information.
There was a problem hiding this comment.
Since #19344 is now merged I've updated the PR and removed the usage of LangOpts.SYCLIsNativeCPU.
uwedolinsky
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
uwedolinsky
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
| @@ -149,7 +149,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch( | |||
| bool isLocalSizeOne = | |||
| ndr.LocalSize[0] == 1 && ndr.LocalSize[1] == 1 && ndr.LocalSize[2] == 1; | |||
| if (isLocalSizeOne && ndr.GlobalSize[0] > numParallelThreads && | |||
| !kernel->hasLocalArgs()) { | |||
| !kernel->hasLocalArgs() && !hKernel->isNDRangeKernel()) { | |||
| // TODO: Check if !kernel->hasLocalArgs() is needed | |||
There was a problem hiding this comment.
You made the logic say that if it has local args, it's automatically necessary an NDRangeKernel. I'm not sure yet whether that's 100% accurate, but if it is, hasLocalArgs() does become redundant.
uwedolinsky
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
uwedolinsky
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
uwedolinsky
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
uwedolinsky
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
uwedolinsky
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
uwedolinsky
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
uwedolinsky
changed the title
uwedolinsky
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
uwedolinsky
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
uwedolinsky
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
uwedolinsky
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
uwedolinsky
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
Both wrappers now call the same function for NativeCPU. More (non-NativeCPU) functions could probably be shared between the wrappers if need be. |
Extends
clang-offload-wrapperandSYCLOffloadWrapper(clang-linker-wrapper), enabling adding kernel properties that are specific only to Native CPU. Adds a compiler pass that checks whether a kernel comes from asycl::nd-rangeand adds a Native CPU - only property for it.This new "kernel property infrastructure" will be extended in the future to encode other kernel capabilities. For example the applied vector width which determines how many workgroups could be executed in one kernel invocation. This could help making the kernel launches more efficient. Another property could encode whether the kernel supports peeling - without peeling the kernel would have less branches and the NativeCPU adapter could schedule peeling invocations (of the scalar kernel) in separate threads which might benefit some performance scenario.
This PR replaces and extends #16152 which still used the old UR api/repo.
This PR also adds testing to ensure the new clang-linker-wrapper integration produces the expected IR on SYCL code.
Update: Since #19550 got merged which removed the unsafe optimization this PR was also trying to make safe, the scope of this PR has now been reduced to emitting properties. The optimization (processing more work items per kernel invocation) removed by #19550 will be added back in a separate PR.