[TT_ERROR] CUDA runtime error: an illegal memory access was encountered TurboTransformers/turbo_transformers/core/cuda_device_context.cpp:33

[auspicious3000]

Below is a rough code to explain what I did.

import torch.multiprocessing as mp

def inference(config):
    data_loader = get_loader(config)
    while True:
        for step in range(128):
            dec_outs, _ = turbo_decoder(current_pred, 
                                        memory_bank, 
                                        step, 
                                        memory_lengths=memory_lengths)
            
ctx = mp.get_context("spawn")
p = ctx.Process(target=inference, args=(config))
p.start()
p.join()

RuntimeError: CUDA error: CUBLAS_STATUS_EXECUTION_FAILED when calling `cublasSgemm( handle, opa, opb, m, n, k, &alpha, a, lda, b, ldb, &beta, c, ldc)`

I got the above error when using turbo_decoder to generate data for training. The error could appear at any iteration. Sometimes after calling inference hundreds of times, sometimes after calling it thousands of times.
It looks similar to #174, but I have not found the real solution for two days.

Hopefully you could shed some light on this.

[feifeibear]

Did huggingface's decoder work for you?

[auspicious3000]

I only tried opennmt transformer decoder, because I have many customized operations and the opennmt looks more flexible.

[feifeibear]

So, your code works when using opennmt, but is failed after switching to turbo?

[auspicious3000]

Yes.
Opennmt decoder works fine. I just used the transformer decoder similar to the one in "Attention is all you need".
Basically, I used the transformer decoder in modeling_decoders.py, and initialized it using the from_onmt method.

[feifeibear]

It looks that you are using turbo correctly. It maybe some bad cases in Turbo decoder.
Can you run OpenNMT decoder and Turbo simultaneously? And check the results of two decoders at each step.

[auspicious3000]

I have done that. The results match correctly.
However, when you call inference many times, it may break at some point.

Specifically, it looks like the inference breaks at the while loop level rather than the inner for loop level as shown in the example code above.

[feifeibear]

Did you see the Device memory usage? The crush may result from excessive GPU memory consumption.
nvidia-smi --id=0 --query-compute-apps=used_memory --format=csv -lms 100

[auspicious3000]

I checked that as well. Only used less than 50%.

[auspicious3000]

By the way, I found after removing the "with torch.no_grad" over the inference function, it is less likely to crash and runs longer. Not sure if this information is helpful.

[feifeibear]

I have some hints for you to debug.

1. use CPU only to check if the results are always the same with ONMT.
2. Try to upgrade your PyTorch version since I am not sure the error is thrown from the turbo.

[auspicious3000]

又debug了两天，gcc版本，pytorch版本，docker镜像，都换过了，其它能想到的一些方法也试了，只要调用次数一多，还是会报错。
基本上发现是，对于turbo decoder的输入和输出做的操作越少，跑得越久，输入和输出操作前都要deepcopy，否则很快就报错了。但是这样也没有从根本上解决问题，最多调用一万次左右，还是会报错。我猜可能是turbo本身不稳定，反复多次调用容易出问题？

[feifeibear]

有可能是内存管理的第三方库cub不稳定。
你把naive allocator里面都改成在显存里直接分配试试。

return allocate_impl(size, kDLGPU);
allocate_free(mem, kDLGPU);

https://github.com/Tencent/TurboTransformers/blob/master/turbo_transformers/core/allocator/naive_allocator.h#L48
https://github.com/Tencent/TurboTransformers/blob/master/turbo_transformers/core/allocator/naive_allocator.h#L73

[auspicious3000]

请问，我能否确认一下，是把49行到63行，以及74到80行，替换成您上面建议的这两行代码，然后重新编译吗？

[feifeibear]

是的

[auspicious3000]

报错了…… 求助~

[180/270] Building CXX object turbo_transformers...eFiles/tt_core.dir/allocator/allocator_api.cpp.o
FAILED: turbo_transformers/core/CMakeFiles/tt_core.dir/allocator/allocator_api.cpp.o
/usr/bin/c++  -DLOGURU_WITH_STREAMS=1 -DTT_BLAS_USE_MKL -DTT_WITH_CUDA -D__CLANG_SUPPORT_DYN_ANNOTATION__ -I/usr/local/cuda/include -I/mnt/TurboTransformers/3rd/cub -I/mnt/TurboTransformers/3rd/FP16/include -I/mnt/TurboTransformers -I/opt/miniconda3/include -I/mnt/TurboTransformers/3rd/abseil -I/mnt/TurboTransformers/3rd/dlpack/include -I/mnt/TurboTransformers/3rd/loguru -Wall -m64 -fopenmp -O3 -DNDEBUG -fPIC   -std=gnu++14 -MD -MT turbo_transformers/core/CMakeFiles/tt_core.dir/allocator/allocator_api.cpp.o -MF turbo_transformers/core/CMakeFiles/tt_core.dir/allocator/allocator_api.cpp.o.d -o turbo_transformers/core/CMakeFiles/tt_core.dir/allocator/allocator_api.cpp.o -c /mnt/TurboTransformers/turbo_transformers/core/allocator/allocator_api.cpp
In file included from /mnt/TurboTransformers/turbo_transformers/core/allocator/allocator_api.cpp:21:0:
/mnt/TurboTransformers/turbo_transformers/core/allocator/model_aware_allocator.h: In constructor 'turbo_transformers::core::allocator::ModelAwareAllocator::ModelAwareAllocator(const string&)':
/mnt/TurboTransformers/turbo_transformers/core/allocator/model_aware_allocator.h:118:13: warning: 'turbo_transformers::core::allocator::ModelAwareAllocator::gpu_chunk_list_' will be initialized after [-Wreorder]
   ChunkList gpu_chunk_list_;
             ^~~~~~~~~~~~~~~
/mnt/TurboTransformers/turbo_transformers/core/allocator/model_aware_allocator.h:114:13: warning:   'turbo_transformers::core::allocator::ChunkList turbo_transformers::core::allocator::ModelAwareAllocator::cpu_chunk_list_' [-Wreorder]
   ChunkList cpu_chunk_list_;
             ^~~~~~~~~~~~~~~
/mnt/TurboTransformers/turbo_transformers/core/allocator/model_aware_allocator.h:34:12: warning:   when initialized here [-Wreorder]
   explicit ModelAwareAllocator(const std::string& model_name)
            ^~~~~~~~~~~~~~~~~~~
In file included from /mnt/TurboTransformers/turbo_transformers/core/allocator/allocator_api.cpp:22:0:
/mnt/TurboTransformers/turbo_transformers/core/allocator/naive_allocator.h: In member function 'virtual void* turbo_transformers::core::allocator::NaiveAllocator::allocate(size_t, DLDeviceType, const string&)':
/mnt/TurboTransformers/turbo_transformers/core/allocator/naive_allocator.h:50:21: error: 'mem' was not declared in this scope
       allocate_free(mem, kDLGPU);
                     ^~~
/mnt/TurboTransformers/turbo_transformers/core/allocator/naive_allocator.h:50:7: error: 'allocate_free' was not declared in this scope
       allocate_free(mem, kDLGPU);
       ^~~~~~~~~~~~~
/mnt/TurboTransformers/turbo_transformers/core/allocator/naive_allocator.h:50:7: note: suggested alternative: 'allocate_impl'
       allocate_free(mem, kDLGPU);
       ^~~~~~~~~~~~~
       allocate_impl
/mnt/TurboTransformers/turbo_transformers/core/allocator/naive_allocator.h: In member function 'virtual void turbo_transformers::core::allocator::NaiveAllocator::free(void*, DLDeviceType, const string&)':
/mnt/TurboTransformers/turbo_transformers/core/allocator/naive_allocator.h:61:28: error: 'size' was not declared in this scope
       return allocate_impl(size, kDLGPU);
                            ^~~~
/mnt/TurboTransformers/turbo_transformers/core/allocator/naive_allocator.h:61:28: note: suggested alternative: 'dysize'
       return allocate_impl(size, kDLGPU);
                            ^~~~
                            dysize
/mnt/TurboTransformers/turbo_transformers/core/allocator/naive_allocator.h:61:40: error: return-statement with a value, in function returning 'void' [-fpermissive]
       return allocate_impl(size, kDLGPU);
                                        ^
/mnt/TurboTransformers/turbo_transformers/core/allocator/naive_allocator.h:62:7: error: 'allocate_free' was not declared in this scope
       allocate_free(mem, kDLGPU);
       ^~~~~~~~~~~~~
/mnt/TurboTransformers/turbo_transformers/core/allocator/naive_allocator.h:62:7: note: suggested alternative: 'allocate_impl'
       allocate_free(mem, kDLGPU);
       ^~~~~~~~~~~~~
       allocate_impl
[187/270] Building CXX object turbo_transformers...iles/catch2_test_main.dir/catch2_test_main.cpp.o
ninja: build stopped: subcommand failed.