search for: func_retval0

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 &gt...
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 >> > ( &gt...
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: