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

[nkavv at physics.auth.gr]

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
>