similar to: [LLVMdev] PTX builtin functions.

Displaying 20 results from an estimated 5000 matches similar to: "[LLVMdev] PTX builtin functions."

2011 Nov 21
1
[LLVMdev] PTX builtin functions.
On Mon, Nov 21, 2011 at 7:01 AM, Alberto Magni <alberto.magni86 at gmail.com>wrote: > Hi Justin, > > attached you find the patch for the integer max instruction. > The multiclass PTX_INTRINSIC_INT3 in file PTXIntrinsicInstrInfo.td > is almost an exact copy of PTX_INT3 in PTXInstrInfo.td, maybe > a modification of this class can be defined in a separate file. > I'm
2011 Nov 21
2
[LLVMdev] PTX builtin functions.
On Mon, Nov 21, 2011 at 11:45 AM, Alberto Magni <alberto.magni86 at gmail.com>wrote: > On Mon, Nov 21, 2011 at 3:36 PM, Justin Holewinski > <justin.holewinski at gmail.com> wrote: > > On Mon, Nov 21, 2011 at 7:01 AM, Alberto Magni < > alberto.magni86 at gmail.com> > > wrote: > >> > >> Hi Justin, > >> > >> attached you find
2011 Nov 22
2
[LLVMdev] PTX builtin functions.
Alberto, The AMDIL backend solves your problem with intrinsic overloading this way: def int_AMDIL_mad : GCCBuiltin<"__amdil_mad">, TernaryIntFloat; Where TernaryIntFloat is defined as: class TernaryIntFloat : Intrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>], []>; This allows us to write a
2011 Nov 21
0
[LLVMdev] PTX builtin functions.
On Mon, Nov 21, 2011 at 3:36 PM, Justin Holewinski <justin.holewinski at gmail.com> wrote: > On Mon, Nov 21, 2011 at 7:01 AM, Alberto Magni <alberto.magni86 at gmail.com> > wrote: >> >> Hi Justin, >> >> attached you find the patch for the integer max instruction. >> The multiclass PTX_INTRINSIC_INT3 in file PTXIntrinsicInstrInfo.td >> is almost
2011 Nov 23
2
[LLVMdev] PTX builtin functions.
On Nov 23, 2011 6:57 AM, "Alberto Magni" <alberto.magni86 at gmail.com> wrote: > > On Tue, Nov 22, 2011 at 5:01 PM, Villmow, Micah <Micah.Villmow at amd.com> wrote: > > Alberto, > > The AMDIL backend solves your problem with intrinsic overloading this way: > > def int_AMDIL_mad : GCCBuiltin<"__amdil_mad">, TernaryIntFloat; >
2011 Dec 08
3
[LLVMdev] PTX builtin functions.
It is my understanding that all you need to do is specify let isTarget = 1 in your .td file and it will generate target specific intrinsics. This should allow you to keep the IntrinsicsPTX.td file in the same location. Micah From: Justin Holewinski [mailto:justin.holewinski at gmail.com] Sent: Monday, December 05, 2011 6:13 AM To: Alberto Magni Cc: Villmow, Micah; LLVM Developers Mailing List
2011 Dec 04
2
[LLVMdev] PTX builtin functions.
Hi Justin, sorry for the delay, I have been busy. Micah's proposal requires to move the definitions of the intrinsics from include/llvm/IntrinsicsPTX.td to lib/Target/PTX/PTXIntrinsics.td thus allowing the generation of the file PTXGenIntrinsics.inc which will be included by PTXIntrinsicInfo.cpp. This is a quite big modification, do you agree with this ? Or do you have a better solution.
2011 Nov 22
0
[LLVMdev] PTX builtin functions.
On Mon, Nov 21, 2011 at 5:31 PM, Justin Holewinski <justin.holewinski at gmail.com> wrote: > On Mon, Nov 21, 2011 at 11:45 AM, Alberto Magni <alberto.magni86 at gmail.com> > wrote: >> >> On Mon, Nov 21, 2011 at 3:36 PM, Justin Holewinski >> <justin.holewinski at gmail.com> wrote: >> > On Mon, Nov 21, 2011 at 7:01 AM, Alberto Magni >> >
2011 Nov 23
0
[LLVMdev] PTX builtin functions.
On Tue, Nov 22, 2011 at 5:01 PM, Villmow, Micah <Micah.Villmow at amd.com> wrote: > Alberto, >  The AMDIL backend solves your problem with intrinsic overloading this way: > def int_AMDIL_mad     : GCCBuiltin<"__amdil_mad">, TernaryIntFloat; > > Where TernaryIntFloat is defined as: > class TernaryIntFloat : >          Intrinsic<[llvm_anyfloat_ty],
2011 Nov 23
0
[LLVMdev] PTX builtin functions.
On Nov 23, 2011 8:33 AM, "Justin Holewinski" <justin.holewinski at gmail.com> wrote: > > > On Nov 23, 2011 6:57 AM, "Alberto Magni" <alberto.magni86 at gmail.com> wrote: > > > > On Tue, Nov 22, 2011 at 5:01 PM, Villmow, Micah <Micah.Villmow at amd.com> wrote: > > > Alberto, > > > The AMDIL backend solves your problem
2011 Dec 05
0
[LLVMdev] PTX builtin functions.
On Sun, Dec 4, 2011 at 1:10 PM, Alberto Magni <alberto.magni86 at gmail.com>wrote: > Hi Justin, > > sorry for the delay, I have been busy. > > Micah's proposal requires to move the definitions of the intrinsics > from include/llvm/IntrinsicsPTX.td to lib/Target/PTX/PTXIntrinsics.td > thus allowing the generation of the file PTXGenIntrinsics.inc which > will be
2011 Dec 08
0
[LLVMdev] PTX builtin functions.
On Thu, Dec 8, 2011 at 11:36 AM, Villmow, Micah <Micah.Villmow at amd.com>wrote: > It is my understanding that all you need to do is specify let isTarget = > 1 in your .td file and it will generate target specific intrinsics. This > should allow you to keep the IntrinsicsPTX.td file in the same location. > So we keep the intrinsics defined in include/llvm/IntrinsicsPTX.td?
2011 Nov 16
0
[LLVMdev] PTX builtin functions.
On Wed, Nov 16, 2011 at 8:05 AM, Alberto Magni <alberto.magni86 at gmail.com>wrote: > Dear Justin, > > I am trying to add the support for some OpenCL builtin functions to > the PTX backend. > The attached file represent the first stub of a patch for the fmax > builtin function. > First off, thanks for helping to improve the PTX back-end! There are really two main
2011 Oct 24
1
[LLVMdev] Function pointer parameters in PTX backend
Hi everybody, I am trying to produce ptx code starting from OpenCL C. I am experiencing a problem concerning pointer parameters. Here follows an example: kernel void function(__global float* parameter1) {} NVIDIA NVCC Compiler: .entry function( .param .u32 *.ptr* .global .align 4 function_param_0 ) { ret; } CLANG + LLVM PTX backend // (skipping builtin functions definitions) .entry
2011 Nov 19
1
[LLVMdev] llvm_anyint_ty clarification
Hello everyone, I am trying to implement the max PTX builtin function. This is defined in the following way: "max.type d, a, b;" where .type can be: .type = { .u16, .u32, .u64, .s16, .s32, .s64 }; The presence of multiple types requires llvm.ptx.max to be overloaded for i16, i32 and i64. So I think that the right way to define the intrinsic would be (as in the
2011 Nov 14
2
[LLVMdev] PTX backend fatal error
Hi everybody, I am testing the PTX backend using the OpenCL NVIDIA SDK benchmarks. Compiling the Histogram64.cl program I get a several backend errors. I isolated one of them in the following kernel program: __kernel void kernel_function(__global int *input) { __local char localArray[16]; for(unsigned int index = 0; index < 16; ++index) localArray[index] = 0; input[0] =
2011 Nov 14
0
[LLVMdev] PTX backend fatal error
On Mon, Nov 14, 2011 at 8:57 AM, Alberto Magni <alberto.magni86 at gmail.com>wrote: > Hi everybody, > > I am testing the PTX backend using the OpenCL NVIDIA SDK benchmarks. > Compiling the Histogram64.cl program I get a several backend errors. > > I isolated one of them in the following kernel program: > > __kernel void kernel_function(__global int *input) { >
2011 Nov 14
1
[LLVMdev] PTX backend fatal error
Justin, Add this to your TargetLowering constructor, this fixes the mem* issue. maxStoresPerMemcpy = 4096; maxStoresPerMemmove = 4096; maxStoresPerMemset = 4096; From: llvmdev-bounces at cs.uiuc.edu [mailto:llvmdev-bounces at cs.uiuc.edu] On Behalf Of Justin Holewinski Sent: Monday, November 14, 2011 7:12 AM To: Alberto Magni Cc: llvmdev at cs.uiuc.edu Subject: Re: [LLVMdev] PTX backend
2011 May 11
3
[LLVMdev] [LLVMDev] Add not instruction to PTX backend
Hi, Dan I add "not" instruction support in PTXInstrInfo.td as you suggested before. multiclass PTX_LOGIC_2OP<string opcstr,PatFrag opnode> { ... } Now I am trying to write test case for logic and shift operations. But I have a trouble in mapping LLVM IR to PTX IR for "not" instruction. The test case I wrote is, define ptx_device i16 @t4_u16(i16 %x) { ; CHECK:
2011 Dec 13
3
[LLVMdev] Changes to the PTX calling conventions
From: Justin Holewinski [mailto:justin.holewinski at gmail.com] Sent: Tuesday, December 13, 2011 10:50 AM To: Villmow, Micah Cc: LLVM Developers Mailing List Subject: Re: [LLVMdev] Changes to the PTX calling conventions On Tue, Dec 13, 2011 at 12:54 PM, Villmow, Micah <Micah.Villmow at amd.com<mailto:Micah.Villmow at amd.com>> wrote: From: Justin Holewinski [mailto:justin.holewinski