thr3ads.net - llvm dev - [LLVMdev] Problem with PTX assembly printing (NVPTX backend) [Feb 2013]

If this information is useful, please help other people find it:
Share via:

nkavv at physics.auth.gr

2013-Feb-04 19:04 UTC

[LLVMdev] Problem with PTX assembly printing (NVPTX backend)

Hi,
> Can you post the llc command line you're using?  Can you post an LLVM
IR
> file that causes this behavior?
yes:

${LLVM_PATH}/bin/llc -o helloworld.s -march=nvptx helloworld.ll

where LLVM_PATH my local installation path for LLVM.

Also attaching helloworld.c:

#include <stdio.h>

int main(void) {
   printf("Hello World!\n");
   return 0;
}



and helloworld.ll:

; ModuleID = 'helloworld.c'
target datalayout =  
"e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-f80:128:128-v64:64:64-v128:128:128-a0:0:64-f80:32:32-n8:16:32-S32"
target triple = "i686-w64-mingw32"

@str = private unnamed_addr constant [13 x i8] c"Hello World!\00"

define i32 @main() nounwind {
   %puts = tail call i32 @puts(i8* getelementptr inbounds ([13 x i8]*  
@str, i32 0, i32 0))
   ret i32 0
}

declare i32 @puts(i8* nocapture) nounwind


For instance, the "mips" target produces this:

	.section .mdebug.abi32
	.previous
	.file	"helloworld.ll"
	.text
	.globl	main
	.align	2
	.type	main, at function
	.set	nomips16                # @main
	.ent	main
main:
	.frame	$sp,24,$ra
	.mask 	0x80000000,-4
	.fmask	0x00000000,0
	.set	noreorder
	.set	nomacro
	.set	noat
# BB#0:
	lui	$2, %hi(_gp_disp)
	addiu	$2, $2, %lo(_gp_disp)
	addiu	$sp, $sp, -24
	sw	$ra, 20($sp)            # 4-byte Folded Spill
	addu	$gp, $2, $25
	lw	$1, %got($str)($gp)
	lw	$25, %call16(puts)($gp)
	jalr	$25
	addiu	$4, $1, %lo($str)
	addiu	$2, $zero, 0
	lw	$ra, 20($sp)            # 4-byte Folded Reload
	jr	$ra
	addiu	$sp, $sp, 24
	.set	at
	.set	macro
	.set	reorder
	.end	main
$tmp2:
	.size	main, ($tmp2)-main

	.type	$str, at object            # @str
	.section	.rodata.str1.4,"aMS", at progbits,1
	.align	2
$str:
	.asciz	 "Hello World!"
	.size	$str, 13





Best regards
Nikolaos Kavvadias

>>  BTW I've built LLVM with mingw (strange mixture of gcc-3.4.5 and
>>>> g++-4.6.1). Native compilation, MIPS backend and other things
that i've
>>>> tested, all work properly.
>>>>
>>>
>>> Perhaps this "strange mixture" is causing some symbol
errors somewhere.
>>> Can you post the exact sequence of steps you used to build LLVM,
along
>>> with
>>> the llc command-line that fails for you?
>>>
>>
>> My configuration options were:
>>
>> --prefix=/my/local/path/or/so --enable-shared --enable-targets=all
>> --enable-optimized --disable-libffi --disable-debug-runtime
>> --disable-assertions --disable-expensive-checks
>>
>> Following configuration i did the typical make and make install (to a
>> selected path by --prefix.
>>
>>
>> Best regards
>> Nikolaos Kavvadias
>>
>>
>>
>>>
>>>
>>>> Any help is appreciated.
>>>>
>>>> Best regards,
>>>> Nikolaos Kavvadias
>>>>
>>>>
>>>>
>>>> ______________________________****_________________
>>>>
>>>> LLVM Developers mailing list
>>>> LLVMdev at cs.uiuc.edu         http://llvm.cs.uiuc.edu
>>>>
http://lists.cs.uiuc.edu/****mailman/listinfo/llvmdev<http://lists.cs.uiuc.edu/**mailman/listinfo/llvmdev>
>>>>
<http:**//lists.cs.uiuc.edu/mailman/**listinfo/llvmdev<http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev>
>>>> >
>>>>
>>>>
>>>
>>>
>>> --
>>>
>>> Thanks,
>>>
>>> Justin Holewinski
>>>
>>>
>>
>>
>>
>
>
> --
>
> Thanks,
>
> Justin Holewinski
>

Justin Holewinski

2013-Feb-04 19:15 UTC

head link

[LLVMdev] Problem with PTX assembly printing (NVPTX backend)

Alright, couple of points here:

1. Address space 0 is invalid for global variables.  This is causing a
crash in llc where we use llvm_unreachable() on this case.  This is most
likely why you're seeing llc run forever.  The fix for this is to use
address space 1 for globals, which puts them into PTX global memory.  On
our side, we should provide a meaningful error message in this case.

2. The stdio.h header you're using has no chance of producing IR that is
valid for NVPTX.  In this case, printf gets translated into a call to
puts(), which will be undefined if you try to execute the PTX.

I'm putting together some documentation on using the back-end; hopefully
I'll have something ready soon.


On Mon, Feb 4, 2013 at 2:04 PM, <nkavv at physics.auth.gr> wrote:
> Hi,
>
>
>  Can you post the llc command line you're using?  Can you post an LLVM
IR
>> file that causes this behavior?
>>
>
> yes:
>
> ${LLVM_PATH}/bin/llc -o helloworld.s -march=nvptx helloworld.ll
>
> where LLVM_PATH my local installation path for LLVM.
>
> Also attaching helloworld.c:
>
> #include <stdio.h>
>
> int main(void) {
>   printf("Hello World!\n");
>   return 0;
> }
>
>
>
> and helloworld.ll:
>
> ; ModuleID = 'helloworld.c'
> target datalayout = "e-p:32:32:32-i1:8:8-i8:8:8-**
> i16:16:16-i32:32:32-i64:64:64-**f32:32:32-f64:64:64-f80:128:**
> 128-v64:64:64-v128:128:128-a0:**0:64-f80:32:32-n8:16:32-S32"
> target triple = "i686-w64-mingw32"
>
> @str = private unnamed_addr constant [13 x i8] c"Hello World!\00"
>
> define i32 @main() nounwind {
>   %puts = tail call i32 @puts(i8* getelementptr inbounds ([13 x i8]* @str,
> i32 0, i32 0))
>   ret i32 0
> }
>
> declare i32 @puts(i8* nocapture) nounwind
>
>
> For instance, the "mips" target produces this:
>
>         .section .mdebug.abi32
>         .previous
>         .file   "helloworld.ll"
>         .text
>         .globl  main
>         .align  2
>         .type   main, at function
>         .set    nomips16                # @main
>         .ent    main
> main:
>         .frame  $sp,24,$ra
>         .mask   0x80000000,-4
>         .fmask  0x00000000,0
>         .set    noreorder
>         .set    nomacro
>         .set    noat
> # BB#0:
>         lui     $2, %hi(_gp_disp)
>         addiu   $2, $2, %lo(_gp_disp)
>         addiu   $sp, $sp, -24
>         sw      $ra, 20($sp)            # 4-byte Folded Spill
>         addu    $gp, $2, $25
>         lw      $1, %got($str)($gp)
>         lw      $25, %call16(puts)($gp)
>         jalr    $25
>         addiu   $4, $1, %lo($str)
>         addiu   $2, $zero, 0
>         lw      $ra, 20($sp)            # 4-byte Folded Reload
>         jr      $ra
>         addiu   $sp, $sp, 24
>         .set    at
>         .set    macro
>         .set    reorder
>         .end    main
> $tmp2:
>         .size   main, ($tmp2)-main
>
>         .type   $str, at object            # @str
>         .section        .rodata.str1.4,"aMS",@**progbits,1
>         .align  2
> $str:
>         .asciz   "Hello World!"
>         .size   $str, 13
>
>
>
>
>
> Best regards
> Nikolaos Kavvadias
>
>
>   BTW I've built LLVM with mingw (strange mixture of gcc-3.4.5 and
>>>
>>>> g++-4.6.1). Native compilation, MIPS backend and other things
that i've
>>>>> tested, all work properly.
>>>>>
>>>>>
>>>> Perhaps this "strange mixture" is causing some symbol
errors somewhere.
>>>> Can you post the exact sequence of steps you used to build
LLVM, along
>>>> with
>>>> the llc command-line that fails for you?
>>>>
>>>>
>>> My configuration options were:
>>>
>>> --prefix=/my/local/path/or/so --enable-shared --enable-targets=all
>>> --enable-optimized --disable-libffi --disable-debug-runtime
>>> --disable-assertions --disable-expensive-checks
>>>
>>> Following configuration i did the typical make and make install (to
a
>>> selected path by --prefix.
>>>
>>>
>>> Best regards
>>> Nikolaos Kavvadias
>>>
>>>
>>>
>>>
>>>>
>>>>  Any help is appreciated.
>>>>>
>>>>> Best regards,
>>>>> Nikolaos Kavvadias
>>>>>
>>>>>
>>>>>
>>>>> ______________________________******_________________
>>>>>
>>>>>
>>>>> LLVM Developers mailing list
>>>>> LLVMdev at cs.uiuc.edu         http://llvm.cs.uiuc.edu
>>>>>
http://lists.cs.uiuc.edu/******mailman/listinfo/llvmdev<http://lists.cs.uiuc.edu/****mailman/listinfo/llvmdev>
>>>>>
<http:**//lists.cs.uiuc.edu/**mailman/**listinfo/llvmdev<http://lists.cs.uiuc.edu/**mailman/listinfo/llvmdev>
>>>>> >
>>>>>
<http:**//lists.cs.uiuc.edu/**mailman/**listinfo/llvmdev<http://lists.cs.uiuc.edu/mailman/**listinfo/llvmdev>
>>>>>
<htt**p://lists.cs.uiuc.edu/mailman/**listinfo/llvmdev<http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev>
>>>>> >
>>>>> >
>>>>>
>>>>>
>>>>>
>>>>
>>>> --
>>>>
>>>> Thanks,
>>>>
>>>> Justin Holewinski
>>>>
>>>>
>>>>
>>>
>>>
>>>
>>
>> --
>>
>> Thanks,
>>
>> Justin Holewinski
>>
>>
>
>

-- 

Thanks,

Justin Holewinski
-------------- next part --------------
An HTML attachment was scrubbed...
URL:
<http://lists.llvm.org/pipermail/llvm-dev/attachments/20130204/b04bf8d1/attachment.html>

nkavv at physics.auth.gr

2013-Feb-04 19:49 UTC

head link

[LLVMdev] Problem with PTX assembly printing (NVPTX backend)

Hi Justin,

thank you for your time, you were right!

This helloworld.c (no I/O calls):

int main(void) {return 0;}

produces the following PTX output:

//
// Generated by LLVM NVPTX Back-End
//

.version 3.1
.target sm_10, texmode_independent
.address_size 32


	// .globl	main
.func  (.reg .b32 func_retval0) main(

)                                       // @main
{
	.reg .pred %p<396>;
	.reg .s16 %rc<396>;
	.reg .s16 %rs<396>;
	.reg .s32 %r<396>;
	.reg .s64 %rl<396>;
	.reg .f32 %f<396>;
	.reg .f64 %fl<396>;

// BB#0:
	mov.u32 	%r0, 0;
	mov.b32	func_retval0, %r0;
	ret;
}



Best regards
Nikolaos Kavvadias



Quoting Justin Holewinski <justin.holewinski at gmail.com>:
> Alright, couple of points here:
>
> 1. Address space 0 is invalid for global variables.  This is causing a
> crash in llc where we use llvm_unreachable() on this case.  This is most
> likely why you're seeing llc run forever.  The fix for this is to use
> address space 1 for globals, which puts them into PTX global memory.  On
> our side, we should provide a meaningful error message in this case.
>
> 2. The stdio.h header you're using has no chance of producing IR that
is
> valid for NVPTX.  In this case, printf gets translated into a call to
> puts(), which will be undefined if you try to execute the PTX.
>
> I'm putting together some documentation on using the back-end;
hopefully
> I'll have something ready soon.
>
>
> On Mon, Feb 4, 2013 at 2:04 PM, <nkavv at physics.auth.gr> wrote:
>
>> Hi,
>>
>>
>>  Can you post the llc command line you're using?  Can you post an
LLVM IR
>>> file that causes this behavior?
>>>
>>
>> yes:
>>
>> ${LLVM_PATH}/bin/llc -o helloworld.s -march=nvptx helloworld.ll
>>
>> where LLVM_PATH my local installation path for LLVM.
>>
>> Also attaching helloworld.c:
>>
>> #include <stdio.h>
>>
>> int main(void) {
>>   printf("Hello World!\n");
>>   return 0;
>> }
>>
>>
>>
>> and helloworld.ll:
>>
>> ; ModuleID = 'helloworld.c'
>> target datalayout = "e-p:32:32:32-i1:8:8-i8:8:8-**
>> i16:16:16-i32:32:32-i64:64:64-**f32:32:32-f64:64:64-f80:128:**
>> 128-v64:64:64-v128:128:128-a0:**0:64-f80:32:32-n8:16:32-S32"
>> target triple = "i686-w64-mingw32"
>>
>> @str = private unnamed_addr constant [13 x i8] c"Hello
World!\00"
>>
>> define i32 @main() nounwind {
>>   %puts = tail call i32 @puts(i8* getelementptr inbounds ([13 x i8]*
@str,
>> i32 0, i32 0))
>>   ret i32 0
>> }
>>
>> declare i32 @puts(i8* nocapture) nounwind
>>
>>
>> For instance, the "mips" target produces this:
>>
>>         .section .mdebug.abi32
>>         .previous
>>         .file   "helloworld.ll"
>>         .text
>>         .globl  main
>>         .align  2
>>         .type   main, at function
>>         .set    nomips16                # @main
>>         .ent    main
>> main:
>>         .frame  $sp,24,$ra
>>         .mask   0x80000000,-4
>>         .fmask  0x00000000,0
>>         .set    noreorder
>>         .set    nomacro
>>         .set    noat
>> # BB#0:
>>         lui     $2, %hi(_gp_disp)
>>         addiu   $2, $2, %lo(_gp_disp)
>>         addiu   $sp, $sp, -24
>>         sw      $ra, 20($sp)            # 4-byte Folded Spill
>>         addu    $gp, $2, $25
>>         lw      $1, %got($str)($gp)
>>         lw      $25, %call16(puts)($gp)
>>         jalr    $25
>>         addiu   $4, $1, %lo($str)
>>         addiu   $2, $zero, 0
>>         lw      $ra, 20($sp)            # 4-byte Folded Reload
>>         jr      $ra
>>         addiu   $sp, $sp, 24
>>         .set    at
>>         .set    macro
>>         .set    reorder
>>         .end    main
>> $tmp2:
>>         .size   main, ($tmp2)-main
>>
>>         .type   $str, at object            # @str
>>         .section        .rodata.str1.4,"aMS",@**progbits,1
>>         .align  2
>> $str:
>>         .asciz   "Hello World!"
>>         .size   $str, 13
>>
>>
>>
>>
>>
>> Best regards
>> Nikolaos Kavvadias
>>
>>
>>   BTW I've built LLVM with mingw (strange mixture of gcc-3.4.5 and
>>>>
>>>>> g++-4.6.1). Native compilation, MIPS backend and other
things that i've
>>>>>> tested, all work properly.
>>>>>>
>>>>>>
>>>>> Perhaps this "strange mixture" is causing some
symbol errors somewhere.
>>>>> Can you post the exact sequence of steps you used to build
LLVM, along
>>>>> with
>>>>> the llc command-line that fails for you?
>>>>>
>>>>>
>>>> My configuration options were:
>>>>
>>>> --prefix=/my/local/path/or/so --enable-shared
--enable-targets=all
>>>> --enable-optimized --disable-libffi --disable-debug-runtime
>>>> --disable-assertions --disable-expensive-checks
>>>>
>>>> Following configuration i did the typical make and make install
(to a
>>>> selected path by --prefix.
>>>>
>>>>
>>>> Best regards
>>>> Nikolaos Kavvadias
>>>>
>>>>
>>>>
>>>>
>>>>>
>>>>>  Any help is appreciated.
>>>>>>
>>>>>> Best regards,
>>>>>> Nikolaos Kavvadias
>>>>>>
>>>>>>
>>>>>>
>>>>>> ______________________________******_________________
>>>>>>
>>>>>>
>>>>>> LLVM Developers mailing list
>>>>>> LLVMdev at cs.uiuc.edu         http://llvm.cs.uiuc.edu
>>>>>>
http://lists.cs.uiuc.edu/******mailman/listinfo/llvmdev<http://lists.cs.uiuc.edu/****mailman/listinfo/llvmdev>
>>>>>>
<http:**//lists.cs.uiuc.edu/**mailman/**listinfo/llvmdev<http://lists.cs.uiuc.edu/**mailman/listinfo/llvmdev>
>>>>>> >
>>>>>>
<http:**//lists.cs.uiuc.edu/**mailman/**listinfo/llvmdev<http://lists.cs.uiuc.edu/mailman/**listinfo/llvmdev>
>>>>>>
<htt**p://lists.cs.uiuc.edu/mailman/**listinfo/llvmdev<http://lists.cs.uiuc.edu/mailman/listinfo/llvmdev>
>>>>>> >
>>>>>> >
>>>>>>
>>>>>>
>>>>>>
>>>>>
>>>>> --
>>>>>
>>>>> Thanks,
>>>>>
>>>>> Justin Holewinski
>>>>>
>>>>>
>>>>>
>>>>
>>>>
>>>>
>>>
>>> --
>>>
>>> Thanks,
>>>
>>> Justin Holewinski
>>>
>>>
>>
>>
>
>
> --
>
> Thanks,
>
> Justin Holewinski
>

Possibly Parallel Threads

Search for more possibly parallel threads

llvm dev - Feb 2013 - [LLVMdev] Problem with PTX assembly printing (NVPTX backend)

[LLVMdev] Problem with PTX assembly printing (NVPTX backend)

[LLVMdev] Problem with PTX assembly printing (NVPTX backend)

[LLVMdev] Problem with PTX assembly printing (NVPTX backend)

Possibly Parallel Threads