Bug 99552

Summary: Make Advanced Simulation Library (ASL) OpenCL GPU support work on Clover and RadeonSI
Product: Mesa Reporter: Vedran Miletić <vedran>
Component: Drivers/Gallium/radeonsiAssignee: Default DRI bug account <dri-devel>
Status: RESOLVED FIXED QA Contact: Default DRI bug account <dri-devel>
Severity: normal    
Priority: medium    
Version: git   
Hardware: All   
OS: All   
Whiteboard:
i915 platform: i915 features:
Bug Depends on:    
Bug Blocks: 99553    

Description Vedran Miletić 2017-01-26 19:33:23 UTC
Advanced Simulation Library (ASL) is a partial differential equation solver that supports running with OpenCL. When using it on Mesa/radeonsi/LLVM/AMDGPU, test/testACL/testKernel.cc [1] partially works but ultimately fails with:

$ ./testKernel
Test of "copy" function... Ok
Test of Kernel with double... Ok
Test of KernelSIMD... Ok
Test of KernelSIMDUA... Ok
Test of kernel with PrivateVariable... Ok
Test of kernel with PrivateArray...
                        BUILD LOG
 ************************************************
<unknown>:0:0: in function compute_4 void (i32 addrspace(1)*, float addrspace(1)*): invalid addrspacecast

 ************************************************

                        KERNEL SOURCE CODE
 ------------------------------------------------
#pragma OPENCL EXTENSION cl_khr_fp64 : disable

__kernel void compute_4(__global int *a_i1,
                        __global float *a_f8)
{
        uint index = get_global_id(0);
        float pa_f1[6] = {-9, 2, 0, 15, 1, 3};
        (a_f8[index]=pa_f1[a_i1[index]]);
}
 ------------------------------------------------
terminate called after throwing an instance of 'std::logic_error'
  what():  ASL ERROR: Program::build() (-11).
Aborted (core dumped)

[1] https://github.com/AvtechScientific/ASL/blob/master/test/testACL/testKernel.cc
Comment 1 Vedran Miletić 2017-02-08 16:05:11 UTC
How to test:

$ mkdir build; cd build
$ cmake -DWITH_TESTS=ON ..
$ make -j
$ make test

Right now, Mesa/LLVM latest git pulled half an hour ago I get:

Running tests...
Test project /users/miletivn/workspace/ASL/build
      Start  1: testABDFormat
 1/10 Test  #1: testABDFormat ....................   Passed    0.07 sec
      Start  2: testVectorOfElements
 2/10 Test  #2: testVectorOfElements .............   Passed    5.27 sec
      Start  3: testMatrixOfElements
 3/10 Test  #3: testMatrixOfElements .............   Passed    7.89 sec
      Start  4: testKernel
 4/10 Test  #4: testKernel .......................***Exception: Other 13.94 sec
      Start  5: testOperators
 5/10 Test  #5: testOperators ....................   Passed    7.88 sec
      Start  6: testKernelMerger
 6/10 Test  #6: testKernelMerger .................   Passed    2.65 sec
      Start  7: testPrivateVar
 7/10 Test  #7: testPrivateVar ...................***Exception: Other  3.22 sec
      Start  8: testASLData
 8/10 Test  #8: testASLData ......................   Passed   13.72 sec
      Start  9: testDistanceFunction
 9/10 Test  #9: testDistanceFunction .............   Passed   42.43 sec
      Start 10: testReductionFunction
10/10 Test #10: testReductionFunction ............   Passed   29.24 sec

80% tests passed, 2 tests failed out of 10

Label Time Summary:
DoublePrecision    =  85.62 sec (3 tests)
Performance        =   3.22 sec (1 test)

Total Test time (real) = 126.34 sec

The following tests FAILED:
	  4 - testKernel (OTHER_FAULT)
	  7 - testPrivateVar (OTHER_FAULT)
Errors while running CTest
Makefile:61: recipe for target 'test' failed
make: *** [test] Error 8
Comment 2 Vedran Miletić 2017-02-08 18:01:37 UTC
(In reply to Vedran Miletić from comment #1)
>       Start  7: testPrivateVar
>  7/10 Test  #7: testPrivateVar ...................***Exception: Other  3.22
> sec

Adding native_rsqrt fixes this, patch will appear here: http://lists.llvm.org/pipermail/libclc-dev/2017-February/
Comment 3 Vedran Miletić 2017-02-14 10:17:47 UTC
Fixed https://reviews.llvm.org/D27283 https://reviews.llvm.org/rL294786

Use of freedesktop.org services, including Bugzilla, is subject to our Code of Conduct. How we collect and use information is described in our Privacy Policy.