similar to: [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

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

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

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

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; >

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

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.

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 >> >

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],

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

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

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?

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

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

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

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] =

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) { >

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

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:

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