Ensure that we can use cuda::std::optional with types that are not __host__ __device__
#1663
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
miscco
force-pushed
the
exec_check_monads
branch
from
22c1418 to
868f28d
Compare
|
I am currently only defining |
miscco
force-pushed
the
exec_check_monads
branch
from
868f28d to
8d5519b
Compare
jrhemstad
reviewed
Apr 24, 2024
| } | ||
| }; | ||
|
|
||
| __global__ void ensure_exec_check_is_disabled() |
There was a problem hiding this comment.
question: Do we need to define a separate __global__ function? The fake_main() thing already makes main() a __global__ function, right?
Can't we do something like:
int main(){
NV_IF_TARGET(NV_IS_DEVICE, assert(optional<DeviceOnly>{} == DeviceOnly{}));
}
?
There was a problem hiding this comment.
I do not really understand why, but that is what I initially did and it did not emit any errors.
Probably the same that using global objects like cuda::std::nullopt only triggers errors in a kernel
There was a problem hiding this comment.
Could be because main() actually expands to a __host__ __device__ function which is then invoked from a __global__. That could interfere with if/when nvcc emits a diagnostic.
cccl/libcudacxx/test/libcudacxx/force_include.h
Lines 49 to 57 in e3d181f
miscco
force-pushed
the
exec_check_monads
branch
from
8d5519b to
7e0adb0
Compare
7 tasks
miscco
added
feature request
🟨 CI Results [ Failed: 1 | Passed: 301 | Total: 302 ]
|
| # | Runner |
|---|---|
| 232 | linux-amd64-cpu16 |
| 28 | linux-amd64-gpu-v100-latest-1 |
| 24 | linux-arm64-cpu16 |
| 18 | windows-amd64-cpu16 |
👃 Inspect Changes
Modifications in project?
| Project | |
|---|---|
| CCCL Infrastructure | |
| +/- | libcu++ |
| CUB | |
| Thrust |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| +/- | libcu++ |
| +/- | CUB |
| +/- | Thrust |
1 similar comment
🟨 CI Results [ Failed: 1 | Passed: 301 | Total: 302 ]
|
| # | Runner |
|---|---|
| 232 | linux-amd64-cpu16 |
| 28 | linux-amd64-gpu-v100-latest-1 |
| 24 | linux-arm64-cpu16 |
| 18 | windows-amd64-cpu16 |
👃 Inspect Changes
Modifications in project?
| Project | |
|---|---|
| CCCL Infrastructure | |
| +/- | libcu++ |
| CUB | |
| Thrust |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| +/- | libcu++ |
| +/- | CUB |
| +/- | Thrust |
🟩 CI Results [ Failed: 0 | Passed: 302 | Total: 302 ]
|
| # | Runner |
|---|---|
| 232 | linux-amd64-cpu16 |
| 28 | linux-amd64-gpu-v100-latest-1 |
| 24 | linux-arm64-cpu16 |
| 18 | windows-amd64-cpu16 |
👃 Inspect Changes
Modifications in project?
| Project | |
|---|---|
| CCCL Infrastructure | |
| +/- | libcu++ |
| CUB | |
| Thrust |
Modifications in project or dependencies?
| Project | |
|---|---|
| CCCL Infrastructure | |
| +/- | libcu++ |
| +/- | CUB |
| +/- | Thrust |
bernhardmgruber
approved these changes
May 7, 2024
| _LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX20 ~__expected_destruct() | ||
| { | ||
| if (__has_val_) | ||
| { | ||
| _CUDA_VSTD::__destroy_at(_CUDA_VSTD::addressof(__union_.__val_)); |
There was a problem hiding this comment.
Question: there is no advantage of destroy_at over calling the dtor explicitely? Both work during constant evaluation, if the dtor is constexpr. destroy_at` only has a convenience advantage because we do not need to spell the destroyed type's dtor, right?
There was a problem hiding this comment.
Right, but it has the disadvantage of an additional function instantiation and that it will only work with std::destroy_at during compile time as that is the standard blessed function
wmaxey
approved these changes
May 8, 2024
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
We do not want to force users to make everything
__host__ __device__so allow types that are either host only or device only with the caveat that the users must ensure that they are used correctly