hipMemcpy2D fails with invalid arg if hipMemcpyDefault used #40
Description
This simple repro fails with ROCm 5.7
#include <hip/hip_runtime.h>
constexpr int width = 64;
constexpr int height = 64;
constexpr int src_sz = width * height;
constexpr int dst_offset = src_sz; // This larger offset fails
// constexpr int dst_offset = 1; // This small offset passes
constexpr int dst_sz = src_sz + dst_offset;
int main() {
float *src, *dst;
// Allocate memory on the device
hipMalloc(&src, src_sz * sizeof(float));
hipMallocManaged(&dst, dst_sz * sizeof(float));
size_t pitch = width * sizeof(float); // no padding
auto err = hipMemcpy2D(dst + dst_offset, pitch, src, pitch,
width * sizeof(float), height, hipMemcpyDefault);
if (err != hipSuccess)
printf("hipMemcpy2DAsync failed: %s\n", hipGetErrorString(err));
hipFree(src);
hipFree(dst);
}Things that fix the invalid argument return code:
- Making the offset smaller
- Changing policy from
hipMemcpyDefaulttohipMemcpyDeviceToHost - Changing hipMallocManaged to hipMalloc
$ hipcc --version
HIP version: 5.7.31921-d1770ee1b
AMD clang version 17.0.0 (https://github.com/RadeonOpenCompute/llvm-project roc-5.7.0 23352 d1e13c532a947d0cbfc94759c00dcf152294aa13)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /opt/rocm-5.7.0/llvm/bin