similar to: [LLVMdev] Build PTX samples with LLVM/Clang/libclc

On Tue, Dec 6, 2011 at 10:17 PM, 陳韋任 <chenwj at iis.sinica.edu.tw> wrote: > Hi Justin, > > I download llvm-ptx-samples [1] and try to build them. I found it seems > lack > of a complete document on how to build them with LLVM/Clang/libclc. Do you > think > it's a good idea to put a complete document/tutorial in _one_ place? > Currently, > there are your

Hi, My goal is to use Clang/LLVM/libclc to compile an OpenCL kernel and eventually generate a PTX code. I already did this but I am not sure 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:

Hi, This is to announce the availability of libclc, an open source, BSD licensed implementation of the library requirements of the OpenCL C programming language, as specified by the OpenCL 1.1 Specification. libclc is intended to be used with Clang's OpenCL frontend. libclc website: http://www.pcc.me.uk/~peter/libclc/ libclc is designed to be portable and extensible. To this end, it

Do we have a list of these open-source LLVM-based OpenCL projects somewhere? Off the top of my head, we have: libclc: http://www.pcc.me.uk/~peter/libclc/ pocl: https://launchpad.net/pocl clover: http://cgit.freedesktop.org/~steckdenis/clover/ (I think that all of these have BSD- or MIT-style licenses). Are there any others? -Hal On Wed, 2011-10-19 at 14:47 +0100, Peter Collingbourne wrote:

Hi Ralf, > The project started as a use-case for our "Whole-Function Vectorization" > library, which allows to transform a function to compute the same as W > executions of the original code by using SIMD instructions (W = 4 for > SSE/AltiVec, 8 for AVX). Quite interesting. We were planning to add "vectorization" to our passes also, but if I understood the

Ralf, What version of the SDK were you using for your analysis? I don't see that in the slides/pdf. Thanks, Micah > -----Original Message----- > From: llvmdev-bounces at cs.uiuc.edu [mailto:llvmdev-bounces at cs.uiuc.edu] > On Behalf Of Ralf Karrenberg > Sent: Wednesday, October 19, 2011 2:13 PM > To: llvmdev at cs.uiuc.edu > Subject: Re: [LLVMdev] ANN: libclc (OpenCL C

Hi Micah, The numbers from the paper were measured with the ATI Stream SDK v2.1 (it's only mentioned in the references I think). The most recent measurements I have were done with the current v2.5. Best, Ralf Am 19.10.2011 23:43, schrieb Villmow, Micah: > Ralf, > What version of the SDK were you using for your analysis? I don't see that in the slides/pdf. > > Thanks, >

Hi everybody, the compiler design lab at Saarland University (chair of Sebastian Hack) is also working on an LLVM-based OpenCL driver. The project started as a use-case for our "Whole-Function Vectorization" library, which allows to transform a function to compute the same as W executions of the original code by using SIMD instructions (W = 4 for SSE/AltiVec, 8 for AVX). The

Hi Carlos, On 10/20/11 9:54 AM, Carlos Sánchez de La Lama wrote: >> The project started as a use-case for our "Whole-Function Vectorization" >> library, which allows to transform a function to compute the same as W >> executions of the original code by using SIMD instructions (W = 4 for >> SSE/AltiVec, 8 for AVX). > > Quite interesting. We were planning to

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

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

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

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

Hello Bergström/Eric, Thanks for the reply. The G80(sm_10) architecture was ported on FPGA by a group of researchers (http://www.ecs.umass.edu/ece/tessier/andryc-fpt13.pdf). Our group have some further research interest on this work. I was working on modifying the Clang-LLVM for a couple of months and achieved the required changes. But Clang-LLVM is only allowing me to generate PTX for sm_20,

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