Skip to content

aiscp: Handle files > 2,147,479,552 bytes #157

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

Closed
kurtmcmillan wants to merge 2 commits into from
Closed

aiscp: Handle files > 2,147,479,552 bytes #157

Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
6de0eaa
aiscp: Handle files > 2,147,479,552 bytes
kurtmcmillan Jan 15, 2026
96aecae
aiscp: Handle empty source files
kurtmcmillan Jan 15, 2026
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
63 changes: 41 additions & 22 deletions examples/aiscp/aiscp.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,9 +10,14 @@
* This _very_basic_ program copies SOURCE to DEST via GPU memory.
*/

#ifndef AISCP_CHUNK_SIZE
Copy link

Copilot AI Jan 16, 2026

Choose a reason for hiding this comment

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

The AISCP_CHUNK_SIZE constant (0x7ffff000lu = 2,147,479,552 bytes) should have a comment explaining why this specific value was chosen and how it relates to the Linux maximum I/O size mentioned in the PR description. This would help future maintainers understand the rationale behind this magic number.

Suggested change
#ifndef AISCP_CHUNK_SIZE
#ifndef AISCP_CHUNK_SIZE
// Use a chunk size just below 2 GiB (2^31 bytes) and one page (4 KiB) smaller than
// that limit: 0x7ffff000 = 2_147_479_552 bytes. This keeps each I/O request within
// the Linux maximum I/O size constraints while remaining page-aligned for O_DIRECT.

Copilot uses AI. Check for mistakes.
#define AISCP_CHUNK_SIZE 0x7ffff000lu
#endif

#include <hipfile.h>
#include <hip/hip_runtime_api.h>

#include <algorithm>
#include <cerrno>
#include <cstdio>
#include <cstdlib>
Expand Down Expand Up @@ -70,6 +75,16 @@ close_file(const char *path, int fd, hipFileHandle_t handle)
return 0;
}

/// @brief Round value to the next multiple of align. Align _must_ be a power of 2.
/// @param value The value to round up.
/// @param align Value will be round up to a multiple of align
Copy link

Copilot AI Jan 16, 2026

Choose a reason for hiding this comment

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

The parameter description for 'align' contains a grammatical error: "Value will be round up" should be "Value will be rounded up".

Suggested change
/// @param align Value will be round up to a multiple of align
/// @param align Value will be rounded up to a multiple of align

Copilot uses AI. Check for mistakes.
/// @return Value rounded up to a multiple of align.
static inline size_t
alignUp(size_t value, size_t align)
{
return (value + align - 1) & ~(align - 1);
}

int
main(int argc, char *argv[])
{
Expand All @@ -80,7 +95,7 @@ main(int argc, char *argv[])
hipError_t hip_err;
int exit_status = EXIT_FAILURE;
size_t buffer_size, file_size, block_size;
ssize_t nbytes;
ssize_t ncopy{}, nwrite{}, nread{}, nbytes{};

if (argc != 3) {
fprintf(stderr, "Usage: %s SOURCE DEST\n", argv[0]);
Expand All @@ -106,42 +121,46 @@ main(int argc, char *argv[])
}

if (0 == file_size) {
exit_status = EXIT_SUCCESS;
goto close_dst;
}

if (open_file(src_path, O_RDONLY, 0, &src_fd, &src_handle)) {
goto close_dst;
}

buffer_size = file_size;
// If needed, round buffer_size up to the next multiple of block_size
if (buffer_size & (block_size - 1)) {
buffer_size = (buffer_size + block_size) & ~(block_size - 1);
}
hip_err = hipMalloc(&devbuf, buffer_size);
buffer_size = alignUp(std::min(file_size, AISCP_CHUNK_SIZE), block_size);
hip_err = hipMalloc(&devbuf, buffer_size);
if (hipSuccess != hip_err) {
fprintf(stderr, "Could not allocate device buffer (%d)", hip_err);
goto close_src;
}

nbytes = hipFileRead(src_handle, devbuf, buffer_size, 0, 0);
if (nbytes < 0 || file_size != static_cast<size_t>(nbytes)) {
fprintf(stderr, "Could not read from %s (%zd) (%s)\n", src_path, nbytes,
IS_HIPFILE_ERR(nbytes) ? HIPFILE_ERRSTR(nbytes) : strerror(errno));
goto free_devbuf;
}
while (static_cast<size_t>(ncopy) < file_size) {
nread = hipFileRead(src_handle, devbuf, buffer_size, ncopy, 0);
Copy link

Copilot AI Jan 16, 2026

Choose a reason for hiding this comment

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

The hipFileRead call always reads buffer_size bytes, even on the last chunk where fewer bytes may remain. While hipFileRead likely handles this correctly by returning the actual bytes read, it would be more explicit and safer to limit the read size to the remaining bytes: std::min(buffer_size, file_size - static_cast<size_t>(ncopy)). This also makes the code's intent clearer.

Suggested change
nread = hipFileRead(src_handle, devbuf, buffer_size, ncopy, 0);
const size_t remaining = file_size - static_cast<size_t>(ncopy);
const size_t read_size = std::min(buffer_size, remaining);
nread = hipFileRead(src_handle, devbuf, read_size, ncopy, 0);

Copilot uses AI. Check for mistakes.
if (nread < 0) {
fprintf(stderr, "Could not read from %s (%zd) (%s)\n", src_path, nread,
IS_HIPFILE_ERR(nread) ? HIPFILE_ERRSTR(nread) : strerror(errno));
goto free_devbuf;
}

nbytes = hipFileWrite(dst_handle, devbuf, buffer_size, 0, 0);
if (nbytes < 0 || buffer_size != static_cast<size_t>(nbytes)) {
fprintf(stderr, "Could not write to %s (%zd) (%s)\n", src_path, nbytes,
IS_HIPFILE_ERR(nbytes) ? HIPFILE_ERRSTR(nbytes) : strerror(errno));
goto free_devbuf;
nwrite = 0;
while (nwrite < nread) {
nbytes =
hipFileWrite(dst_handle, devbuf, alignUp(static_cast<size_t>(nread - nwrite), block_size),
static_cast<hoff_t>(ncopy + nwrite), static_cast<hoff_t>(nwrite));
if (nbytes < 0) {
fprintf(stderr, "Could not write to %s (%zd) (%s)\n", src_path, nbytes,
Copy link

Copilot AI Jan 16, 2026

Choose a reason for hiding this comment

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

The error message incorrectly references src_path instead of dst_path. When hipFileWrite fails, the error should indicate the destination file, not the source file.

Suggested change
fprintf(stderr, "Could not write to %s (%zd) (%s)\n", src_path, nbytes,
fprintf(stderr, "Could not write to %s (%zd) (%s)\n", dst_path, nbytes,

Copilot uses AI. Check for mistakes.
IS_HIPFILE_ERR(nbytes) ? HIPFILE_ERRSTR(nbytes) : strerror(errno));
goto free_devbuf;
}
nwrite += nbytes;
}
ncopy += nread;
}
Comment on lines +139 to 160
Copy link

Copilot AI Jan 16, 2026

Choose a reason for hiding this comment

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

The outer while loop continues until ncopy equals file_size, but there's no handling for the case when hipFileRead returns 0 (EOF) before reaching file_size. If hipFileRead returns 0 prematurely, this will result in an infinite loop since ncopy will never reach file_size. Consider adding a check: if nread is 0, break from the loop or handle it as an error.

Copilot uses AI. Check for mistakes.

if (file_size < buffer_size) {
if (-1 == ftruncate(dst_fd, static_cast<off_t>(file_size))) {
fprintf(stderr, "Could not truncate %s (%zu) (%s)\n", dst_path, file_size, strerror(errno));
}
if (-1 == ftruncate(dst_fd, static_cast<off_t>(file_size))) {
fprintf(stderr, "Could not truncate %s (%zu) (%s)\n", dst_path, file_size, strerror(errno));
}

exit_status = EXIT_SUCCESS;
Expand Down