NVIDIA / NVIDIA/cccl

[BUG]: `cuda::dynamic_shared_memory(config)` may returned misaligned pointer

Open
#7,867 0 comments 0 reactions 0 assignees View on GitHub
Dominant language
C++
Stars
2.5k
Forks
486
Avg merge
2d 6h
Merged PRs (30d)
295

Description

### Is this a duplicate?

- [x] I confirmed there appear to be no [duplicate issues](https://github.com/NVIDIA/cccl/issues) for this bug and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)

### Type of Bug

Silent Failure

### Component

libcu++

### Describe the bug

When using `cuda::launch` with dynamic shared memory, we allocate the proper size, but we don't take into account the alignment.

In `cuda::dynamic_shared_memory(config)`, we do:
```cpp
{
extern __shared__ unsigned char dyn_smem[];
return *reinterpret_cast(dyn_smem);
}
```
but there is no guarantee that `dyn_smem` will have the right alignment.

### How to Reproduce

no-repro

### Expected behavior

`cuda::dynamic_shared_memory(config)` should return a view to properly aligned memory address.

### Reproduction link

_No response_

### Operating System

_No response_

### nvidia-smi output

_No response_

### NVCC version

_No response_

Contributor guide

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.