[llvm-dev] OpenCL toolset (for AMD GPU)
Leslie Zhai via llvm-dev
llvm-dev at lists.llvm.org
Tue Jun 27 20:40:16 PDT 2017
Hi Tom,
I found it https://clang.llvm.org/docs/FAQ.html for
clang/include/clang/Driver/CC1Options.td
在 2017年06月27日 11:44, Leslie Zhai 写道:
> 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!
$ clang -x cl -Dcl_clang_storage_class_specifiers -target amdgcn--amdhsa
-mcpu=carrizo -Xclang -mlink-bitcode-file /usr/lib/clc/amdgcn--amdhsa.bc
-include clc/clc.h -o cos.co cos.cl
error: error reading '/tmp/amdgcn--amdhsa-ab2e8b.o'
error: unable to open output file '': 'Permission denied'
2 errors generated.
it failed to generate amdgcn--amdhsa-ab2e8b.o for
/usr/lib/clc/amdgcn--amdhsa.bc?
>
>
> 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
works!
$ clang -x assembler -target amdgcn--amdhsa -mcpu=carrizo -c -o cos.o
cos.amdhsa.s
$ clang -target amdgcn--amdhsa cos.o -o cos.co
$ llvm-readobj -sections -symbols -program-headers cos.co
File: cos.co
Format: ELF64-amdgpu-hsacobj
Arch: amdgcn
AddressSize: 64bit
LoadName:
Sections [
Section {
Index: 0
Name: (0)
Type: SHT_NULL (0x0)
Flags [ (0x0)
]
Address: 0x0
Offset: 0x0
Size: 0
Link: 0
Info: 0
AddressAlignment: 0
EntrySize: 0
}
Section {
Index: 1
Name: .note (1)
Type: SHT_NOTE (0x7)
Flags [ (0x2)
SHF_ALLOC (0x2)
]
Address: 0x200
Offset: 0x200
Size: 1136
Link: 0
Info: 0
AddressAlignment: 4
EntrySize: 0
}
Section {
Index: 2
Name: .dynsym (7)
Type: SHT_DYNSYM (0xB)
Flags [ (0x2)
SHF_ALLOC (0x2)
]
Address: 0x670
Offset: 0x670
Size: 48
Link: 4
Info: 1
AddressAlignment: 8
EntrySize: 24
}
Section {
Index: 3
Name: .hash (15)
Type: SHT_HASH (0x5)
Flags [ (0x2)
SHF_ALLOC (0x2)
]
Address: 0x6A0
Offset: 0x6A0
Size: 24
Link: 2
Info: 0
AddressAlignment: 4
EntrySize: 4
}
Section {
Index: 4
Name: .dynstr (21)
Type: SHT_STRTAB (0x3)
Flags [ (0x2)
SHF_ALLOC (0x2)
]
Address: 0x6B8
Offset: 0x6B8
Size: 5
Link: 0
Info: 0
AddressAlignment: 1
EntrySize: 0
}
Section {
Index: 5
Name: .text (29)
Type: SHT_PROGBITS (0x1)
Flags [ (0x6)
SHF_ALLOC (0x2)
SHF_EXECINSTR (0x4)
]
Address: 0x1000
Offset: 0x1000
Size: 8132
Link: 0
Info: 0
AddressAlignment: 256
EntrySize: 0
}
Section {
Index: 6
Name: .dynamic (35)
Type: SHT_DYNAMIC (0x6)
Flags [ (0x3)
SHF_ALLOC (0x2)
SHF_WRITE (0x1)
]
Address: 0x3000
Offset: 0x3000
Size: 96
Link: 4
Info: 0
AddressAlignment: 8
EntrySize: 16
}
Section {
Index: 7
Name: .AMDGPU.csdata (44)
Type: SHT_PROGBITS (0x1)
Flags [ (0x0)
]
Address: 0x0
Offset: 0x3060
Size: 0
Link: 0
Info: 0
AddressAlignment: 1
EntrySize: 0
}
Section {
Index: 8
Name: .comment (59)
Type: SHT_PROGBITS (0x1)
Flags [ (0x30)
SHF_MERGE (0x10)
SHF_STRINGS (0x20)
]
Address: 0x0
Offset: 0x3060
Size: 186
Link: 0
Info: 0
AddressAlignment: 1
EntrySize: 0
}
Section {
Index: 9
Name: .symtab (68)
Type: SHT_SYMTAB (0x2)
Flags [ (0x0)
]
Address: 0x0
Offset: 0x3120
Size: 264
Link: 11
Info: 10
AddressAlignment: 8
EntrySize: 24
}
Section {
Index: 10
Name: .shstrtab (76)
Type: SHT_STRTAB (0x3)
Flags [ (0x0)
]
Address: 0x0
Offset: 0x3228
Size: 94
Link: 0
Info: 0
AddressAlignment: 1
EntrySize: 0
}
Section {
Index: 11
Name: .strtab (86)
Type: SHT_STRTAB (0x3)
Flags [ (0x0)
]
Address: 0x0
Offset: 0x3286
Size: 64
Link: 0
Info: 0
AddressAlignment: 1
EntrySize: 0
}
]
Symbols [
Symbol {
Name: (0)
Value: 0x0
Size: 0
Binding: Local (0x0)
Type: None (0x0)
Other: 0
Section: Undefined (0x0)
}
Symbol {
Name: BB0_11 (1)
Value: 0x25D0
Size: 0
Binding: Local (0x0)
Type: None (0x0)
Other: 0
Section: .text (0x5)
}
Symbol {
Name: BB0_12 (8)
Value: 0x2A74
Size: 0
Binding: Local (0x0)
Type: None (0x0)
Other: 0
Section: .text (0x5)
}
Symbol {
Name: BB0_2 (15)
Value: 0x12E8
Size: 0
Binding: Local (0x0)
Type: None (0x0)
Other: 0
Section: .text (0x5)
}
Symbol {
Name: BB0_3 (21)
Value: 0x177C
Size: 0
Binding: Local (0x0)
Type: None (0x0)
Other: 0
Section: .text (0x5)
}
Symbol {
Name: BB0_5 (27)
Value: 0x191C
Size: 0
Binding: Local (0x0)
Type: None (0x0)
Other: 0
Section: .text (0x5)
}
Symbol {
Name: BB0_6 (33)
Value: 0x1DC0
Size: 0
Binding: Local (0x0)
Type: None (0x0)
Other: 0
Section: .text (0x5)
}
Symbol {
Name: BB0_8 (39)
Value: 0x1F94
Size: 0
Binding: Local (0x0)
Type: None (0x0)
Other: 0
Section: .text (0x5)
}
Symbol {
Name: BB0_9 (45)
Value: 0x2430
Size: 0
Binding: Local (0x0)
Type: None (0x0)
Other: 0
Section: .text (0x5)
}
Symbol {
Name: _DYNAMIC (55)
Value: 0x3000
Size: 0
Binding: Local (0x0)
Type: None (0x0)
Other [ (0x2)
STV_HIDDEN (0x2)
]
Section: .dynamic (0x6)
}
Symbol {
Name: foo (51)
Value: 0x1000
Size: 8132
Binding: Global (0x1)
Type: AMDGPU_HSA_KERNEL (0xA)
Other: 0
Section: .text (0x5)
}
]
ProgramHeaders [
ProgramHeader {
Type: PT_PHDR (0x6)
Offset: 0x40
VirtualAddress: 0x40
PhysicalAddress: 0x40
FileSize: 448
MemSize: 448
Flags [ (0x4)
PF_R (0x4)
]
Alignment: 8
}
ProgramHeader {
Type: PT_LOAD (0x1)
Offset: 0x0
VirtualAddress: 0x0
PhysicalAddress: 0x0
FileSize: 1725
MemSize: 1725
Flags [ (0x4)
PF_R (0x4)
]
Alignment: 4096
}
ProgramHeader {
Type: PT_LOAD (0x1)
Offset: 0x1000
VirtualAddress: 0x1000
PhysicalAddress: 0x1000
FileSize: 8132
MemSize: 8132
Flags [ (0x5)
PF_R (0x4)
PF_X (0x1)
]
Alignment: 4096
}
ProgramHeader {
Type: PT_LOAD (0x1)
Offset: 0x3000
VirtualAddress: 0x3000
PhysicalAddress: 0x3000
FileSize: 96
MemSize: 96
Flags [ (0x6)
PF_R (0x4)
PF_W (0x2)
]
Alignment: 4096
}
ProgramHeader {
Type: PT_DYNAMIC (0x2)
Offset: 0x3000
VirtualAddress: 0x3000
PhysicalAddress: 0x3000
FileSize: 96
MemSize: 96
Flags [ (0x6)
PF_R (0x4)
PF_W (0x2)
]
Alignment: 8
}
ProgramHeader {
Type: PT_GNU_RELRO (0x6474E552)
Offset: 0x3000
VirtualAddress: 0x3000
PhysicalAddress: 0x3000
FileSize: 96
MemSize: 4096
Flags [ (0x4)
PF_R (0x4)
]
Alignment: 1
}
ProgramHeader {
Type: PT_GNU_STACK (0x6474E551)
Offset: 0x0
VirtualAddress: 0x0
PhysicalAddress: 0x0
FileSize: 0
MemSize: 0
Flags [ (0x6)
PF_R (0x4)
PF_W (0x2)
]
Alignment: 0
}
ProgramHeader {
Type: PT_NOTE (0x4)
Offset: 0x200
VirtualAddress: 0x200
PhysicalAddress: 0x200
FileSize: 1136
MemSize: 1136
Flags [ (0x4)
PF_R (0x4)
]
Alignment: 4
}
]
>
>
> $ 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/
More information about the llvm-dev
mailing list