Is cuda::std::bitset designed to be thread-safe?

[SagaciousFish]

I might have encountered data racing during set().
It seems that during massive set() calls, some results set before would return to unset.
The following script shows it.

__global__ void func(
	cuda::std::bitset<SOME_SIZE>* bitset_ptr,
	size_t pos
) {
	...
	size_t SOME_POS = 6;
	bitset_ptr->set(dst);
	if (!bitset_ptr->test(dst)) {
	    printf("dst did not set\n");
	}
	if (dst == SOME_POS) {
	    if (!bitset_ptr->test(SOME_POS)) {
	        printf("SOME_POS did not set\n");
	    }
	    printf("check SOME_POS marked: %s\n",
	           curr_frontier_bitset_ptr->test(SOME_POS) ? "true" : "false"
	    );
	}
}

The output could be (shown in my runs):

dst did not set
SOME_POS did not set
check SOME_POS marked: false

or

SOME_POS did not set
check SOME_POS marked: false

Even after checking the source code, I could not see a good lock-free method to avoid the racing condition on user side. Is there any good practice on this?

[miscco]

Unfortunately, this is not unexpected.

We cannot add internal locking that would negatively affect performance for common use cases

[SagaciousFish]

What if we have a std::byte chunk array and use atomic updates? One major use case for a bitset is bool array with smaller memory usage. If thread-safety is not guaranteed, then what's the use of this in device programs?

[miscco]

Our data structures especially those in cuda::std:: are meant to provide access to familiar C++ features but are still standard conforming. They are also usable on both host and device and can be copied around.

Using them does not prevent any data races that would still happen with std:: data structures or even a plain bool array. That is on the user to ensure that there are no racy operations

[SagaciousFish]

Thank you for your answer