NVIDIA / NVIDIA/cccl

Restrict in-place execution

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

Description

There's a blind spot in the Thrust/CUB in-place execution guarantees that I believe should be addressed. Thrust/CUB allow iterators to point to the same memory meanwhile there's no restriction on these iterators value types.

For instance, it's legal to reinterpret iterators as pointing to char. The following code illustrates a useful scenario of converting masks into offsets. It's not at all clear from the documentation why this is not allowed:

```cpp
const bool alias = true;
const int nc = sizeof(std::size_t) * 512 * 1024;
const int ns = nc / sizeof(std::size_t);

char *d_in{};
cudaMalloc(&d_in, nc);
cudaMemset(d_in, 1, nc);

std::size_t *d_out{};

if (alias)
{
d_out = reinterpret_cast(d_in);
}
else
{
cudaMalloc(&d_out, nc);
cudaMemset(d_out, 1, nc);
}

cub::Sum op{};
std::size_t init{};

std::size_t temp_storage_bytes{};
std::uint8_t *d_temp_storage{};

cub::DeviceScan::ExclusiveScan(
d_temp_storage, temp_storage_bytes,
d_in, d_out, op, init, ns);

cudaMalloc(&d_temp_storage, temp_storage_bytes);

cub::DeviceScan::ExclusiveScan(
d_temp_storage, temp_storage_bytes,
d_in, d_out, op, init, ns);

std::vector h_out(ns);
cudaMemcpy(h_out.data(), d_out, nc, cudaMemcpyDeviceToHost);

for (int i = 0; i < ns; i++)
{
if (h_out[i] != i)
{
std::cout << "case 1 out[" << i << "] = " << h_out[i] << " != " << i << std::endl;
break;
}
}
```

The code above introduces a data race leading to incorrect result in the case of aliasing. The more casual example would be usage of unions:

```cpp
union U
{
char in[sizeof(std::size_t) * 512 * 1024];
std::size_t out[sizeof(in) / sizeof(std::size_t)];
};

const std::size_t nc = sizeof(U::in);
const std::size_t ns = sizeof(U::out) / sizeof(std::size_t);

U *d_u{};
cudaMalloc(&d_u, sizeof(U));

char *d_in = d_u->in;
std::size_t *d_out = d_u->out;

cudaMemset(d_in, 1, nc);

cub::Sum op{};
std::size_t init{};

std::size_t temp_storage_bytes{};
std::uint8_t *d_temp_storage{};

cub::DeviceScan::ExclusiveScan(
d_temp_storage, temp_storage_bytes,
d_in, d_out, op, init, ns);

cudaMalloc(&d_temp_storage, temp_storage_bytes);

cub::DeviceScan::ExclusiveScan(
d_temp_storage, temp_storage_bytes,
d_in, d_out, op, init, ns);

std::vector h_out(ns);
cudaMemcpy(h_out.data(), d_out, nc, cudaMemcpyDeviceToHost);

for (int i = 0; i < ns; i++)
{
if (h_out[i] != i)
{
std::cout << "case 2 out[" << i << "] = " << h_out[i] << " != " << i << std::endl;
break;
}
}

cudaFree(d_temp_storage);
cudaFree(d_u);
```
This also leads to the previous issues. Thrust inherits the very same issues.

Since the change of documentation might be considered as a breaking change, I'd like to get this addressed in 2.0.

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.