Displaying 17 results from an estimated 17 matches for "func_retval0".
2013 Mar 01
4
[LLVMdev] NVPTX CUDA_ERROR_NO_BINARY_FOR_GPU
...a 640M GPU.
> >
> > PTX Code (for a mandelbrot calculation):
> >
> > //
> > // Generated by LLVM NVPTX Back-End
> > //
> >
> > .version 3.1
> > .target sm_10, texmode_independent
> > .address_size 64
> >
> > .func (.reg .b32 func_retval0) INT_PTX_SREG_CTAID_X
> > (
> >
> > )
> > ;
> > .func (.reg .b32 func_retval0) INT_PTX_SREG_CTAID_Y
> > (
> >
> > )
> > ;
> > .func (.reg .b32 func_retval0) INT_PTX_SREG_TID_X
> > (
> >
> > )
> > ;
> > .fu...
2013 Mar 01
0
[LLVMdev] NVPTX CUDA_ERROR_NO_BINARY_FOR_GPU
...brot calculation):
> > >
> > > //
> > > // Generated by LLVM NVPTX Back-End
> > > //
> > >
> > > .version 3.1
> > > .target sm_10, texmode_independent
> > > .address_size 64
> > >
> > > .func (.reg .b32 func_retval0) INT_PTX_SREG_CTAID_X
> > > (
> > >
> > > )
> > > ;
> > > .func (.reg .b32 func_retval0) INT_PTX_SREG_CTAID_Y
> > > (
> > >
> > > )
> > > ;
> > > .func (.reg .b32 func_retval0) INT_PTX_SREG_TID_X
>...
2013 Mar 01
0
[LLVMdev] NVPTX CUDA_ERROR_NO_BINARY_FOR_GPU
...e (for a mandelbrot calculation):
>> >
>> > //
>> > // Generated by LLVM NVPTX Back-End
>> > //
>> >
>> > .version 3.1
>> > .target sm_10, texmode_independent
>> > .address_size 64
>> >
>> > .func (.reg .b32 func_retval0) INT_PTX_SREG_CTAID_X
>> > (
>> >
>> > )
>> > ;
>> > .func (.reg .b32 func_retval0) INT_PTX_SREG_CTAID_Y
>> > (
>> >
>> > )
>> > ;
>> > .func (.reg .b32 func_retval0) INT_PTX_SREG_TID_X
>> > (
>...
2013 Mar 01
1
[LLVMdev] NVPTX CUDA_ERROR_NO_BINARY_FOR_GPU
...t;
> >> > //
> >> > // Generated by LLVM NVPTX Back-End
> >> > //
> >> >
> >> > .version 3.1
> >> > .target sm_10, texmode_independent
> >> > .address_size 64
> >> >
> >> > .func (.reg .b32 func_retval0) INT_PTX_SREG_CTAID_X
> >> > (
> >> >
> >> > )
> >> > ;
> >> > .func (.reg .b32 func_retval0) INT_PTX_SREG_CTAID_Y
> >> > (
> >> >
> >> > )
> >> > ;
> >> > .func (.reg .b32 fu...
2013 Mar 01
2
[LLVMdev] NVPTX CUDA_ERROR_NO_BINARY_FOR_GPU
...y load the module using CUDA, I get an
error: CUDA_ERROR_NO_BINARY_FOR_GPU. I'm running this on a 2012 MBP with
a 640M GPU.
PTX Code (for a mandelbrot calculation):
//
// Generated by LLVM NVPTX Back-End
//
.version 3.1
.target sm_10, texmode_independent
.address_size 64
.func (.reg .b32 func_retval0) INT_PTX_SREG_CTAID_X
(
)
;
.func (.reg .b32 func_retval0) INT_PTX_SREG_CTAID_Y
(
)
;
.func (.reg .b32 func_retval0) INT_PTX_SREG_TID_X
(
)
;
.func (.reg .b32 func_retval0) INT_PTX_SREG_NTID_X
(
)
;
.func (.reg .b32 func_retval0) INT_PTX_SREG_NTID_Y
(
)
;
// .globl examples_...
2013 Mar 01
0
[LLVMdev] NVPTX CUDA_ERROR_NO_BINARY_FOR_GPU
...RY_FOR_GPU. I'm running this on a 2012 MBP
> with a 640M GPU.
>
> PTX Code (for a mandelbrot calculation):
>
> //
> // Generated by LLVM NVPTX Back-End
> //
>
> .version 3.1
> .target sm_10, texmode_independent
> .address_size 64
>
> .func聽 (.reg .b32 func_retval0) INT_PTX_SREG_CTAID_X
> (
>
> )
> ;
> .func聽 (.reg .b32 func_retval0) INT_PTX_SREG_CTAID_Y
> (
>
> )
> ;
> .func聽 (.reg .b32 func_retval0) INT_PTX_SREG_TID_X
> (
>
> )
> ;
> .func聽 (.reg .b32 func_retval0) INT_PTX_SREG_NTID_X
> (
>
> )
&g...
2012 Jul 10
2
[LLVMdev] [NVPTX] CUDA inline PTX asm definitions scoping "{" "}" is broken
...i32 242, i32 285, i32
327}
> llc -march=nvptx64 test.ll -o test.ptx
> cat test.ptx
//
// Generated by LLVM NVPTX Back-End
//
.version 3.0
.target sm_10, texmode_independent
.address_size 64
// .globl _Z5__anyi
.visible .global .align 4 .b8 __local_depot0[8];
.func (.reg .b32 func_retval0) _Z5__anyi(
.reg .b32 _Z5__anyi_param_0
) // @_Z5__anyi
{
.reg .b64 %SP;
.reg .b64 %SPL;
.reg .pred %p<396>;
.reg .s16 %rc<396>;
.reg .s16 %rs<396>;
.reg .s32 %r<396>;
.reg .s64 %rl<396>;...
2012 Jul 10
0
[LLVMdev] [NVPTX] CUDA inline PTX asm definitions scoping "{" "}" is broken
...tx
> > cat test.ptx
> //
> // Generated by LLVM NVPTX Back-End
> //
>
> .version 3.0
> .target sm_10, texmode_independent
> .address_size 64
>
>
> // .globl _Z5__anyi
> .visible .global .align 4 .b8 __local_depot0[8];
>
> .func (.reg .b32 func_retval0) _Z5__anyi(
> .reg .b32 _Z5__anyi_param_0
> ) // @_Z5__anyi
> {
> .reg .b64 %SP;
> .reg .b64 %SPL;
> .reg .pred %p<396>;
> .reg .s16 %rc<396>;
> .reg .s16 %rs<396>;
> .reg .s32 %r&...
2012 Jul 10
1
[LLVMdev] [NVPTX] CUDA inline PTX asm definitions scoping "{" "}" is broken
...st.ptx
> > cat test.ptx
> //
> // Generated by LLVM NVPTX Back-End
> //
>
> .version 3.0
> .target sm_10, texmode_independent
> .address_size 64
>
>
> // .globl _Z5__anyi
> .visible .global .align 4 .b8 __local_depot0[8];
>
> .func (.reg .b32 func_retval0) _Z5__anyi(
> .reg .b32 _Z5__anyi_param_0
> ) // @_Z5__anyi
> {
> .reg .b64 %SP;
> .reg .b64 %SPL;
> .reg .pred %p<396>;
> .reg .s16 %rc<396>;
> .reg .s16 %rs<396>;
> .reg .s32 %r&...
2016 Apr 07
2
[GPUCC] how to remove _ZL21__nvvm_reflect_anchorv() automatically?
...complaint:
*ptxas c.s, line 171; error : Duplicate definition of function
'_ZL21__nvvm_reflect_anchorv'*
*ptxas c.s, line 171; fatal : Parsing error near '.2': syntax error*
So I inspected c.s and found the issue above was caused by the following
line:
.*func** (.param .b32 func_retval0) _ZL21__nvvm_reflect_anchorv.2() //
@_ZL21__nvvm_reflect_anchorv.2*
After I manually deleted the definition of this function in c.s, the
compilation works file. I wonder how could I force llc to remove *`*
_ZL21__nvvm_reflect_anchorv.2*()`*? Or is that possible to prevent
_ZL21__nvvm_reflect_anc...
2016 Apr 08
2
[GPUCC] how to remove _ZL21__nvvm_reflect_anchorv() automatically?
...unction
>> '_ZL21__nvvm_reflect_anchorv'*
>>
>> *ptxas c.s, line 171; fatal : Parsing error near '.2': syntax error*
>>
>> So I inspected c.s and found the issue above was caused by the following
>> line:
>>
>> .*func** (.param .b32 func_retval0) _ZL21__nvvm_reflect_anchorv.2() //
>> @_ZL21__nvvm_reflect_anchorv.2*
>>
>> After I manually deleted the definition of this function in c.s, the
>> compilation works file. I wonder how could I force llc to remove *`*
>> _ZL21__nvvm_reflect_anchorv.2*()`*? Or is that...
2015 Feb 03
2
[LLVMdev] Example for usage of LLVM/Clang/libclc
...if the
PTX code I am generating is correct (is the one that is supposed to be
generated).
For example, currently,
In OpenCL : get_global_id(0) translates to
In LLVM : %call = tail call i32 @get_global_id(i32 0) which translates
to
In PTX:
// .globl blur2d
.func (.param .b32 func_retval0) get_global_id
(
.param .b32 get_global_id_param_0
)
;
mov.u32 %r2, 0;
.param .b32 param0;
st.param.b32 [param0+0], %r2;
.param .b32 retval0;
call.uni (retval0),
get_global_id,
(
param0
);
Is this what...
2016 Apr 09
2
[GPUCC] how to remove _ZL21__nvvm_reflect_anchorv() automatically?
...>>>>
>>>> *ptxas c.s, line 171; fatal : Parsing error near '.2': syntax error*
>>>>
>>>> So I inspected c.s and found the issue above was caused by the
>>>> following line:
>>>>
>>>> .*func** (.param .b32 func_retval0) _ZL21__nvvm_reflect_anchorv.2()
>>>> // @_ZL21__nvvm_reflect_anchorv.2*
>>>>
>>>> After I manually deleted the definition of this function in c.s, the
>>>> compilation works file. I wonder how could I force llc to remove *`*
>>>> _ZL21_...
2012 May 16
2
[LLVMdev] NVPTX: __iAtomicCAS support ?
...ret void
}
declare ptx_device i32 @_Z12__iAtomicCASPiii(i32*, i32, i32)
CODEGEN
=========
dmikushin at hp2:~> llc < kernelgen_monitor.ll -march=nvptx -mcpu=sm_20
//
// Generated by LLVM NVPTX Back-End
//
.version 3.0
.target sm_20, texmode_independent
.address_size 32
.func (.param .b32 func_retval0) _Z12__iAtomicCASPiii
(
.param .b32 _Z12__iAtomicCASPiii_param_0,
.param .b32 _Z12__iAtomicCASPiii_param_1,
.param .b32 _Z12__iAtomicCASPiii_param_2
)
;
Not Implemented
UNREACHABLE executed at
/tmp/rpmbuild_debug/BUILD/llvm/build/include/llvm/Target/TargetLowering.h:1249!
0 libLLVM-3.2svn.so 0...
2013 Feb 04
0
[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
2012 May 16
0
[LLVMdev] NVPTX: __iAtomicCAS support ?
...2)
>
> CODEGEN
> =========
>
> dmikushin at hp2:~> llc < kernelgen_monitor.ll -march=nvptx -mcpu=sm_20
> //
> // Generated by LLVM NVPTX Back-End
> //
>
> .version 3.0
> .target sm_20, texmode_independent
> .address_size 32
>
> .func (.param .b32 func_retval0) _Z12__iAtomicCASPiii
> (
> .param .b32 _Z12__iAtomicCASPiii_param_0,
> .param .b32 _Z12__iAtomicCASPiii_param_1,
> .param .b32 _Z12__iAtomicCASPiii_param_2
> )
> ;
>
> Not Implemented
> UNREACHABLE executed at
> /tmp/rpmbuild_debug/BUILD/llvm/build/include/llvm/Ta...
2013 Feb 04
2
[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: