Skip to content

Enable more adaptive cpp - #1247

Open
TApplencourt wants to merge 7 commits into
mainfrom
enable-more-adaptive-cpp
Open

TApplencourt wants to merge 7 commits into
mainfrom
enable-more-adaptive-cpp

Conversation

@TApplencourt

Copy link
Copy Markdown
Contributor
  • Bump acpp
  • Enable more tests (who both compile and run) -- minus the id who have a PR open. But I guess doesn't matter, we just want the CTS-CI to pass, and the CTS doesn't run acpp.

TApplencourt and others added 7 commits October 6, 2026 23:12
Several test categories were disabled for AdaptiveCpp with FIXMEs naming
features that are in fact implemented. Verified on AdaptiveCpp 25.10.0
(NVIDIA A100, CUDA backend) by building and running each category.

local_accessor: was 23/23 test cases compile-time disabled. The api,
linearization and access_among_work_items tests now pass in full (1095,
10 and 15 assertions respectively). The constructors test still fails, so
its guard is kept.

optional_kernel_features: was 12/12 test cases compile-time disabled. The
device_has tests whose kernels carry no attribute now pass (216
assertions), as does the speculative compilation test (66 assertions).
Guards are kept for the cases that rely on SYCL kernel attributes:
AdaptiveCpp parses [[sycl::device_has()]], [[sycl::reqd_work_group_size]]
and [[sycl::reqd_sub_group_size]] but ignores them, so the exception the
tests require is never thrown. The FIXMEs are updated to say so.

platform_info, queue_info: the guards claimed struct information
descriptors were missing, but every descriptor these tests query is
available. Enabling them adds passing assertions (platform 28 -> 36,
queue 490 -> 494) with no new failures.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
The pinned revision fcd26a61 dates from 2025-10-24 and is not a release;
it sits five commits before the v25.10.0 tag. Move to a current develop
revision, which is what the newly enabled tests were verified against.

Among the changes picked up is a refactor of the group algorithms to
work in place (#1903), covering inclusive_scan, exclusive_scan, the
joint_ variants and group_broadcast, plus a sub_group fix so that
get_local_range() reports the real size of an incomplete sub-group
instead of always returning the warp size. These affect test categories
that are still guarded for AdaptiveCpp, so testing against a current
revision is needed to tell which of those guards can be lifted.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
The joint scan and scan-over-group tests were disabled with
DISABLED_FOR_TEST_CASE(AdaptiveCpp). Verified on AdaptiveCpp
25.10.0+git.5f94ef59 (NVIDIA A100, CUDA backend).

Both test cases carried a second, inner layer of AdaptiveCpp workarounds
inside the test body: an #if that substituted same-type-only variants of
the joint scan combinations, and WARN messages describing further
limitations. That code was unreachable. DISABLED_FOR_TEST_CASE replaces
the whole body with an auto-failing stub, so under AdaptiveCpp the inner
branches could never run. Both layers were added in the same commit
(fbb2ca7), so the workarounds had never executed.

With the outer guard lifted, the inner claims turn out not to hold:
the full mixed-type combinations pass, taking the category from 2593 to
22168 assertions. The WARN about joint scans not working over several
sub-groups has no code behind it at all. Those are removed.

What does not work is scan without an init value: exclusive_scan_over_group
returns a zero-initialized value rather than the operator's identity
(0 instead of -inf for sycl::maximum on float), so that test case keeps
its guard, with the FIXME updated to describe the actual defect. The
variant taking an explicit init passes in full, which also disproves the
note about inclusive_scan_over_group taking init and op in the wrong
order.

Category result: 91 test cases, 88 passed, 3 failed, where all three
failures are the test case still deliberately disabled above.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
The two sub-group checks in the joint reduce tests were replaced by an
unconditional `res_acc[1] = true` for AdaptiveCpp, on the grounds that it
had no implementation over sub-groups. It does: the ASSERT_RETURN_TYPE
immediately above each stub already called
joint_reduce(sub_group, ...) unconditionally, so only the value
comparison was being skipped.

Running the real comparison passes: 540 assertions over 15 joint reduce
test cases, and the category as a whole is unchanged at 88 of 91 test
cases passing, the three failures being a test case disabled elsewhere.
To confirm the restored assertion is actually evaluated rather than
vacuously true, it was temporarily made to compare against a wrong
value, which produced 54 additional failures.

Verified on AdaptiveCpp 25.10.0+git.5f94ef59 (NVIDIA A100, CUDA backend).

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Drop the stale known_identity FIXMEs and WARN from group_joint_scan: both
test cases are enabled now, so the comments described a disable that no
longer exists.

Restore the original one-line wording on the remaining
group_scan_over_group guard, and give the two still-disabled
device_has_exceptions cases a one-line reason of their own now that the
file-wide errc guard they used to sit behind is gone.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
vec deduction guides are implemented; the test passes on AdaptiveCpp
develop (80 assertions, A100, --acpp-targets=generic).

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
All 13 guards in tests/id/id.cpp are lifted. The suite passes in full
(89 test cases, 2960 assertions, A100, --acpp-targets=generic).

The 9 cases that previously failed (`id OP bool`, `id OP char16_t`,
`id OP char32_t`) were blocked by a bug in AdaptiveCpp, not by a missing
feature: HIPSYCL_ID_BINARY_OP_OUT_OF_PLACE_SIZE_T in
hipSYCL/sycl/libkernel/id.hpp casts each result element to T, the operand
type, rather than std::size_t, so a narrow T truncates the result. With
T = bool every result collapses to 0 or 1.

SYCL 2020 4.9.1.3 Table 85 specifies the result as an id, whose
components are std::size_t (Table 84), and the only conversion it
prescribes runs the other way: "If the operator returns a bool the result
is then cast to std::size_t". The sibling macro for the id-op-id form
already casts to std::size_t.

NOTE: requires the AdaptiveCpp fix; CI will fail until it lands upstream
and the pinned version in .github/workflows/cts_ci.yml is bumped past it.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
@TApplencourt
TApplencourt requested a review from a team as a code owner October 7, 2026 17:42
@TApplencourt

Copy link
Copy Markdown
Contributor Author

@illuhad \o/ lot of new passes!

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant