QUESTION about memory management

[vakopian]

The following loop should be able to succeed on my nvidia GPU with 4GB of memory, because each step in the loop should only use less than 800MB. However this loop aborts after a few iterations with CUDA Error (2): out of memory:

    array res = constant(0, 10000, 20000, f32);
    for (int i = 0; i < 20; ++i) {
        array R = randu(10000, 20000, f32);
        res += R;
    }

I added printMemInfo() in each step and this is what it shows:

starting step 0

---------------------------------------------------------
|     POINTER      |    SIZE    |  AF LOCK  | USER LOCK |
---------------------------------------------------------
---------------------------------------------------------
starting step 1

---------------------------------------------------------
|     POINTER      |    SIZE    |  AF LOCK  | USER LOCK |
---------------------------------------------------------
|    0x1302f40000  |   762.9 MB |       Yes |        No |
---------------------------------------------------------
starting step 2

---------------------------------------------------------
|     POINTER      |    SIZE    |  AF LOCK  | USER LOCK |
---------------------------------------------------------
|    0x1332a40000  |   762.9 MB |       Yes |        No |
|    0x1302f40000  |   762.9 MB |       Yes |        No |
---------------------------------------------------------
starting step 3

---------------------------------------------------------
|     POINTER      |    SIZE    |  AF LOCK  | USER LOCK |
---------------------------------------------------------
|    0x1362540000  |   762.9 MB |       Yes |        No |
|    0x1332a40000  |   762.9 MB |       Yes |        No |
|    0x1302f40000  |   762.9 MB |       Yes |        No |
---------------------------------------------------------
starting step 4

---------------------------------------------------------
|     POINTER      |    SIZE    |  AF LOCK  | USER LOCK |
---------------------------------------------------------
|    0x1392040000  |   762.9 MB |       Yes |        No |
|    0x1362540000  |   762.9 MB |       Yes |        No |
|    0x1332a40000  |   762.9 MB |       Yes |        No |
|    0x1302f40000  |   762.9 MB |       Yes |        No |
---------------------------------------------------------
starting step 5

---------------------------------------------------------
|     POINTER      |    SIZE    |  AF LOCK  | USER LOCK |
---------------------------------------------------------
|    0x13c1b40000  |   762.9 MB |       Yes |        No |
|    0x1302f40000  |   762.9 MB |       Yes |        No |
|    0x1332a40000  |   762.9 MB |       Yes |        No |
|    0x1362540000  |   762.9 MB |       Yes |        No |
|    0x1392040000  |   762.9 MB |       Yes |        No |
---------------------------------------------------------
terminate called after throwing an instance of 'af::exception'
  what():  ArrayFire Exception (Device out of memory:101):
In function virtual void* cuda::MemoryManager::nativeAlloc(size_t)
In file src/backend/cuda/memory.cpp:97
CUDA Error (2): out of memory


In function af::array af::randu(const af::dim4&, af::dtype)
In file src/api/cpp/random.cpp:96

I have tried inserting deviceGC in each step in the loop, but the result is the same. I've also tried to sync() (and plus deviceGC) after each step, still the same result.

If I understand this correctly, even though the R array goes out of scope, for some reason AF still keeps the lock on it and does not release the memory back to the pool. It keeps accumulating these until it runs out of device memory.

Is this a bug or expected behavior? If it's expected, what can I do to make such a loop work?

This is of course an artificially constructed example to demonstrate the problem. In my real code, I have a function that uses a lot of GPU memory. I call this function many times, and the first few times it's able to finish properly, but eventually it results in the same out of memory error as above. And printMemInfo() shows the same pattern - temporary arrays that are allocated and used within the function show up as locked memory even after the function is returned.

[pavanky]

The reason for this is ArrayFire's lazy evaluation. In many cases this is a good thing because it avoids reading and writing arrays multiple times. In your example, if it did not run out of memory, the entire loop would have resulted in 20 reads (one for each R) and 1 write (for the final res). But this requires all 20 R arrays to exist in memory at the same time.

If arrayfire was not performing lazy evaluation, each loop in your example would take 2 reads (1 for R, 1 for old res) and 1 write (for the new res after addition). This results in 40 reads and 20 writes which is ~3x more memory operations.

So ArrayFire's JIT compiler, through lazy evaluation, sometimes trades increased memory requirements for performance. There are some heuristics in place to ensure that the tradeoffs are not too high, but they sort of fail when you are iteratively adding large arrays.

Calling af::deviceGC does not solve anything because the arrays are locked by the JIT compiler.

So in short, this is sort of expected behavior but arrayfire should have better heuristics for detecting these kinds of situations.

One way to resolve this is to use af::eval(res); at the end of the loop. This would mean that each iteration of the loop is going to evaluate the result immediately so you are not using too much memory. This however may result in performance degradation. One solution is to do someting like if (i % n == 0) af::eval(res); where you can choose the best n after a few trials.

[vakopian]

@pavanky First of all, I'd like to thank you and the whole team for such an outstanding support; your commitment and availability is unparalleled and is highly appreciated.

I have a quick followup question: what is the "USER LOCK"? Is it determined by the reference count of the array?

The reason for this question is that after I added some calls to eval it fixed (most of) the cases where "AF LOCK" is "Yes" and "USER LOCK" is "No". However I still have lots of unexpected memory usages that keep accumulating, where both of "AF LOCK" and "USER LOCK" are "Yes". By "unexpected" I mean that all my array allocations are already out of scope so I would expect them to be "unlocked" and available for GC. How can I "release" the "USER LOCK" to prevent this cumulative effect?

[pavanky]

USER_LOCK only turns to YES if you have arr.device<>() in your code or if you explicitly use arr.lock(). Are you getting device pointers in your code base anywhere?

[vakopian]

Yes I do use arr.device<>(): I pass the GPU memory pointers to Caffe. What should I call to unlock it after caffe returns?

[vakopian]

I just saw arr.unlock(). Is that it?

[pavanky]

@vakopian yeah arr.unlock() should do it. But if caffe and arrayfire are using different streams, you'll need to synchronize arrayfire's streams (done using af::sync()) before calling device and synchronize caffe's streams before calling unlock().

This ensures memory is not reused by arrayfire when operations are being done on another library.

[vakopian]

Yes, unlock made everything run properly, no more "out of memory"! Thank you very much.
Thank you for reminding me about sync as well. I now have to find out how to synchronize from caffe side before unlocking.

One last question about memory management. Sometimes I have code where I can't rely on c++ scope to release memory allocated in large arrays. So I created a simple function called clear (a-la matlab), that does this:

static const array EMPTY_ARRAY;

void clear(array& arr) {
    arr = EMPTY_ARRAY;
}

Is this a good idea? Is there a better way or a built-in function that does this?

[pavanky]

@vakopian There is no other way right now. But you don't need to use a static EMPTY_ARRAY. You can just do arr = af::array(). It is light weight enough to clear the memory.

I am closing the issue if this answers all the issues.

[pavanky]

@vakopian BTW if you call cudaDeviceSynchronize() it synchronizes all calls on the device. It is costlier than syncing just the stream caffe is using, but it can be a work around for now.