Really bad performance with TensorFlow

[DarjanKrijan]

Hi,

first of all good job on getting the ROCm stack up on Arch, it was a huge hassle a couple months back when I tried.

Anyway, I'm getting a real bad performance with TensorFlow using the rocm-arch 3.8 backend on my Radeon VII.
Tried with both 2.3.1 versions from AUR and the tensorflow-rocm pip package with sort of same results.
With both I only get a low 40 images per second on the resnet50 benchmark from
https://github.com/tensorflow/benchmarks/blob/master/scripts/tf_cnn_benchmarks/tf_cnn_benchmarks.py

Benchmark logs for arch and another run on Ubuntu 20.04 with more than 5x difference are attached.
tf_benchmark_arch_rocm3.8_tf2.3.1.txt
tf_benchmark_ubuntu20.04_rocm3.8_tf2.3.1.txt

I also installed rocm-dkms from AUR and booted into 5.4.70-1-lts with no difference.

Is someone able to explain what could be throttling this?

[acxz]

@DarjanKrijan thanks for reporting this.
The performance degradation is quite concerning and with a stack this big it'll be hard to pinpoint where the bottleneck is happening on arch.

Based on your observations the issue is most likely not tensorflow so then we should go one level deeper. If you could go through the dependencies that tensorflow requires, like hip-rocclr, etc., run the test suites for them on both Arch and Ubuntu, and report the results (timing) we can take this conversation further. I'm afraid without this information we can't get to the bottom.

Well, before that let's try to resolve the differences as seen by your log file.

[acxz]

Differences on your log files:
Ubuntu:

coreClock: 1.801GHz coreCount: 60 deviceMemorySize: 15.98GiB deviceMemoryBandwidth: 953.67GiB/s
.
.
.
2020-10-15 19:30:14.633698: I tensorflow/core/common_runtime/gpu/gpu_device.cc:1402] Created TensorFlow device (/job:localhost/replica:0/task:0/device:GPU:0 with 15385 MB memory) -> physical GPU (device: 0, name: Vega 20 [Radeon VII], pci bus id: 0000:0b:00.0)
2020-10-15 19:30:14.858718: I tensorflow/core/common_runtime/gpu_fusion_pass.cc:505] ROCm Fusion is enabled.
INFO:tensorflow:Running local_init_op.
I1015 19:30:15.038749 140173098755904 session_manager.py:505] Running local_init_op.
2020-10-15 19:30:15.057322: I tensorflow/core/common_runtime/gpu_fusion_pass.cc:505] ROCm Fusion is enabled.
INFO:tensorflow:Done running local_init_op.
I1015 19:30:15.084958 140173098755904 session_manager.py:508] Done running local_init_op.
2020-10-15 19:30:15.185120: I tensorflow/core/common_runtime/gpu_fusion_pass.cc:505] ROCm Fusion is enabled.
2020-10-15 19:30:15.205727: I tensorflow/core/common_runtime/gpu_fusion_pass.cc:505] ROCm Fusion is enabled.
2020-10-15 19:30:15.216218: I tensorflow/core/common_runtime/gpu_fusion_pass.cc:505] ROCm Fusion is enabled.
Running warm up
2020-10-15 19:30:15.629997: I tensorflow/core/common_runtime/gpu_fusion_pass.cc:505] ROCm Fusion is enabled.
2020-10-15 19:30:15.656588: I tensorflow/stream_executor/platform/default/dso_loader.cc:48] Successfully opened dynamic library librocblas.so
2020-10-15 19:30:15.781469: I tensorflow/stream_executor/platform/default/dso_loader.cc:48] Successfully opened dynamic library libMIOpen.so
Done warm up

Arch:

coreClock: 1.801GHz coreCount: 60 deviceMemorySize: 15.98GiB deviceMemoryBandwidth: -1B/s
.
.
.
2020-10-15 17:33:07.395915: I tensorflow/core/common_runtime/gpu/gpu_device.cc:1402] Created TensorFlow device (/job:localhost/replica:0/task:0/device:GPU:0 with 15385 MB memory) -> physical GPU (device: 0, name: Vega 20 [Radeon VII], pci bus id: 0000:0b:00.0)
INFO:tensorflow:Running local_init_op.
I1015 17:33:08.046584 140119934445376 session_manager.py:505] Running local_init_op.
INFO:tensorflow:Done running local_init_op.
I1015 17:33:08.074286 140119934445376 session_manager.py:508] Done running local_init_op.
Running warm up
2020-10-15 17:33:09.917566: I tensorflow/stream_executor/platform/default/dso_loader.cc:48] Successfully opened dynamic library librocblas.so
2020-10-15 17:33:10.025179: I tensorflow/stream_executor/platform/default/dso_loader.cc:48] Successfully opened dynamic library libMIOpen.so
Done warm up

So basically we have an incorrect deviceMemoryBandwidth and ROCm Fusion is not enabled.

I guess we should start with these two. deviceMemoryBandwidth is particularly concerning.

Going through tf's codebase, here is where the value should be properly set:
https://github.com/tensorflow/tensorflow/blob/2b623518efa48daf88ccad6edcaf63c7d8e804c9/tensorflow/stream_executor/rocm/rocm_driver.cc#L1230

Looks like an issue with hip. I guess we should start making sure hip tests work.

[acxz]

@tpkessler do you have any idea on why the hipDeviceTotalMem is being incorrectly reported?

[DarjanKrijan]

A bit more input:
The GPU goes off like a rocket just like in Ubuntu, see screenshot
Figured maybe it's a code efficiency problem if the performance degrades by a factor of 5x.
During compilation of TensorFlow I noticed lots of warnings of loops not being unrolled, maybe the compiler backend could be somehow misconfigured? Not sure if these messages are expected or not.

Also could you point me to some tests that would somehow check the separate hip/rocm packages if there are some?
I hope they are contained in the binary release for Ubuntu for comparison...

[acxz]

Figured maybe it's a code efficiency problem if the performance degrades by a factor of 5x.

hmm it's reassuring the GPU is being used, but the code is the same between Arch and Ubuntu (they are on the same versions). It is most likely discrepancy on how the rocm packages are build on Arch vs Ubuntu imo.
If you do want to check take a look at the files installed by the Arch package and the files installed by Ubuntu package and if there are any missing files on the Arch side, that could be an issue. pacman -Ql hip-rocclr not sure what it is on Ubuntu.

During compilation of TensorFlow I noticed lots of warnings of loops not being unrolled, maybe the compiler backend could be somehow misconfigured? Not sure if these messages are expected or not.

Let's try to limit the independent variables and do our tests with the tensorflow-rocm pip package instead of using the AUR package. Although for this particular issue we should try to run the same HIP function tensorflow-rocm is running in a HIP C++ file and see if we can reproduce it outside the tensorflow environment.

Also could you point me to some tests that would somehow check the separate hip/rocm packages if there are some?

https://github.com/ROCm-Developer-Tools/HIP/blob/master/CONTRIBUTING.md#unit-testing-environment

Ideally we want these unit tests to happen in the PKGBUILD.

[tpkessler]

With my Vega 56 I get the correct bandwidth:

pciBusID: 0000:1f:00.0 name: Vega 10 XL/XT [Radeon RX Vega 56/64]     ROCm AMD GPU ISA: gfx900
coreClock: 1.59GHz coreCount: 56 deviceMemorySize: 7.98GiB deviceMemoryBandwidth: 381.47GiB/s

Do you guys both use a Radeon VII?

Tensorflow also reports

2020-10-16 09:31:52.623054: I tensorflow/core/common_runtime/gpu_fusion_pass.cc:505] ROCm Fusion is enabled.

But: 98% of my VRAM is being blocked by tensorflow, thus allocation of tensors needed for the benchmark fails.

[tpkessler]

Ideally we want these unit tests to happen in the PKGBUILD.

I had a look at the source code and have no idea on how to compile and run them. The directory units contains a single file which only checks for some command line arguments. Additionally, the README is outdated and points to files which are not present anymore.

[DarjanKrijan]

The --batch-size=128 is probably too much for the 8GB Vega 56 if you used that. --batch-size=64 may work.

Also another log with the pip package attached - ROCm Fusion is enabled here but doesn't seem to be the culprit:
tf_benchmark_arch_rocm3.8_tf2.3.1_pip.txt
deviceMemoryBandwidth is also printed properly here.

I will check if I it somehow works when I symlink (some of) the Ubuntu packages into /opt/rocm to see if the performance changes (probably later today/tonight).

[tpkessler]

I see. Thanks for the hint. These are my results with batch-size=64:

Step	Img/sec	total_loss
1	images/sec: 129.3 +/- 0.0 (jitter = 0.0)	7.608
10	images/sec: 129.4 +/- 0.1 (jitter = 0.3)	7.849
20	images/sec: 129.0 +/- 0.2 (jitter = 0.2)	8.013
30	images/sec: 129.1 +/- 0.1 (jitter = 0.2)	7.940
40	images/sec: 128.2 +/- 0.3 (jitter = 0.4)	8.137
50	images/sec: 127.9 +/- 0.3 (jitter = 0.6)	8.051
60	images/sec: 127.9 +/- 0.3 (jitter = 0.5)	7.783
70	images/sec: 127.8 +/- 0.2 (jitter = 0.6)	7.855
80	images/sec: 127.9 +/- 0.2 (jitter = 0.6)	8.006
90	images/sec: 127.9 +/- 0.2 (jitter = 0.7)	7.839
100	images/sec: 127.9 +/- 0.2 (jitter = 0.7)	8.091
----------------------------------------------------------------
total images/sec: 127.89
----------------------------------------------------------------

Far better than your results on Arch Linux and, extrapolating the results of the Radeon VII, on par with Ubuntu.

[tpkessler]

I forgot to mention: I'm using the pip package as suggested by @acxz

[DarjanKrijan]

Interesting, yours seems to work properly

Maybe compare the installed rocm-arch packages? I haven't modified any PKGBUILDs except the tensorflow one slightly (but using the pip package for now anyway)

$ pacman -Q | grep "roc\|hip\|miopen\|rccl"
hip-rocclr 3.8.0-1
hipblas 3.8.0-1
hipcub 3.8.0-1
hipsparse 3.8.0-1
hsa-rocr 3.8.0-1
hsakmt-roct 3.8.0-1
libcroco 0.6.13-2
libutf8proc 2.5.0-2
miopen-hip 3.8.0-1
procps-ng 3.3.16-2
python-keras-preprocessing 1.1.2-1
python-ptyprocess 0.6.0-4
python-setproctitle 1.1.10-5
python-tensorflow-opt-rocm 2.3.1-2
python2-ptyprocess 0.6.0-4
rccl 3.8.0-1
rocalution 3.7.0-1
rocblas 3.8.0-1
rocclr 3.8.0-1
rocfft 3.8.0-1
rock-dkms-bin 3.8-1
rock-dkms-firmware-bin 3.8-1
rocm-clang-ocl 3.8.0-1
rocm-cmake 3.8.0-1
rocm-dbgapi 3.8.0-1
rocm-dev 3.8.0-1
rocm-device-libs 3.8.0-1
rocm-dkms 3.8.0-1
rocm-gdb 3.8.0-1
rocm-libs 3.8.0-1
rocm-opencl-runtime 3.8.0-1
rocm-smi 3.8.0-1
rocm-smi-lib64 3.8.0-1
rocm-utils 3.8.0-1
rocminfo 3.8.0-1
rocprim 3.8.0-1
rocprofiler 3.8.0-1
rocrand 3.8.0-1
rocsolver 3.8.0-1
rocsparse 3.8.0-1
rocthrust 3.8.0-1
roctracer 3.8.0-1
tensorflow-opt-rocm 2.3.1-2
webrtc-audio-processing 0.3.1-2

[acxz]

Do you guys both use a Radeon VII?

I haven't test on my machine (not near my AMD machine) was just trying to debug based on the logs.

I had a look at the source code and have no idea on how to compile and run them. The directory units contains a single file which only checks for some command line arguments. Additionally, the README is outdated and points to files which are not present anymore.

hmm interesting, I gave a shot at it and couldn't find the targets. I guess we need to look deeper into this or might need to make an issue over at the HIP repo.

It is reassuring that it works for you @tpkessler, @DarjanKrijan can you try the same batch size?

Also another log with the pip package attached - ROCm Fusion is enabled here but doesn't seem to be the culprit:
tf_benchmark_arch_rocm3.8_tf2.3.1_pip.txt
deviceMemoryBandwidth is also printed properly here.

well I guess the logs are now exactly the same then, which is good, but that means the problem is somewhere deeper. Gonna put the title back to what it was previously since the log differences are not an issue anymore.
@DarjanKrijan do you know what you did to resolve those differences?

Maybe compare the installed rocm-arch packages? I haven't modified any PKGBUILDs except the tensorflow one slightly (but using the pip package for now anyway)

Also, even if the PKGBUILDs are the same version there could still be differences depending on the environment it was built in. It is recommended to build the packages in clean chroot or to use the arch4edu binary packages to keep consistency in the build packages. Although I doubt that could be an issue, but it is still important to keep in mind. Coincidentally tensorflow-rocm is now on arch4edu.

[DarjanKrijan]

I just used the tensorflow-rocm pip package that time around, can verify by the ${HOME}/.local/lib/python3.8/site-packages in the verbose log messages.

Reducing the batch size actually improved things, though not to expected levels:
batch_size_64.txt
batch_size_32.txt

Back to my theory that the bad compute kernels are produced with the backend on my Radeon VII - some bad access patterns might be hogging the memory bandwidth or clogging it with requests... would explain why the throughput is better with smaller batch sizes.

[tpkessler]

These are the packages I have installed:

packages

aomp-amdgpu 11.9.0-1
comgr 3.8.0-1
hip-rocclr 3.8.0-1
hipblas 3.8.0-1
hipcub 3.8.0-1
hipsparse 3.8.0-1
hsa-amd-aqlprofile-bin 3.8.0-1
hsa-rocr 3.8.0-1
hsakmt-roct 3.8.0-1
llvm-amdgpu 3.8.0-1
miopen-hip 3.8.0-1
miopengemm 3.8.0-1
rccl 3.8.0-1
rocalution 3.8.0-1
rocblas 3.8.0-1
rocclr 3.8.0-1
rocfft 3.8.0-1
rocm-clang-ocl 3.8.0-1
rocm-cmake 3.8.0-1
rocm-dbgapi 3.8.0-1
rocm-device-libs 3.8.0-1
rocm-gdb 3.8.0-1
rocm-opencl-runtime 3.8.0-1
rocm-smi-lib64 3.8.0-1
rocminfo 3.8.0-1
rocprim 3.8.0-1
rocprofiler 3.8.0-1
rocrand 3.8.0-1
rocsolver 3.8.0-1
rocsparse 3.8.0-1
rocthrust 3.8.0-1
roctracer 3.8.0-1

[DarjanKrijan]

So I did some tinkering.

Symlinking /opt/rocm with the Ubuntu stack on my other boot didn't work out, this segfaults left and right.

Reinstalled with the packages provided by arch4edu (before I built the rocm stack myself, directly from AUR), but now getting other errors about missing ISA or something:

$ python ./tf_cnn_benchmarks.py --num_gpus=1 --batch_size=128 --model=resnet50
WARNING:tensorflow:From /home/daren/.local/lib/python3.8/site-packages/tensorflow/python/compat/v2_compat.py:96: disable_resource_variables (from tensorflow.python.ops.variable_scope) is deprecated and will be removed in a future version.
Instructions for updating:
non-resource variables are not supported in the long term
2020-10-29 15:50:18.425549: I tensorflow/core/platform/cpu_feature_guard.cc:142] This TensorFlow binary is optimized with oneAPI Deep Neural Network Library (oneDNN)to use the following CPU instructions in performance-critical operations:  SSE3 SSE4.1 SSE4.2 AVX AVX2 FMA
To enable them in other operations, rebuild TensorFlow with the appropriate compiler flags.
2020-10-29 15:50:18.443221: I tensorflow/core/platform/profile_utils/cpu_utils.cc:104] CPU Frequency: 3994340000 Hz
2020-10-29 15:50:18.443301: I tensorflow/compiler/xla/service/service.cc:168] XLA service 0x56062a4c9c10 initialized for platform Host (this does not guarantee that XLA will be used). Devices:
2020-10-29 15:50:18.443308: I tensorflow/compiler/xla/service/service.cc:176]   StreamExecutor device (0): Host, Default Version
2020-10-29 15:50:18.444445: I tensorflow/stream_executor/platform/default/dso_loader.cc:48] Successfully opened dynamic library libamdhip64.so
2020-10-29 15:50:18.459474: I tensorflow/compiler/xla/service/service.cc:168] XLA service 0x56062a49abb0 initialized for platform ROCM (this does not guarantee that XLA will be used). Devices:
2020-10-29 15:50:18.459492: I tensorflow/compiler/xla/service/service.cc:176]   StreamExecutor device (0): Vega 20 [Radeon VII], AMDGPU ISA version: gfx906
2020-10-29 15:50:18.701926: I tensorflow/core/common_runtime/gpu/gpu_device.cc:1734] Found device 0 with properties: 
pciBusID: 0000:0b:00.0 name: Vega 20 [Radeon VII]     ROCm AMD GPU ISA: gfx906
coreClock: 1.801GHz coreCount: 60 deviceMemorySize: 15.98GiB deviceMemoryBandwidth: 953.67GiB/s
2020-10-29 15:50:18.705297: I tensorflow/stream_executor/platform/default/dso_loader.cc:48] Successfully opened dynamic library librocblas.so
2020-10-29 15:50:18.706275: I tensorflow/stream_executor/platform/default/dso_loader.cc:48] Successfully opened dynamic library libMIOpen.so
/build/hip-rocclr/src/HIP-rocm-3.8.0/rocclr/hip_code_object.cpp:92: guarantee(false && "hipErrorNoBinaryForGpu: Coudn't find binary for current devices!")
Fatal Python error: Aborted

Current thread 0x00007fac04bbf740 (most recent call first):
  File "/home/daren/.local/lib/python3.8/site-packages/tensorflow/python/client/session.py", line 701 in __init__
  File "/home/daren/.local/lib/python3.8/site-packages/tensorflow/python/client/session.py", line 1586 in __init__
  File "/home/daren/code/benchmarks/scripts/tf_cnn_benchmarks/benchmark_cnn.py", line 3538 in setup
  File "./tf_cnn_benchmarks.py", line 61 in main
  File "/home/daren/.local/lib/python3.8/site-packages/absl/app.py", line 251 in _run_main
  File "/home/daren/.local/lib/python3.8/site-packages/absl/app.py", line 300 in run
  File "./tf_cnn_benchmarks.py", line 73 in <module>
[1]    8120 abort (core dumped)  python ./tf_cnn_benchmarks.py --num_gpus=1 --batch_size=128 --model=resnet50

Checking the error messages led me to this:
ROCm/tensorflow-upstream#1106

Would some rocm-arch/arch4edu packages be missing the gfx906 (Radeon VII) ISA?

[acxz]

Would some rocm-arch/arch4edu packages be missing the gfx906 (Radeon VII) ISA?

That could be the case, if arch4edu isn't compiling on gfx906 hardware (so that it can be autodetected) or if gfx906 is not a specified build target, then I would expect you to see these errors. I guess the way around this is to figure out which packages are missing the gfx906 target and add them or to build the stack on a gfx906 machine so that it can be autodetected when building from source.

The first solution is the appropriate one, although the second solution is prob the one you want to go with, even though you just removed your source builds and would have to rebuild the entire stack (feels bad).

[tarcey]

Having the same issue with a Vega FE: Bad performance, and a performance improvement with smaller batch sizes. Tensorflow uses the entire VRAM according to both 'radeontop' and 'rocm-smi --showmeminfo all', regardless of batch size. No significant difference when using XLA, and/or eager vs graph mode, and/or TF_ROCM_FUSION_ENABLE.

TF version is 2.4, build from source from the PKGBUILD provided here.
ROCm version is 4.0.0 using arch4edu packages.

I'll try building all necessary packages locally instead of using arch4edu and report back if there is any difference.

rocminfo output

===================== HSA System Attributes ===================== Runtime Version: 1.1 System Timestamp Freq.: 1000.000000MHz Sig. Max Wait Duration: 18446744073709551615 (0xFFFFFFFFFFFFFFFF) (timestamp count) Machine Model: LARGE System Endianness: LITTLE

==========
HSA Agents

---

Agent 1

---

Name: Intel(R) Xeon(R) CPU E5-2678 v3 @ 2.50GHz
Uuid: CPU-XX
Marketing Name: Intel(R) Xeon(R) CPU E5-2678 v3 @ 2.50GHz
Vendor Name: CPU
Feature: None specified
Profile: FULL_PROFILE
Float Round Mode: NEAR
Max Queue Number: 0(0x0)
Queue Min Size: 0(0x0)
Queue Max Size: 0(0x0)
Queue Type: MULTI
Node: 0
Device Type: CPU
Cache Info:
L1: 32768(0x8000) KB
Chip ID: 0(0x0)
Cacheline Size: 64(0x40)
Max Clock Freq. (MHz): 3300
BDFID: 0
Internal Node ID: 0
Compute Unit: 24
SIMDs per CU: 0
Shader Engines: 0
Shader Arrs. per Eng.: 0
WatchPts on Addr. Ranges:1
Features: None
Pool Info:
Pool 1
Segment: GLOBAL; FLAGS: KERNARG, FINE GRAINED
Size: 65709732(0x3eaa6a4) KB
Allocatable: TRUE
Alloc Granule: 4KB
Alloc Alignment: 4KB
Accessible by all: TRUE
Pool 2
Segment: GLOBAL; FLAGS: COARSE GRAINED
Size: 65709732(0x3eaa6a4) KB
Allocatable: TRUE
Alloc Granule: 4KB
Alloc Alignment: 4KB
Accessible by all: TRUE
ISA Info:
N/A

---

Agent 2

---

Name: gfx900
Uuid: GPU-XX
Marketing Name: Vega 10 XTX [Radeon Vega Frontier Edition]
Vendor Name: AMD
Feature: KERNEL_DISPATCH
Profile: BASE_PROFILE
Float Round Mode: NEAR
Max Queue Number: 128(0x80)
Queue Min Size: 4096(0x1000)
Queue Max Size: 131072(0x20000)
Queue Type: MULTI
Node: 1
Device Type: GPU
Cache Info:
L1: 16(0x10) KB
Chip ID: 26723(0x6863)
Cacheline Size: 64(0x40)
Max Clock Freq. (MHz): 1600
BDFID: 1280
Internal Node ID: 1
Compute Unit: 64
SIMDs per CU: 4
Shader Engines: 4
Shader Arrs. per Eng.: 1
WatchPts on Addr. Ranges:4
Features: KERNEL_DISPATCH
Fast F16 Operation: FALSE
Wavefront Size: 64(0x40)
Workgroup Max Size: 1024(0x400)
Workgroup Max Size per Dimension:
x 1024(0x400)
y 1024(0x400)
z 1024(0x400)
Max Waves Per CU: 40(0x28)
Max Work-item Per CU: 2560(0xa00)
Grid Max Size: 4294967295(0xffffffff)
Grid Max Size per Dimension:
x 4294967295(0xffffffff)
y 4294967295(0xffffffff)
z 4294967295(0xffffffff)
Max fbarriers/Workgrp: 32
Pool Info:
Pool 1
Segment: GLOBAL; FLAGS: COARSE GRAINED
Size: 16760832(0xffc000) KB
Allocatable: TRUE
Alloc Granule: 4KB
Alloc Alignment: 4KB
Accessible by all: FALSE
Pool 2
Segment: GROUP
Size: 64(0x40) KB
Allocatable: FALSE
Alloc Granule: 0KB
Alloc Alignment: 0KB
Accessible by all: FALSE
ISA Info:
ISA 1
Name: amdgcn-amd-amdhsa--gfx900
Machine Models: HSA_MACHINE_MODEL_LARGE
Profiles: HSA_PROFILE_BASE
Default Rounding Mode: NEAR
Default Rounding Mode: NEAR
Fast f16: TRUE
Workgroup Max Size: 1024(0x400)
Workgroup Max Size per Dimension:
x 1024(0x400)
y 1024(0x400)
z 1024(0x400)
Grid Max Size: 4294967295(0xffffffff)
Grid Max Size per Dimension:
x 4294967295(0xffffffff)
y 4294967295(0xffffffff)
z 4294967295(0xffffffff)
FBarrier Max Size: 32
*** Done ***

[tarcey]

Something i noticed, but i don't know if that is the issue yet:
There's a choice between hip-nvccl and hip-rocclr. Can someone who is getting reasonable performance answer with which of the options you've built the software stack? I noticed that arch4edu uses hip-rocclr by default, whereas when using yay the default is hip-nvccl.

[tpkessler]

Hi @tarcey, have you compared the performance of tensorflow on Arch Linux with that on Ubuntu?

There's a choice between hip-nvccl and hip-rocclr.

The first is a wrapper around the CUDA compiler designed by nVidia, the latter is what you have to use if you have an AMD GPU.

Does the problem of full VRAM also appears in other applications? I'm thinking of OpenCL code or pure ROCm code like matrix factorisation by rocsolver.

[tarcey]

@tpkessler

have you compared the performance of tensorflow on Arch Linux with that on Ubuntu?

I measure the performance using tf_cnn_benchmarks.py and compare with the results other people have been getting with the same or similar hardware at the same settings. I get only around 60-70 images/sec on resnet50, even though the GPU is fully utilized and pinned at 1.6GHz shader clock (i adjusted the power limit). If all else fails i will probably resort to dual booting another distro or trying docker first.

Does the problem of full VRAM also appears in other applications? I'm thinking of OpenCL code or pure ROCm code like matrix factorisation by rocsolver.

I have dug a bit deeper in the memory usage issue, and it seems that this is not an issue at all but expected behavior from tensorflow. There are options to disable this behavior, e.g. by setting the TF_FORCE_GPU_ALLOW_GROWTH envvar to true. But it does not resolve the performance issue.

To address something that you mentioned earlier here, i do not get the memory bandwidth issue you and also others described. So the bad performance in TF might not be related to that. But at the same time, the circumstance that smaller batch sizes have better throughput does point to a memory access issue.
I'm currently in the process of compiling TF, which i downloaded directly from AMD's ROCm repo rather than the official TF one. I also noticed that there are a number of circular dependencies in the ROCm stack, notably llvm-amdgpu depending on an existing TF installation. I'm hoping the performance issue will be resolved after building everything ROCm related plus tf-rocm twice, to ensure that everything builds against the correct versions of the libraries.

[tarcey]

I found the problem/solution:
https://github.com/ROCmSoftwarePlatform/MIOpen/blob/master/doc/src/perfdatabase.md

It seems miopen ships with outdated or miscalculated tuning parameters for algorithms for some GPUs. I stumbled on this when comparing the 56cu and 64cu variant of the gfx900 in /opt/rocm/miopen/share/miopen/db/. According to the measurements there, one would get the impression that the 56cu variant is much faster than the 64cu variant, which can't be right.

Running the benchmark again with MIOPEN_FIND_ENFORCE=DB_UPDATE yields much more sensible results:

$ MIOPEN_FIND_ENFORCE=DB_UPDATE python tf_cnn_benchmarks.py --num_gpus=1 --batch_size=64 --model=resnet50

Step	Img/sec	total_loss
1	images/sec: 157.5 +/- 0.0 (jitter = 0.0)	7.608
10	images/sec: 157.4 +/- 0.1 (jitter = 0.5)	7.849
20	images/sec: 157.5 +/- 0.1 (jitter = 0.4)	8.013
30	images/sec: 157.4 +/- 0.1 (jitter = 0.4)	7.939
40	images/sec: 157.4 +/- 0.1 (jitter = 0.5)	8.137
50	images/sec: 157.5 +/- 0.1 (jitter = 0.5)	8.050
60	images/sec: 157.5 +/- 0.1 (jitter = 0.4)	7.785
70	images/sec: 157.6 +/- 0.1 (jitter = 0.4)	7.855
80	images/sec: 157.6 +/- 0.1 (jitter = 0.4)	8.007
90	images/sec: 157.6 +/- 0.1 (jitter = 0.4)	7.840
100	images/sec: 157.5 +/- 0.1 (jitter = 0.4)	8.093
----------------------------------------------------------------
total images/sec: 157.52
----------------------------------------------------------------

After this, running it again at a different batch size (128) but without the environment variable leads to very low performance again (~50-60). So it seems the tuning parameters are specific to batch size.
WITH the parameter and a batch size of 128:

$ MIOPEN_FIND_ENFORCE=DB_UPDATE python tf_cnn_benchmarks.py --num_gpus=1 --batch_size=128 --model=resnet50

Step	Img/sec	total_loss
1	images/sec: 159.1 +/- 0.0 (jitter = 0.0)	7.782
10	images/sec: 159.4 +/- 0.1 (jitter = 0.2)	7.740
20	images/sec: 159.2 +/- 0.1 (jitter = 0.4)	7.827
30	images/sec: 159.0 +/- 0.1 (jitter = 0.5)	7.967
40	images/sec: 158.9 +/- 0.1 (jitter = 0.6)	7.880
50	images/sec: 158.8 +/- 0.1 (jitter = 0.7)	7.792
60	images/sec: 158.8 +/- 0.1 (jitter = 0.7)	8.010
70	images/sec: 158.7 +/- 0.1 (jitter = 0.7)	7.869
80	images/sec: 158.6 +/- 0.1 (jitter = 0.6)	7.925
90	images/sec: 158.7 +/- 0.1 (jitter = 0.7)	7.735
100	images/sec: 158.8 +/- 0.1 (jitter = 0.6)	8.002
----------------------------------------------------------------
total images/sec: 158.74
----------------------------------------------------------------

Rerunning both tests without the environment parameter leads to expected (good) results. Which is good, because the tuning takes a rather long time.

Overall it is very close to what i was expecting (~180 img/sec). I already tested this with rock-dkms and a compatible kernel, and get ~170 img/sec, and further, enabling XLA with the TF_XLA_FLAGS=--tf_xla_auto_jit=2 environment variable brings me up to ~177 img/sec. I suppose the remaining performance discrepancy compared to other's tests comes down to hardware differences or minor performance regressions in newer versions of either ROCm or TF (which seems to be the case judging from these tests).

Notes:
I deleted .config/miopen and .cache/miopen before re-testing, which is where the user databases of tuning results seem to be. I also disabled power management on the card to get more accurate tuning results.
According to the documentation linked above, it seems that MIOPEN_FIND_ENFORCE=SEARCH should tune every algorithm regardless of whether it is called. But if i understand correctly, setting the parameter to DB_UPDATE will still fine-tune even further for each given task.

[tpkessler]

Wow, that's an impressive analysis! Thank you very much, @tarcey. Currently, tensorflow does not work on my PC (maybe related to python 3.9) so I cannot check this with my Vega 56. Can you reproduce the results @DarjanKrijan?

[DarjanKrijan]

Did not have the time to check again.

Last time I tried the version numbers of the packages were mixed (3.X and 4.0) and something didn't work out. This is problematic in general, maybe arch4edu should point towards a stable branch where the versions are updated at once and then build a clean stack.

Will try again towards the weekend.

[acxz]

Considering it has been more than a year since the last comment on this thread, I'm marking this solved with @tarcey's analysis.
If anyone is still have tensorflow issues unique to Arch, feel free to post another discussions thread.