[Libclc-dev] [llvm-dev] OpenCL toolset (for AMD GPU)

[Leslie Zhai via Libclc-dev]

Hi Tom,

I found the correct mailing list finally :)


在 2017年06月26日 19:21, Tom Stellard 写道:
> On 06/26/2017 04:20 AM, Leslie Zhai wrote:
>> Hi Tom,
>>
>> Thanks for your kind response!
>>
>>> clang -include /path/to/libclc/headers/clc.h -I /path/to/libclc/headers -Dcl_clang_storage_class_specifiers -target amdgcn--amdhsa -mcpu=carrizo $INPUT_FILE -o $OUTPUT_FILE
>> When I tried to build cos.cl testcase https://github.com/llvm-mirror/libclc/blob/master/test/cos.cl
>>
>> $ clang -include clc/clc.h -Dcl_clang_storage_class_specifiers -target amdgcn--amdhsa -mcpu=carrizo cos.cl -o cos.out
>>
>> cos.cl:1:27: error: pointer arguments to kernel functions must reside in '__global', '__constant' or
>>        '__local' address space
>> __kernel void foo(float4 *f) {
>>                            ^
>> 1 error generated.
>>
>> Then I added __local before float4 *f like this:
>>
>> Index: cos.cl
>> ===================================================================
>> --- cos.cl      (revision 306265)
>> +++ cos.cl      (working copy)
>> @@ -1,3 +1,3 @@
>> -__kernel void foo(float4 *f) {
>> +__kernel void foo(__local float4 *f) {
>>     *f = cos(*f);
>>   }
>>
>> because [OpenCL] Improve address space diagnostics https://reviews.llvm.org/D27671
>>
>> then rebuilt again,
>>
>> cos.cl:1:15: error: unsupported call to function _Z3cosDv4_f
>> __kernel void foo(__local float4 *f) {
>>                ^
>> 1 error generated.
>>
>> why not found *cos* function? please give me some hint, thanks a lot!
>>
> Hi,
>
> This is becuase you aren't linking the kernel with the libclc bitcode library,
> try passing the path to the libclc library to clang using this option:
>   -mlink-bitcode-file
>
> -Tom
I should at first read 
http://lists.llvm.org/pipermail/libclc-dev/2016-May/002203.html

$ clang -x cl -Dcl_clang_storage_class_specifiers -target amdgcn--amdhsa 
-mcpu=carrizo -B -Xclang -mlink-bitcode-file -Xclang 
clc/amdgcn--amdhsa.bc -include clc/clc.h -o cos.co cos.cl

but,

clang-5.0: error: unknown argument: '-mlink-bitcode-file'

Allow linking multiple bitcode files https://reviews.llvm.org/D13913 so 
buildbot and other developers had already checked the CC1Options.td for 
'mlink-bitcode-file' argument, but how to use it correctly, please give 
me some hint, thanks a lot!

PS: it works as Ricardo mentioned!

$ clang -Dcl_clang_storage_class_specifiers -isystem /usr/include 
-include clc/clc.h -target amdgcn--amdhsa -mcpu=carrizo -S -emit-llvm 
-xcl -o cos.ll cos.cl
$ llvm-link cos.ll /usr/lib/clc/amdgcn--amdhsa.bc -o cos.linked.bc
$ clang -target amdgcn--amdhsa -mcpu=carrizo cos.linked.bc -S -o 
cos.amdhsa.s

$ clang --version
iSoft clang version 5.0.0 (trunk 305877) (based on LLVM 5.0.0svn)
Target: x86_64-isoft-linux
Thread model: posix
InstalledDir: /bin

   Registered Targets:
     ...
     amdgcn     - AMD GCN GPUs
     ...
     avr        - Atmel AVR Microcontroller
     ...
     mips64el   - Mips64el [experimental]
     ...
     nvptx64    - NVIDIA PTX 64-bit
     ...
     r600       - AMD GPUs HD2XXX-HD6XXX
     ...

-- 
Regards,
Leslie Zhai https://reviews.llvm.org/p/xiangzhai/