Unhandled exception, "Failed to build program" error in TensorFlow sample

[torokati44]

I have a Radeon RX580, running Ubuntu 18.04, with amdgpu driver (18.20-579836) from AMD. Simple OpenCL works fine on the GPU, both in Blender Cycles, LuxMark, and darktable-cltest. (clinfo output)
ComputeCpp version is CE 0.8.0. (computecpp_info output)

After building TensorFlow (on the dev/amd_gpu branch, at tensorflow/tensorflow@00b0040 plus tensorflow/tensorflow#17508 cherry-picked on top of it), by following the instructions here, the models/tutorials/image/mnist sample outputs this error:

2018-05-17 01:15:34.090983: I ./tensorflow/core/common_runtime/sycl/sycl_device.h:70] Found following OpenCL devices:
2018-05-17 01:15:34.091027: I ./tensorflow/core/common_runtime/sycl/sycl_device.h:72] id: 0, type: GPU, name: Ellesmere, vendor: Advanced Micro Devices, Inc., profile: FULL_PROFILE
terminate called after throwing an instance of 'cl::sycl::detail::exception_implementation<(cl::sycl::detail::exception_types)7, cl::sycl::detail::exception_implementation<(cl::sycl::detail::exception_types)6, cl::sycl::exception> >'
Aborted (core dumped)

All I could figure out is that the actual exception is:

Error: [ComputeCpp:RT0100] Failed to build program

And it is thrown from somewhere inside PhiloxRandomOp::Compute, at least in this specific case. Trying to run the default generator from the textgenrnn project also results in the same error. I think I've also seen a similar exception coming from CastOpBase::Compute in a different sample.

I couldn't find out any more details about the error so far.

I think #77 might be related, but not quite the same.

[torokati44]

Also, this is not with ROCm, but their older (legacy?) OpenCL driver.
I'm not sure if this is a TensorFlow or a ComputeCpp issue, but a more useful and informational error message in the sycl::exception from ComputeCpp sure would help!

EDIT: I'm slightly confused now that AMD apparently has three different OpenCL stacks (plus Clover of course, but let's not talk about that). I'm absolutely certain that I'm not running ROCm, as I couldn't get that to work properly yet, however hard I tried. As for which one of the other two is actually running, I'm not sure, but I'd guess it's the PAL-based, not Orca. Maybe you can figure it out from the clinfo output, and once I get home in a couple hours, I'll be glad to provide any info about the system.

[lukeiwanski]

@torokati44 could you provide us with clinfo output?

[torokati44]

@lukeiwanski It is already linked in the first line of the report, but here it is again: [link]

[lukeiwanski]

@torokati44 ah sorry, I missed that. cl_khr_spir is present! good first sign!
On my set-up I experienced issues when using AMDGPU-PRO (17.50.511655).
The last version that worked for me is amdgpu-pro-17.40-501128 - however that's for FirePro w8100.

At present we have no access to r5xx series - therefore debugging this issue might be challenging.

Have you tried running ComputeCpp SDK examples? Are they failing too?

[torokati44]

@lukeiwanski It's okay. I'll try some ComputeCpp examples and report back, once I have access to that machine again, which will be about 3-4 hours from now.

[torokati44]

All samples (that I've tried) stop with the same error: hello-world backtrace

Except some of them catch it, like images:

Image size: 4096
SYCL exception caught: Error: [ComputeCpp:RT0100] Failed to build program SYCL Runtime closed with the following errors:
SYCL objects are still alive while the runtime is shutting down

This probably indicates that a SYCL object was created but not properly destroyed.

Can I get my hands on a debug build of libComputeCpp.so somehow, or is that a closely guarded secret?

[DuncanMcBain]

We don't distribute a debug libComputeCpp.so at the moment. If every sample is failing, it doesn't look good, but we might get more information if you apply a patch looking something like the following:

diff --git a/samples/simple-vector-add/simple-vector-add.cpp b/samples/simple-vector-add/simple-vector-add.cpp
index 8ebd7d5..89c3e39 100644
--- a/samples/simple-vector-add/simple-vector-add.cpp
+++ b/samples/simple-vector-add/simple-vector-add.cpp
@@ -49,6 +49,28 @@ void simple_vadd(const std::array<T, N>& VA, const std::array<T, N>& VB,
   cl::sycl::buffer<T, 1> bufferB(VB.data(), numOfItems);
   cl::sycl::buffer<T, 1> bufferC(VC.data(), numOfItems);
 
+  cl::sycl::program prog(deviceQueue.get_context());
+  try {
+    prog.build_with_kernel_type<SimpleVadd<T>>("");
+  } catch (cl::sycl::exception& e) {
+    // we expect this
+    std::cout << "Expected exception caught: " << e.what();
+  }
+
+  auto clprog = prog.get();
+  size_t logsize = 0;
+  auto err = clGetProgramBuildInfo(clprog, deviceQueue.get_device().get(),
+                                   CL_PROGRAM_BUILD_LOG, sizeof(size_t), NULL,
+                                   &logsize);
+  if (err)  std::cout << "Err was nonzero! " << err << "\n";
+  std::cout << logsize << " bytes long\n";
+  auto log = std::vector<char>(logsize + 1);
+  log[logsize] = '\0';
+  err = clGetProgramBuildInfo(clprog, deviceQueue.get_device().get(),
+                              CL_PROGRAM_BUILD_LOG, logsize, log.data(), NULL);
+  if (err)  std::cout << "Err was nonzero! " << err << "\n";
+  std::cout << "Device log: " << log.data() << "\n";
+
   deviceQueue.submit([&](cl::sycl::handler& cgh) {
     auto accessorA = bufferA.template get_access<sycl_read>(cgh);
     auto accessorB = bufferB.template get_access<sycl_read>(cgh);

This patch will make the sample "simple-vector-add" output the build log in all cases, hopefully catching the build error.

When I run this using Intel CPU, I get some useful output indicating that vectorisation was successful. Unfortunately when I run this on AMD GPU, it has no output whatsoever, but I am running a very different driver version to you. I am hoping that we get useful information from your driver. (You're lucky - when my AMD driver hits a build error, it just crashes, so I'd take an actual error status as an improvement!)

[torokati44]

This is what I get with the patch:

Expected exception caught: Error: [ComputeCpp:RT0100] Failed to build program122 bytes long
Device log: Error: HSAIL doesn't support OpenCL extension spir.
An invalid option was specified.
Error: BRIG code generation failed.

terminate called after throwing an instance of 'cl::sycl::detail::exception_implementation<(cl::sycl::detail::exception_types)7, cl::sycl::detail::exception_implementation<(cl::sycl::detail::exception_types)6, cl::sycl::exception> >'
Aborted (core dumped)

Is there something wrong with my driver setup after all?

[mirh]

I think #77 might be related, but not quite the same.

Ehrm, no, I mean.. Seriously, my crappy half assed setup couldn't be farther from yours.

As for which one of the other two is actually running, I'm not sure, but I'd guess it's the PAL-based, not Orca.

Really, there's only AMDGPU-PRO driver providing top-notch compute performance for your card..
And it will already select whatever driver fits it best.

p.s. 5xx series is basically a rebranded 4xx one

Device log: Error: HSAIL doesn't support OpenCL extension spir.

HSAIL is HSA Intermediate Layer.. and HSA and ROCm are basically the same thing...
And last time I checked ROCm was still kinda on the sucky side... Idk.
Check your setup

[DuncanMcBain]

Yeah it looks like your driver is the HSA driver, and it doesn't actually support SPIR. Unfortunately, at the moment the only IRs we can emit are SPIR 1.2, SPIR-V and PTX, which don't appear to be supported by your driver. It might be that this driver is too new (or SPIR is not supported on your hardware). You might be able to get more information from AMD directly.

[mirh]

There is this in that case btw.

[torokati44]

Thanks, I could almost get that to work once, so it did run, on the GPU, except it produced garbage results (even more garbage than usual from a neural network 😄 ).
I'm afraid I still have some fragments of ROCm in operation on this system somehow, as remains of these past experimentations, hence the HSAIL error (does this make sense?).
What I think I'll try to do is install a fresh OS, and see what a simple installation of the recommended "legacy" driver will do.