Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Re-use device buffer in yuy2 cuda helper #11737

Open
wants to merge 1 commit into
base: development
Choose a base branch
from
+64 −14
Open

Re-use device buffer in yuy2 cuda helper #11737

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
30 changes: 16 additions & 14 deletions src/cuda/cuda-conversion.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -233,10 +233,12 @@ void rscuda::unpack_yuy2_cuda_helper(const uint8_t* h_src, uint8_t* h_dst, int n

// How many super pixels do we have?
int superPix = n / 2;
std::shared_ptr<uint8_t> d_dst;
std::shared_ptr<uint8_t> d_src = alloc_dev<uint8_t>(superPix * 4);

auto result = cudaMemcpy(d_src.get(), h_src, superPix * sizeof(uint8_t) * 4, cudaMemcpyHostToDevice);
static DeviceBuffer<uint8_t> d_dst;
static DeviceBuffer<uint8_t> d_src;
d_src.reserve(superPix, 4);

auto result = cudaMemcpy(d_src.data(), h_src, superPix * sizeof(uint8_t) * 4, cudaMemcpyHostToDevice);
assert(result == cudaSuccess);

int numBlocks = superPix / RS2_CUDA_THREADS_PER_BLOCK;
Expand All @@ -253,28 +255,28 @@ void rscuda::unpack_yuy2_cuda_helper(const uint8_t* h_src, uint8_t* h_dst, int n
*/
case RS2_FORMAT_Y16:
size = 2;
d_dst = alloc_dev<uint8_t>(n * size);
kernel_unpack_yuy2_y16_cuda << <numBlocks, RS2_CUDA_THREADS_PER_BLOCK >> > (d_src.get(), d_dst.get(), superPix);
d_dst.reserve(n * size);
kernel_unpack_yuy2_y16_cuda << <numBlocks, RS2_CUDA_THREADS_PER_BLOCK >> > (d_src.data(), d_dst.data(), superPix);
break;
case RS2_FORMAT_RGB8:
size = 3;
d_dst = alloc_dev<uint8_t>(n * size);
kernel_unpack_yuy2_rgb8_cuda << <numBlocks, RS2_CUDA_THREADS_PER_BLOCK >> > (d_src.get(), d_dst.get(), superPix);
d_dst.reserve(n * size);
kernel_unpack_yuy2_rgb8_cuda << <numBlocks, RS2_CUDA_THREADS_PER_BLOCK >> > (d_src.data(), d_dst.data(), superPix);
break;
case RS2_FORMAT_BGR8:
size = 3;
d_dst = alloc_dev<uint8_t>(n * size);
kernel_unpack_yuy2_bgr8_cuda << <numBlocks, RS2_CUDA_THREADS_PER_BLOCK >> > (d_src.get(), d_dst.get(), superPix);
d_dst.reserve(n * size);
kernel_unpack_yuy2_bgr8_cuda << <numBlocks, RS2_CUDA_THREADS_PER_BLOCK >> > (d_src.data(), d_dst.data(), superPix);
break;
case RS2_FORMAT_RGBA8:
size = 4;
d_dst = alloc_dev<uint8_t>(n * size);
kernel_unpack_yuy2_rgba8_cuda << <numBlocks, RS2_CUDA_THREADS_PER_BLOCK >> > (d_src.get(), d_dst.get(), superPix);
d_dst.reserve(n * size);
kernel_unpack_yuy2_rgba8_cuda << <numBlocks, RS2_CUDA_THREADS_PER_BLOCK >> > (d_src.data(), d_dst.data(), superPix);
break;
case RS2_FORMAT_BGRA8:
size = 4;
d_dst = alloc_dev<uint8_t>(n * size);
kernel_unpack_yuy2_bgra8_cuda << <numBlocks, RS2_CUDA_THREADS_PER_BLOCK >> > (d_src.get(), d_dst.get(), superPix);
d_dst.reserve(n * size);
kernel_unpack_yuy2_bgra8_cuda << <numBlocks, RS2_CUDA_THREADS_PER_BLOCK >> > (d_src.data(), d_dst.data(), superPix);
break;
default:
assert(false);
Expand All @@ -284,7 +286,7 @@ void rscuda::unpack_yuy2_cuda_helper(const uint8_t* h_src, uint8_t* h_dst, int n

cudaDeviceSynchronize();
Copy link
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

In my opinion, this is unnecessary, but better to use async copy and use the same stream across the whole function and do a stream syncronization before we return from this function. By that we can avoid extra syncing, however the gain would be probably not very significant.

https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__MEMORY.html#group__CUDART__MEMORY_1g85073372f776b4c4d5f89f7124b7bf79


result = cudaMemcpy(h_dst, d_dst.get(), n * sizeof(uint8_t) * size, cudaMemcpyDeviceToHost);
result = cudaMemcpy(h_dst, d_dst.data(), n * sizeof(uint8_t) * size, cudaMemcpyDeviceToHost);
assert(result == cudaSuccess);

/* cudaEventRecord(stop);
Expand Down
48 changes: 48 additions & 0 deletions src/cuda/rscuda_utils.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -15,6 +15,54 @@

namespace rscuda
{

template <typename T>
class DeviceBuffer final
{
public:
DeviceBuffer() = default;

explicit DeviceBuffer(std::size_t const num_elements) :
data_{DeviceBuffer<T>::allocateBuffer(num_elements)},
size_{num_elements}
{}

DeviceBuffer(std::size_t const num_elements, std::size_t const number_of_channels = 1U) :
DeviceBuffer{num_elements * number_of_channels}
{}

void reserve(std::size_t const reserve_size)
{
if (size_ < reserve_size)
{
cudaFree(data_);
data_ = DeviceBuffer<T>::allocateBuffer(reserve_size);
size_ = reserve_size;
}
}

void reserve(std::size_t const reserve_size, std::size_t const reserve_channels)
{
reserve(reserve_size * reserve_channels);
}

std::size_t size() const { return size_; }

T* data() { return data_; }

private:
static T* allocateBuffer(std::size_t const reserve_size)
{
T* datatemp{nullptr};
cudaMalloc(&datatemp, reserve_size * sizeof(T));
Copy link
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

maybe i can add an assert to the return value of cudaMalloc and cudaFree above. What do you think?

return datatemp;
}

T* data_{};
// Size is in number of elements, not bytes.
std::size_t size_{};
};

template<typename T>
std::shared_ptr<T> alloc_dev(int elements)
{
Expand Down