Displaying 20 results from an estimated 37 matches for "address_size".
2012 Sep 03
2
[LLVMdev] [NVPTX] Backend cannot handle array-of-arrays constant
...quot;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 0x00007f1bc9aa8445 gsignal...
2012 Sep 04
2
[LLVMdev] [NVPTX] Backend cannot handle array-of-arrays constant
...;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] =
> {0...
2012 Sep 04
0
[LLVMdev] [NVPTX] Backend cannot handle array-of-arrays constant
...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,0x0,0x54,0x48,0x...
2020 Sep 24
2
cuda __shfl_sync problem
...3 and according to
https://developer.nvidia.com/cuda-gpus#compute my GTX950 has Compute
Capability 5.2.
Also according to
https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#release-notes
PTX ISA 6.0 does support sm_52.
However llc generates:
.version 4.1
.target sm_52, debug
.address_size 64
Any ideas why this is happening? Or am i doing something wrong?
PS. I'm using CUDA 10, driver 440
~George
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-dev/attachments/20200924/4fa6c248/attachment.html>
2012 Sep 06
0
[LLVMdev] [NVPTX] Backend cannot handle array-of-arrays constant
...0",
>> [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
>>
>>...
2020 Sep 25
2
cuda __shfl_sync problem
...gt; Capability 5.2.
>>
>> Also according to
>> https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#release-notes
>> PTX ISA 6.0 does support sm_52.
>>
>> However llc generates:
>>
>> .version 4.1
>> .target sm_52, debug
>> .address_size 64
>>
>> Any ideas why this is happening? Or am i doing something wrong?
>>
>> PS. I'm using CUDA 10, driver 440
>>
>> ~George
>>
>>
>>
>> _______________________________________________
>> LLVM Developers mailing list
>> l...
2012 Jul 18
2
[LLVMdev] [NVPTX] PTXAS - Unimplemented feature: labels as initial values
...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 .align 4096 .b8 _2E_c...
2012 Jun 12
2
[LLVMdev] [NVPTX] For linkonce_odr NVPTX generates .weak, but even newest PTXAS can't handle it
...)], 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>;
.reg .f64 %fl<3...
2012 Jul 10
2
[LLVMdev] [NVPTX] CUDA inline PTX asm definitions scoping "{" "}" is broken
...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<396>;
....
2012 Jul 18
0
[LLVMdev] [NVPTX] PTXAS - Unimplemented feature: labels as initial values
...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 .align 4096 .b8 _2E_c...
2012 Jul 10
0
[LLVMdev] [NVPTX] CUDA inline PTX asm definitions scoping "{" "}" is broken
...t; }
>
> !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;
>...
2012 Jun 13
0
[LLVMdev] [NVPTX] For linkonce_odr NVPTX generates .weak, but even newest PTXAS can't handle it
...ce 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...
2012 Jul 10
1
[LLVMdev] [NVPTX] CUDA inline PTX asm definitions scoping "{" "}" is broken
...}
>
> !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;
> .re...
2012 Jul 11
2
[LLVMdev] [NVPTX] llc -march=nvptx64 -mcpu=sm_20 generates invalid zero align for device function params
...at @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 .pred %p<396&g...
2012 Dec 12
2
[PATCH v7 1/2] xen: unify domain locking in domctl code
These two patches were originally part of the XSM series that I have
posted, and remain prerequisites for that series. However, they are
independent of the XSM changes and are a useful simplification
regardless of the use of XSM.
The Acked-bys on these patches were provided before rebasing them over
the copyback changes in 26268:1b72138bddda, which had minor conflicts
that I resolved.
[PATCH
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/llvm/Target/TargetLo...
2012 Nov 09
0
[LLVMdev] [NVPTX] llc -march=nvptx64 -mcpu=sm_20 generates invalid zero align for device function params
...< " .align " << GVar->getAlignment();
Could you please review and commit? Do you think it needs a test case?
Thanks,
- D.
dmikushin at hp2:~/forge/align0> llc -march=nvptx64 -mcpu=sm_20 align0.ll -o -
//
// Generated by LLVM NVPTX Back-End
//
.version 3.1
.target sm_20
.address_size 64
// .globl __internal_dsmul
.visible .func __internal_dsmul(
.param .b64 __internal_dsmul_param_0,
.param .align 4 .b8 __internal_dsmul_param_1[8],
.param .align 4 .b8 __internal_dsmul_param_2[8]
) // @__internal_dsmul
{
.reg .pred %p<396>;
.reg...
2013 Mar 01
4
[LLVMdev] NVPTX CUDA_ERROR_NO_BINARY_FOR_GPU
..._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
>...
2012 May 16
0
[LLVMdev] NVPTX: __iAtomicCAS support ?
..._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/...
2011 May 17
1
[LLVMdev] [LLVMDev] Add not instruction to PTX backend
> Though, I have to agree with Dan on assessing whether the selection logic is
> needed. Do you have an example where the PTX back-end cannot generate code
> for some piece of LLVM IR because of the lack of 'not' selection?
Honestly, I don't have such example yet. Just want to try to implement
some instructions by myself. :p
Regards,
chenwj
--
Wei-Ren Chen (陳韋任)