https://bugs.freedesktop.org/show_bug.cgi?id=96881

Bug ID: 96881 Summary: ViennaCL fails dense_blas-bench-opencl benchmark with doubles on AMD CYPRESS (DRM 2.43.0, LLVM 3.8.0) Product: Mesa Version: 11.2 Hardware: x86-64 (AMD64) OS: Linux (All) Status: NEW Severity: normal Priority: medium Component: Drivers/Gallium/r600 Assignee: dri-devel@lists.freedesktop.org Reporter: ubizjak@gmail.com QA Contact: dri-devel@lists.freedesktop.org

The dense_blas-bench-opencl benchmark from ViennaCL suite fails with doubles on AMD CYPRESS (DRM 2.43.0, LLVM 3.8.0):

$ ./dense_blas-bench-opencl

---------------------------------------------- Device Info ----------------------------------------------

Name: AMD CYPRESS (DRM 2.43.0, LLVM 3.8.0) Vendor: AMD Type: GPU Available: 1 Max Compute Units: 10 Max Work Group Size: 256 Global Mem Size: 1073741824 Local Mem Size: 32768 Local Mem Type: 1 Host Unified Memory: 1

Benchmark : BLAS ---------------- sCOPY : 64.3 GB/s sAXPY : 95.4 GB/s sDOT : 85.3 GB/s sGEMV-N : 20.8 GB/s sGEMV-T : 44.3 GB/s sGEMM-NN : 126 GFLOPs/s sGEMM-NT : 87.6 GFLOPs/s sGEMM-TN : 90.5 GFLOPs/s sGEMM-TT : 72.3 GFLOPs/s ---- Build Status = -2 ( Err = -11 ) Log: unsupported call to function __subdf3 in av_cpu Sources: #pragma OPENCL EXTENSION cl_khr_fp64 : enable

__kernel void av_cpu( __global double * vec1, uint4 size1, ...

It looks like DFmode (double) instructions are not enabled correctly in LLVM for targets that report cl_khr_fp64 extension.

clinfo reports:

Number of platforms 1 Platform Name Clover Platform Vendor Mesa Platform Version OpenCL 1.1 MESA 11.2.2 Platform Profile FULL_PROFILE Platform Extensions cl_khr_icd Platform Extensions function suffix MESA

Platform Name Clover Number of devices 1 Device Name AMD CYPRESS (DRM 2.43.0, LLVM 3.8.0) Device Vendor AMD Device Vendor ID 0x1002 Device Version OpenCL 1.1 MESA 11.2.2 Driver Version 11.2.2 Device OpenCL C Version OpenCL C 1.1 Device Type GPU Device Profile FULL_PROFILE Max compute units 10 Max clock frequency 850MHz Max work item dimensions 3 Max work item sizes 256x256x256 Max work group size 256 Preferred work group size multiple 64 Preferred / native vector sizes char 16 / 16 short 8 / 8 int 4 / 4 long 2 / 2 half 0 / 0 (n/a) float 4 / 4 double 2 / 2 (cl_khr_fp64) Half-precision Floating-point support (n/a) Single-precision Floating-point support (core) Denormals No Infinity and NANs Yes Round to nearest Yes Round to zero No Round to infinity No IEEE754-2008 fused multiply-add No Support is emulated in software No Correctly-rounded divide and sqrt operations No Double-precision Floating-point support (cl_khr_fp64) Denormals Yes Infinity and NANs Yes Round to nearest Yes Round to zero Yes Round to infinity Yes IEEE754-2008 fused multiply-add Yes Support is emulated in software No Correctly-rounded divide and sqrt operations No ...