cuda::(counting/binary)_semaphore unusable in __shared__ memory

Attempting to use a `__shared__ cuda::(counting/binary)_sempahore` will result in a hard compile error.

```
error: initializer not allowed for __shared__ variable
```

https://godbolt.org/z/68Wq3Eon1

In contrast, a `__shared__ cuda::barrier` only issues a warning because `__shared__` variables are not initialized:

```
warning: dynamic initialization is not supported for a function-scope static
__shared__ variable within a __device__/__global__ function
```

This warning can be suppressed with `#pragma diag_suppress static_var_with_dynamic_init`.

The root cause of the `cuda::(counting/binary)_semaphore` issue is that it doesn't have a trivial default constructor. Instead, it has a constructor with a defaulted parameter:
https://github.com/NVIDIA/libcudacxx/blob/ec5727daab0b21878870e72cc62988bd352d85e3/include/cuda/std/semaphore#L33

Simpler repro: https://godbolt.org/z/a3fjWMvas

**The easy solution is to just follow the pattern of `cuda::barrier` and add an explicitly defaulted constructor with a friend `init` function.**

Start in include/cuda/std/semaphore and compare its constructor and initialization pattern with cuda::barrier. Use the linked Godbolt repro to verify that a __shared__ cuda::counting_semaphore or cuda::binary_semaphore compiles without the hard initializer error; the issue is done when shared-memory semaphore declarations are usable.

Written by the indexing model from the issue text.