Bug 99552 - Make Advanced Simulation Library (ASL) OpenCL GPU support work on Clover and RadeonSI
Summary: Make Advanced Simulation Library (ASL) OpenCL GPU support work on Clover and ...
Status: RESOLVED FIXED
Alias: None
Product: Mesa
Classification: Unclassified
Component: Drivers/Gallium/radeonsi (show other bugs)
Version: git
Hardware: All All
: medium normal
Assignee: Default DRI bug account
QA Contact: Default DRI bug account
URL:
Whiteboard:
Keywords:
Depends on:
Blocks: 99553
  Show dependency treegraph
 
Reported: 2017-01-26 19:33 UTC by Vedran Miletić
Modified: 2017-02-14 10:17 UTC (History)
0 users

See Also:
i915 platform:
i915 features:


Attachments

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.