diff --git a/apps/oidnTest.cpp b/apps/oidnTest.cpp
index 73bf87a2..8d6f5a9c 100644
--- a/apps/oidnTest.cpp
+++ b/apps/oidnTest.cpp
@@ -146,17 +146,18 @@ TEST_CASE("device creation", "[device]")
 
 TEST_CASE("buffer creation", "[buffer]")
 {
+  const size_t bufferSize = 1024*768*3*4; // Metal requires shared buffer size to be multiple of 16K
   DeviceRef device = makeAndCommitDevice();
 
   SECTION("default buffer")
   {
-    BufferRef buffer = device.newBuffer(1234567);
+    BufferRef buffer = device.newBuffer(bufferSize);
     REQUIRE(device.getError() == Error::None);
   }
 
   SECTION("device buffer")
   {
-    BufferRef buffer = device.newBuffer(1234567, Storage::Device);
+    BufferRef buffer = device.newBuffer(bufferSize, Storage::Device);
     REQUIRE(device.getError() == Error::None);
   }
 
@@ -167,14 +168,14 @@ TEST_CASE("buffer creation", "[buffer]")
 
     if (managedMemorySupported)
     {
-      BufferRef buffer = device.newBuffer(1234567, Storage::Managed);
+      BufferRef buffer = device.newBuffer(bufferSize, Storage::Managed);
       REQUIRE(device.getError() == Error::None);
     }
   }
 
   SECTION("shared buffer")
   {
-    BufferRef buffer = device.newBuffer(1234567);
+    BufferRef buffer = device.newBuffer(bufferSize);
     REQUIRE(device.getError() == Error::None);
 
     BufferRef sharedBuffer = device.newBuffer(buffer.getData(), buffer.getSize());
@@ -195,7 +196,7 @@ TEST_CASE("buffer creation", "[buffer]")
 
   SECTION("zero-sized shared buffer")
   {
-    BufferRef buffer = device.newBuffer(1234567);
+    BufferRef buffer = device.newBuffer(bufferSize);
     REQUIRE(device.getError() == Error::None);
 
     BufferRef sharedBuffer = device.newBuffer(buffer.getData(), 0);
@@ -218,13 +219,13 @@ TEST_CASE("buffer creation", "[buffer]")
 
   SECTION("invalid buffer storage")
   {
-    BufferRef buffer = device.newBuffer(10000, static_cast<Storage>(-42));
+    BufferRef buffer = device.newBuffer(bufferSize, static_cast<Storage>(-42));
     REQUIRE(device.getError() == Error::InvalidArgument);
   }
 
   SECTION("device released before buffer")
   {
-    BufferRef buffer = device.newBuffer(123456);
+    BufferRef buffer = device.newBuffer(bufferSize);
     REQUIRE(device.getError() == Error::None);
     device.release();
   }
@@ -715,6 +716,7 @@ TEST_CASE("shared image", "[shared_image]")
 
 TEST_CASE("inplace filter", "[inplace_filter]")
 {
+  // Metal requires shared buffer size to be multiple of 16K
   const int W = 1920;
   const int H = 1080;
 
diff --git a/devices/metal/metal_buffer.mm b/devices/metal/metal_buffer.mm
index 5c7fe7be..f3f94789 100644
--- a/devices/metal/metal_buffer.mm
+++ b/devices/metal/metal_buffer.mm
@@ -64,7 +64,7 @@
                                   deallocator: nil];
 
     if (!buffer)
-      throw Exception(Error::OutOfMemory, "failed to create buffer");
+      throw Exception(Error::InvalidArgument, "failed to create buffer");
   }
 
   MetalBuffer::MetalBuffer(MetalEngine* engine, id<MTLBuffer> buffer)
diff --git a/doc/api.md b/doc/api.md
index 2d37fc7d..0ef9a59d 100644
--- a/doc/api.md
+++ b/doc/api.md
@@ -714,8 +714,9 @@ its size in bytes. At buffer construction time no buffer data is allocated, but
 the buffer data provided by the user is used. The buffer data must remain valid
 for as long as the buffer may be used, and the user is responsible to free the
 buffer data when no longer required. The user must also ensure that the memory is
-accessible by the device by using allocation functions supported by the device
-(e.g. `sycl::malloc_device`, `cudaMalloc`, `hipMalloc`).
+accessible to the device by using a supported allocation function (e.g.
+`sycl::malloc_device`, `cudaMalloc`, `hipMalloc`) and alignment (e.g. Metal
+requires the allocation to be page-aligned).
 
 Buffers can be also imported from graphics APIs as external memory, to avoid
 expensive copying of data through host memory. Different types of external