NVIDIA / NVIDIA/cccl

Misaligned access in `thrust::mr::new_delete_resource`

Open
#726 0 comments 0 reactions 0 assignees View on GitHub
thrust
Dominant language
C++
Stars
2.5k
Forks
487
Avg merge
2d 7h
Merged PRs (30d)
296

Description

The `thrust::mr::new_delete_resource` implements an aligned allocator by storing the number of offset bytes from a root allocation in order to satisfy the alignment request.

For an allocation request of `n` bytes, it stores the offset at `p + n`.

https://github.com/thrust/thrust/blob/master/thrust/mr/new.h#L56

However, if `n % sizeof(std::size_t) != 0`, then this will end up with a misaligned store.

Similarly, a misaligned load will occur in `do_deallocate` here: https://github.com/thrust/thrust/blob/master/thrust/mr/new.h#L70

While most common CPU archs can handle this (potentially with performance loss), it is undefined behavior according to the C++ standard.

The implementation should be modified such that the offset is stored at an aligned location. One solution is to store it at an aligned location _after_ the allocation. Alternatively, it could be stored in the 8B immediately _preceding_ the returned allocation pointer.

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.