     // Register caller-allocated host memory with CUDA.
     // This requires that the caller allocated the memory properly vis-a-vis
     // the requirements of cudaHostRegister!
-    cudaHostRegister(context->array_h, compiletime_info.vecLength*sizeof(ComplexInput), 0);
+    size_t length = compiletime_info.vecLength*sizeof(ComplexInput);
+    // Round length up to next multiple of page size
+    length = (length+page_size-1) / page_size * page_size;
+    fprintf(stderr, "context->array_h = %p\n", context->array_h);
+    fprintf(stderr, "length = %lx\n", length);
+    cudaHostRegister(context->array_h, length, 0);
     internal->unregister_array_h = context->array_h;
     internal->free_array_h = NULL;
+    checkCudaError();
   } else {
     // allocate host memory
     cudaMallocHost(&(context->array_h), compiletime_info.vecLength*sizeof(ComplexInput));

static int writeMatrix = 1;
// this must be enabled for this option to work though, slightly hurts performance
//#define WRITE_OPTION 

typedef struct XGPUInternalContextStruct {
  //memory pointers on the device

    // Register caller-allocated host memory with CUDA.
    // This requires that the caller allocated the memory properly vis-a-vis
    // the requirements of cudaHostRegister!
    size_t length = compiletime_info.matLength*sizeof(Complex);
    // Round length up to next multiple of page size
    length = (length+page_size-1) / page_size * page_size;
    fprintf(stderr, "context->matrix_h = %p\n", context->matrix_h);
    fprintf(stderr, "length = %lx\n", length);
    cudaHostRegister(context->matrix_h, length, 0);
    internal->unregister_matrix_h = context->matrix_h;
    internal->free_matrix_h = NULL;
    checkCudaError();
  } else {
    // allocate host memory
    cudaMallocHost(&(context->matrix_h), compiletime_info.matLength*sizeof(Complex));
