Blender crashes when rendering on GPU with RadeonSI / GCN 1.0 card

I recently attempted to try OpenCL rendering in Cycles once more, using a recent nightly build of Blender 2.79 (2.78.5). Unlike previous attempts, recent versions of Blender seem to get a lot further with my GPU: The Kernel appears to be successfully compiled, however Blender crashes immediately after that and before rendering.

Notice: It’s possible that my GPU might not be officially supported, I have not recently found the list of supported cards. However I still considered this a problem worth reporting, as the cause of the crash might be unrelated whereas fixing support for any video card (especially one that appears close to working) seems like a valid proposal. By default I don’t get a GPU device listed, and need to launch Blender with the following variables which then lists it:

CYCLES_OPENCL_TEST=all CYCLES_OPENCL_SPLIT_KERNEL_TEST=1 ./blender

My video card is a Radeon R7 370 from Gigabyte (RadeonSI, GCN 1.0, AMD Pitcairn Islands). My OS is Linux openSUSE Tumbleweed x64. I have the following relevant system components: Kernel 4.11.6, Mesa 17.1.3, DRM 2.49.0, LLVM 4.0.0. Bug report and logs:

https://developer.blender.org/T51911


mircea@linux-qz0r:~/Software/Blender> CYCLES_OPENCL_TEST=all CYCLES_OPENCL_SPLIT_KERNEL_TEST=1 ./blender
Read prefs: /home/mircea/.config/blender/2.78/config/userpref.blend
found bundled python: /home/mircea/Software/Blender/2.78/python
Device init success
Compiling OpenCL program split
Kernel compilation of split finished in 10.60s.

Compiling OpenCL program base
Kernel compilation of base finished in 4.38s.

OpenCL error (AMD PITCAIRN (DRM 2.49.0 / 4.11.6-1-default, LLVM 4.0.0)): Shader Stats: SGPRS: 22 VGPRS: 59 Code Size: 6568 LDS: 4 Scratch: 0 Max Waves: 4 Spilled SGPRs: 0 Spilled VGPRs: 0 PrivMem VGPRs: 0
OpenCL error (AMD PITCAIRN (DRM 2.49.0 / 4.11.6-1-default, LLVM 4.0.0)): Shader Stats: SGPRS: 14 VGPRS: 6 Code Size: 371972 LDS: 0 Scratch: 0 Max Waves: 10 Spilled SGPRs: 0 Spilled VGPRs: 0 PrivMem VGPRs: 0
OpenCL error (AMD PITCAIRN (DRM 2.49.0 / 4.11.6-1-default, LLVM 4.0.0)): Shader Stats: SGPRS: 14 VGPRS: 6 Code Size: 371972 LDS: 0 Scratch: 0 Max Waves: 10 Spilled SGPRs: 0 Spilled VGPRs: 0 PrivMem VGPRs: 0
Writing: /tmp/blender.crash.txt
Segmentation fault (core dumped)
mircea@linux-qz0r:~/Software/Blender> 
# Blender 2.78 (sub 5), Commit date: 2017-06-09 14:30, Hash e0bc5b533cb
bpy.context.scene.render.engine = 'CYCLES'  # Property
bpy.context.scene.cycles.device = 'GPU'  # Property

# backtrace
./blender(BLI_system_backtrace+0x20) [0x1a60e60]
./blender() [0x1072878]
/lib64/libc.so.6(+0x34ae0) [0x7faa46625ae0]
/usr/lib64/dri/radeonsi_dri.so(+0x6947cd) [0x7faa2895e7cd]
/usr/lib64/libMesaOpenCL.so.1(+0x47e3c) [0x7faa1a0d0e3c]
/usr/lib64/libMesaOpenCL.so.1(+0x439d4) [0x7faa1a0cc9d4]
/usr/lib64/libMesaOpenCL.so.1(+0x442be) [0x7faa1a0cd2be]
/usr/lib64/libMesaOpenCL.so.1(+0x2a4c1) [0x7faa1a0b34c1]
/usr/lib64/libOpenCL.so.1(clEnqueueNDRangeKernel+0x62) [0x7faa1a3854a2]
./blender(_ZN3ccl17OpenCLSplitKernel30enqueue_split_kernel_data_initERKNS_16KernelDimensionsERNS_10RenderTileEiRNS_13device_memoryES7_S7_S7_S7_S7_S7_+0x13a1) [0x1fe0f71]
./blender(_ZN3ccl17DeviceSplitKernel10path_traceEPNS_10DeviceTaskERNS_10RenderTileERNS_13device_memoryES6_+0x250) [0x1fd87e0]
./blender(_ZN3ccl23OpenCLDeviceSplitKernel10thread_runEPNS_10DeviceTaskE+0x34c) [0x1fe239c]
./blender(_ZN3ccl17DedicatedTaskPool10thread_runEv+0x87) [0x2999ec7]
./blender(_ZN3ccl6thread3runEPv+0xe) [0x299b52e]
/lib64/libpthread.so.0(+0x74e7) [0x7faa47c094e7]
/lib64/libc.so.6(clone+0x3f) [0x7faa466dea2f]