Displaying 20 results from an estimated 23 matches for "texmode_independ".
2012 Sep 03
2
[LLVMdev] [NVPTX] Backend cannot handle array-of-arrays constant
...ON", [3 x
i8] c"TUE", [3 x i8] c"WED", [3 x i8] c"THU", [3 x i8] c"FRI", [3 x
i8] c"SAT", [3 x i8] c"SUN"], align 4096
$ llc -march="nvptx" test.ll -o -
//
// Generated by LLVM NVPTX Back-End
//
.version 3.0
.target sm_10, texmode_independent
.address_size 32
Unexpected Constant type
UNREACHABLE executed at
/home/marcusmae/rpmbuild/BUILD/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp:1903!
0 libLLVM-3.2svn.so 0x00007f1bcb71bf0e
1 libLLVM-3.2svn.so 0x00007f1bcb71bd0a
2 libpthread.so.0 0x00007f1bca33ccb0
3 libc.so.6 0x00007f1...
2012 Sep 04
2
[LLVMdev] [NVPTX] Backend cannot handle array-of-arrays constant
...\00", [4 x i8] c"THU\00", [4 x i8] c"FRI\00",
> [4 x i8] c"SAT\00", [4 x i8] c"SUN\00"], align 16
> $ llc -march="nvptx" dayofweek.ll -o -
> //
> // Generated by LLVM NVPTX Back-End
> //
>
> .version 3.0
> .target sm_10, texmode_independent
> .address_size 32
>
> Unexpected Constant type
> UNREACHABLE executed at
> /home/marcusmae/rpmbuild/BUILD/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp:1903!
>
> NVCC
> =====
>
> $ nvcc -c -keep dayofweek.cu
> $ cat dayofweek.ptx
>
> .global .align 1 ....
2012 Sep 04
0
[LLVMdev] [NVPTX] Backend cannot handle array-of-arrays constant
..."TUE\00", [4 x i8] c"WED\00", [4 x i8] c"THU\00", [4 x i8] c"FRI\00",
[4 x i8] c"SAT\00", [4 x i8] c"SUN\00"], align 16
$ llc -march="nvptx" dayofweek.ll -o -
//
// Generated by LLVM NVPTX Back-End
//
.version 3.0
.target sm_10, texmode_independent
.address_size 32
Unexpected Constant type
UNREACHABLE executed at
/home/marcusmae/rpmbuild/BUILD/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp:1903!
NVCC
=====
$ nvcc -c -keep dayofweek.cu
$ cat dayofweek.ptx
.global .align 1 .b8 yweek[28] =
{0x4d,0x4f,0x4e,0x0,0x54,0x55,0x45,0x0,0x57,0x45,0x44...
2012 Sep 06
0
[LLVMdev] [NVPTX] Backend cannot handle array-of-arrays constant
...0", [4 x i8] c"FRI\00",
>> [4 x i8] c"SAT\00", [4 x i8] c"SUN\00"], align 16
>> $ llc -march="nvptx" dayofweek.ll -o -
>> //
>> // Generated by LLVM NVPTX Back-End
>> //
>>
>> .version 3.0
>> .target sm_10, texmode_independent
>> .address_size 32
>>
>> Unexpected Constant type
>> UNREACHABLE executed at
>> /home/marcusmae/rpmbuild/BUILD/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp:1903!
>>
>> NVCC
>> =====
>>
>> $ nvcc -c -keep dayofweek.cu
>> $ cat dayofw...
2012 Jul 18
2
[LLVMdev] [NVPTX] PTXAS - Unimplemented feature: labels as initial values
...hidden constant [3 x i8] c"aa\00", align 4096
@.cst2 = hidden constant [26 x i8] c"Usage: %s <nx> <ny> <nz>\0A\00", align
4096
> 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
.visible .global .align 4096 .b8 _2E_cst[11] = {114, 101, 100, 117, 99,
101, 46, 102, 57, 48, 0};
.visible .global .align 4096 .u64 z = _2E_cst1;
.visible .global .align 4096 .b8 _2E_cst1[3] = {122, 122, 0};
.visible .global .align 4096 .u64 a = _2E_cst12;
.visible .global .ali...
2012 Jun 12
2
[LLVMdev] [NVPTX] For linkonce_odr NVPTX generates .weak, but even newest PTXAS can't handle it
...()* @_Z4testv to
i8*)], section "llvm.metadata"
define linkonce_odr ptx_device void @_Z4testv() nounwind inlinehint {
entry:
ret void
}
> llc -march=nvptx64 -mcpu=sm_20 test3.ll -o test3.ptx
> cat test3.ptx
//
// Generated by LLVM NVPTX Back-End
//
.version 3.0
.target sm_20, texmode_independent
.address_size 64
.weak _Z4testv
.func _Z4testv(
) // @_Z4testv
{
.reg .pred %p<396>;
.reg .s16 %rc<396>;
.reg .s16 %rs<396>;
.reg .s32 %r<396>;
.reg .s64 %rl<396>;
.reg .f32 %f<396>;
....
2012 Jul 10
2
[LLVMdev] [NVPTX] CUDA inline PTX asm definitions scoping "{" "}" is broken
...store i32 %1, i32* %result, align 4
%2 = load i32* %result, align 4
ret i32 %2
}
!0 = metadata !{i32 127, i32 132, i32 166, i32 200, 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&...
2012 Jul 18
0
[LLVMdev] [NVPTX] PTXAS - Unimplemented feature: labels as initial values
...hidden constant [3 x i8] c"aa\00", align 4096
@.cst2 = hidden constant [26 x i8] c"Usage: %s <nx> <ny> <nz>\0A\00", align 4096
> 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
.visible .global .align 4096 .b8 _2E_cst[11] = {114, 101, 100, 117, 99, 101, 46, 102, 57, 48, 0};
.visible .global .align 4096 .u64 z = _2E_cst1;
.visible .global .align 4096 .b8 _2E_cst1[3] = {122, 122, 0};
.visible .global .align 4096 .u64 a = _2E_cst12;
.visible .global .ali...
2012 Jul 10
0
[LLVMdev] [NVPTX] CUDA inline PTX asm definitions scoping "{" "}" is broken
...ign 4
> ret i32 %2
> }
>
> !0 = metadata !{i32 127, i32 132, i32 166, i32 200, 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 ....
2012 Jun 13
0
[LLVMdev] [NVPTX] For linkonce_odr NVPTX generates .weak, but even newest PTXAS can't handle it
...fine linkonce_odr ptx_device void @_Z4testv() nounwind inlinehint {
> entry:
> ret void
> }
>
> > llc -march=nvptx64 -mcpu=sm_20 test3.ll -o test3.ptx
> > cat test3.ptx
>
> //
> // Generated by LLVM NVPTX Back-End
> //
>
> .version 3.0
> .target sm_20, texmode_independent
> .address_size 64
>
>
> .weak _Z4testv
> .func _Z4testv(
>
> ) // @_Z4testv
> {
> .reg .pred %p<396>;
> .reg .s16 %rc<396>;
> .reg .s16 %rs<396>;
> .reg .s32 %r<396>;
>...
2012 Jul 10
1
[LLVMdev] [NVPTX] CUDA inline PTX asm definitions scoping "{" "}" is broken
...n 4
> ret i32 %2
> }
>
> !0 = metadata !{i32 127, i32 132, i32 166, i32 200, 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...
2012 May 16
2
[LLVMdev] NVPTX: __iAtomicCAS support ?
...; preds = %while.cond
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/llv...
2013 Mar 01
4
[LLVMdev] NVPTX CUDA_ERROR_NO_BINARY_FOR_GPU
...or: 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_retva...
2012 May 16
0
[LLVMdev] NVPTX: __iAtomicCAS support ?
...; }
>
> 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...
2013 Mar 01
0
[LLVMdev] NVPTX CUDA_ERROR_NO_BINARY_FOR_GPU
...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
> > > (
> > >
> &...
2013 Mar 01
0
[LLVMdev] NVPTX CUDA_ERROR_NO_BINARY_FOR_GPU
...39;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
>> > (
>> >
>> > )
>>...
2012 Jul 11
2
[LLVMdev] [NVPTX] llc -march=nvptx64 -mcpu=sm_20 generates invalid zero align for device function params
...eclare ptx_device float @llvm.nvvm.add.rn.f(float, float) nounwind readnone
declare ptx_device float @llvm.nvvm.mul.rn.f(float, float) nounwind readnone
> llc -march=nvptx64 -mcpu=sm_20 test.ll -o test.ptx
> cat test.ptx
//
// Generated by LLVM NVPTX Back-End
//
.version 3.0
.target sm_20, texmode_independent
.address_size 64
// .globl __internal_dsmul
.func __internal_dsmul(
.param .b64 __internal_dsmul_param_0,
.param .align 0 .b8 __internal_dsmul_param_1[8],
.param .align 0 .b8 __internal_dsmul_param_2[8]
) // @__internal_dsmul
{
.reg...
2013 Mar 01
1
[LLVMdev] NVPTX CUDA_ERROR_NO_BINARY_FOR_GPU
...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
> >>...
2013 Mar 01
2
[LLVMdev] NVPTX CUDA_ERROR_NO_BINARY_FOR_GPU
...are getting compiled correctly.
In addition, when I try 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_S...
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