develop branch broken on oneAPI
#2124
Comments
|
As a cross check, I have different failures using
|
fwyzard
changed the title
develop branch broken on oneAPI GPUdevelop branch broken on oneAPI
|
The CPU tests were already broken, but they seem in much worse shape now: In particular Tested with |
|
The commit that has broken the tests is this, related to #2093. Before this comment:
After it:
Tested both on Intel(R) Data Center GPU Flex 170 and Intel(R) Data Center GPU Max 1100 with Intel(R) oneAPI DPC++/C++ Compiler 2023.2.0 with: or |
I do not see it in the tests, but I have seen pixeltrack failing in the same way after this commit |
|
Thx for reporting this! From a visual inspection of the diff in PR #2093 there is a clear problem: No changes were made to the SYCL backend :) And since the change did not break API, the backend kept compiling. I will put up a quick PR with the required changes. However, I don't have a machine to test at the moment, since I am out of town. |
|
Ok thanks! I can run the tests on the Intel GPU if needed :) |
These are changes missed as part of alpaka-group#2093. Fixes: alpaka-group#2124
These are changes missed as part of alpaka-group#2093. Fixes: alpaka-group#2124
|
Coming up: #2125. The CI says the changes compile for the SYCL backend. I have not run any tests though. Please give it a try, thx! |
|
|
|
The current HEAD is better, but still has issues in debug mode (tested on Ponte Vecchio gpu): |
|
Regarding are used to build the tests. I've not looked into |
This fixes a compilation error with the SYCL backend in debug mode. Fixes: alpaka-group#2124
This fixes a compilation error with the SYCL backend in debug mode. Fixes: alpaka-group#2124
This fixes a compilation error with the SYCL backend in debug mode. Fixes: #2124
|
Reopening for the warp test. |
|
With |
|
Great! That bug was fixed in #2127, but good to know that the CI can save us in the future now! |
no idea why, but it is broken since this commit |
|
@j-stephan please have a look at the remaining |
|
I've identified the issue but have no idea how to solve it: The behaviour of SYCL's sub-group (= warp) functionality differs between template<typename TDim>
struct Activemask<warp::WarpGenericSycl<TDim>>
{
static auto activemask(warp::WarpGenericSycl<TDim> const& warp) -> std::uint32_t
{
auto const sub_group = warp.m_item_warp.get_sub_group();
auto const mask = sycl::ext::oneapi::group_ballot(sub_group, true);
std::uint32_t bits = 0;
mask.extract_bits(bits);
return bits;
}
};always returns |
|
Does this mean that |
|
I'm not sure. I think technically it is used the wrong way here since But then again, the |
I was searching for some documentation for that! Because I was thinking whether we may just trigger some undefined behavior here and
Well, if the vendor can hack, maybe so can we? :) |
|
Well, there you have it: https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:group-functions
Seems like we cannot implement |
|
I propose to either make |
|
I would favour |
* Switch to 1.0.0-rc1 * Silence more nvcc warnings * Add missing ALPAKA_UNREACHABLE * Silence more warnings Co-authored-by: Bernhard Manfred Gruber <[email protected]> * Disable activemask for SYCL Fixes: #2124 * refactor template order `allocMappedBufIfSupported` Move template `TPlatform` as the last template. There is no need to provide the platform template signature if we pass the platform as an instance. * Remove unused aliases * Only add clang warning flag when supported * Fix a warning with nvcc * Workaround gcc warning on uninitialized PlatformCpu * Fix warnings by clang * Suppress clang warnings in nvcc generated code * Pass alpaka_ENABLE_WERROR from environment to CMake * Disable GCC warning in nvcc generated code Fixes the following warning: ``` /builds/hzdr/crp/alpaka/test/unit/math/src/Defines.hpp:52:35: error: left operand of comma operator has no effect [-Werror=unused-value] 52 | for(size_t i = 0; i < argsItem.arity_nr; ++i) | ~~~~~~~~~^~~~~~~~~~ ``` * Workaround gcc warning on uninitialized PlatformUniformCudaHipRt * Fix OpenMP 5.1 atomics * Workaround clang not recognizing ternay expression * Implement atomicInc/atomicDec via omp critical Fixes: #2170 * Add clang-17 to CI Fixes: #2169 * Rename lambda captures to workaround warnings clang warns that the names of the captures shadow the outside variables. E.g.: QueueUniformCudaHipRt.hpp:215:57: error: declaration shadows a local variable [-Werror,-Wshadow-uncaptured-local] 215 | auto f = queue.m_callbackThread.submit([data = std::move(data)] { data->t(); }); | ^ * Exclude clang CUDA Debug builds from the CI They fail with: ptxas /tmp/randomStrategies-sm_61-768a17.s, line 14415; fatal : Parsing error near '.': syntax error ptxas fatal : Ptx assembly aborted due to errors clang++: error: ptxas command failed with exit code 255 (use -v to see invocation) * Add clang-17 to README.md CUDA/HIP/SYCL is marked as untested for now. * Fix clang-format version in docs (#2176) * Remove unnecessary -fintelfpga flag * Add changelog for alpaka v1.0.0 * Update author lists Co-authored-by: Bernhard Manfred Gruber <[email protected]> * Set release date in changelog --------- Co-authored-by: Jan Stephan <[email protected]> Co-authored-by: René Widera <[email protected]>
(initially reported by @AuroraPerego)
Hi,
looks like the current HEAD of the
developbranch is broken for the SYCL/oneAPI GPU target:This is one an Intel Data Center GPU Max 1100 (Ponte Vecchio) with oneAPI DPC++/C++ Compiler 2023.2.0 (or 2023.2.1, there seems to be some confusion with minor versions):
memBufTestzeroDimBufferTestfails to build withparallelLoopPatternsparallelLoopPatternsfails at runtime with an assertion:bufSlicingTestbufSlicingTestfails at runtime with an assertion:memViewTestmemViewTestfails at runtime with the same assertion:randomCells2DrandomCells2Dfails at runtime:matMulTestmatMulTestfails at runtime:warpTestwarpTestfails at runtime:The text was updated successfully, but these errors were encountered: