similar to: [LLVMdev] [PATCH ]Add Subtarget ptx23

Hi, Justin Thansk, but I have a little concern though. I saw Chiou mentioned the compatibility issue of PTX on the mailing list. http://lists.cs.uiuc.edu/pipermail/llvmdev/2011-March/038654.html In my PTX.td patch, def FeaturePTX23 : SubtargetFeature<"ptx23", "PTXVersion", PTX_VERSION_2_3", "Use PTX Language Version

On Sat, May 7, 2011 at 4:17 AM, 陳韋任 <chenwj at iis.sinica.edu.tw> wrote: > Hi, Justin > > Thansk, but I have a little concern though. I saw Chiou mentioned the > compatibility issue of PTX on the mailing list. > http://lists.cs.uiuc.edu/pipermail/llvmdev/2011-March/038654.html > > In my PTX.td patch, > > def FeaturePTX23 : SubtargetFeature<"ptx23",

On Fri, May 6, 2011 at 5:37 AM, 陳韋任 <chenwj at iis.sinica.edu.tw> wrote: > Hi, Justin > > PTX version in CUDA 4.0 has changed from 2.2 to 2.3. > I add ptx23 subtarget and update a testcase. > > Is that O.K.? > Looks good. Committed in r130980. Thanks! In the future, patches should be sent to llvm-commits instead of llvmdev. > > Regards, > chenwj >

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

How do I make my headset work. Before I try it with anything else, I'd like to make it work with arecord and aplay. VLC and movie player can make sounds with my external speakers. though not in my headset yet. I've got an insignia NS-PAH5101 headset. As directed, the red connector is in the pink spot and the green connector in green spot in the back of my computer. arecord folloed by aplay

Hi Justin, There are some backward incompatible features of PTX; for example, special registers are redefined as v4i32 (they were v4i16) in PTX 2.0. And CUDA 4.0 was rolled out last week. I heard that some instructions are deprecated. I am not sure how stable (or unstable) PTX specification is. Do you have a rough assessment of its stability? If PTX specification is still fast evolving, I would

Hi all, Does PTX backend support llvm sitofp instruction? I failed to compile my llvm source when use llc -march=ptx32. The reason is that the source has a sitofp instruction. After i changed the instruction into uitofp manually, it passed. Thanks in advance, best, Yabin -------------- next part -------------- An HTML attachment was scrubbed... URL:

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, Tanya and I also prefer the extended TBAA solution as it naturally fits with LLVM. From my understanding of TBAA, it seems to provide the power to describe the relationship between address spaces for alias analysis, i.e., it can describe if two address spaces are disjoint or one may nest within another. For OpenCL, it is most useful to indicate that address spaces are disjoint from the

The problem I want to address in this discussion is the representation of OpenCL/CUDA memory spaces in LLVM IR. As support for OpenCL and CUDA mature within Clang, it is important that we provide a way to represent memory spaces in a way that is (1) sufficiently generic that other language front-ends can easily emit the needed annotations, and (2) sufficiently specific that LLVM optimization

Justin, Out of these options, I would take the metadata approach for AA support. This doesn't solve the problem of different frontend/backends choosing different address space representations for the same language, but is the correct approach for providing extra information to the optimizations. The issue about memory spaces in general is a little different. For example, based on the

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 website [2], LLVM [3], Clang and libclc websites [5] over there. I feel people might get lost among those websites. ;-) Here

I was hoping you would come to the collaborative, joint solution. I've been waiting for Clang to have a settled OpenCL implementation to start working on OpenCL. Dealing with 3 parallel projects would be just that, a pain in the rear. - Marc On 10/20/2011 03:45 AM, Ralf Karrenberg wrote: > Hi Carlos, > > On 10/20/11 9:54 AM, Carlos Sánchez de La Lama wrote: >>> The

On Thu, Oct 13, 2011 at 11:57 AM, Peter Collingbourne <peter at pcc.me.uk>wrote: > Hi Justin, > > Thanks for bringing this up, I think it's important to discuss > these issues here. > > On Thu, Oct 13, 2011 at 09:46:28AM -0400, Justin Holewinski wrote: > > It is becoming increasingly clear to me that LLVM address spaces are not > the > > general solution

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:

On Thu, Oct 13, 2011 at 04:21:23PM -0400, Justin Holewinski wrote: > On Thu, Oct 13, 2011 at 4:16 PM, Peter Collingbourne <peter at pcc.me.uk>wrote: > > > On Thu, Oct 13, 2011 at 06:59:47PM +0000, Villmow, Micah wrote: > > > Justin, > > > Out of these options, I would take the metadata approach for AA support. > > > > > > This doesn't

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