Skip to content

Commit

Permalink
test passing a const uninitialized_buffer to launch
Browse files Browse the repository at this point in the history
  • Loading branch information
ericniebler committed Sep 2, 2024
1 parent ac98ee7 commit e696cfd
Showing 1 changed file with 26 additions and 5 deletions.
31 changes: 26 additions & 5 deletions cudax/test/containers/uninitialized_buffer.cu
Original file line number Diff line number Diff line change
Expand Up @@ -168,13 +168,34 @@ __global__ void kernel(_CUDA_VSTD::span<int> data)
data[0] = 42;
}

__global__ void const_kernel(_CUDA_VSTD::span<const int> data)
{
// Touch the memory to be sure it's accessible
CUDAX_CHECK(data.size() == 1024);
}

TEST_CASE("uninitialized_buffer is usable with cudax::launch", "[container]")
{
const int grid_size = 4;
cudax::uninitialized_buffer<int, ::cuda::mr::device_accessible> buffer{cuda::mr::device_memory_resource{}, 1024};
auto dimensions = cudax::make_hierarchy(cudax::grid_dims(grid_size), cudax::block_dims<256>());
SECTION("non-const")
{
const int grid_size = 4;
cudax::uninitialized_buffer<int, ::cuda::mr::device_accessible> buffer{cuda::mr::device_memory_resource{}, 1024};
auto dimensions = cudax::make_hierarchy(cudax::grid_dims(grid_size), cudax::block_dims<256>());

cudax::stream stream;
cudax::stream stream;

cudax::launch(stream, dimensions, kernel, buffer);
cudax::launch(stream, dimensions, kernel, buffer);
}

SECTION("const")
{
const int grid_size = 4;
const cudax::uninitialized_buffer<int, ::cuda::mr::device_accessible> buffer{
cuda::mr::device_memory_resource{}, 1024};
auto dimensions = cudax::make_hierarchy(cudax::grid_dims(grid_size), cudax::block_dims<256>());

cudax::stream stream;

cudax::launch(stream, dimensions, const_kernel, buffer);
}
}

0 comments on commit e696cfd

Please sign in to comment.