similar to: [RFC] Add TargetTransformInfo::isAllocaPtrValueNonZero and let ValueTracking depend on TargetTransformInfo

Hal, Thanks for your suggestion. I think that makes sense. Currently, non-zero alloca address space is already represented by data layout, e.g., the last component of the data layout of amdgcn---amdgiz target is -A5, which means alloca is in address space 5. How about adding a letter z to -A5 to indicate alloca may have zero value? i.e. -A5 means alloca is in address space 5 and always has

> On Dec 14, 2017, at 20:28, Hal Finkel via llvm-dev <llvm-dev at lists.llvm.org> wrote: > > Would that address your use case? Or can you have null dereferenceable pointers in that address space, just not ones from alloca? I would like to clarify what “null” means exactly. One related thing I would like in the future is for the DataLayout to specify what numeric value is the

Hi dev list, I am currently exploring the integration of AMDGPU/ROCm into the PACXX project and observing some strange behavior of the AMDGPU backend. The following IR is generated for a simple address space test that copies from global to shared memory and back to global after a barrier synchronization. Here is the IR is attached as as1.ll The output is as follows: 0 0 0 0 0 0 0 0 0 0 0 0 0

Hello, As I can see, recently LLVM switched to using assume bundles to encode alignment information: https://reviews.llvm.org/rG78de7297abe2e8fa782682168989c70e3cb34a5c However, it seems that the ValueTracking cannot understand the new format. As an example, consider compilation of the following reproducer with clang-11 (old assume format) and clang-trunk (assume bundles): #include

Hi, i am playing the ScalarEvolution these days. i found the the ScalarEvolution will simply return a SCEVUnknow for a ICmpInst, so i think maybe great to add a new kind of SCEV to the ScalarEvolution framework. for example, if i run ScalarEvolution on the bc file generate from the following C source file: int f(int a, int b, int c, int d) { return (2 * a + 5 * c + 2) > (4 * d - 3*b

On Tue, Jun 17, 2014 at 2:33 PM, Jingyue Wu <jingyue at google.com> wrote: > Hi Eric, > > In the IR, besides "target datalayout" and "target triple", we have a > special "target cpu" string which is set by the Clang front-end according to > its -target-cpu flag. We also write a Module::getTargetCPU() method to > retrieve this string from the

Chandler Carruth wrote: > This seems fine to me, but I'd like to make sure it looks OK to Nick as > well. I strongly prefer baking in knowledge about the intrinsics themselves into the passes if possible. Metadata will always be secondary. Separately, should value tracking look use range metadata when it's available? Absolutely. I think it should apply to all CallInst not just

On 17 June 2014 06:41, Eli Bendersky <eliben at google.com> wrote: > On Tue, Jun 17, 2014 at 1:38 AM, Nick Lewycky <nicholas at mxc.ca> wrote: > >> Chandler Carruth wrote: >> >>> This seems fine to me, but I'd like to make sure it looks OK to Nick as >>> well. >>> >> >> I strongly prefer baking in knowledge about the

Eh? How do you envision this? -eric On Tue, Jun 17, 2014 at 2:09 PM, Jingyue Wu <jingyue at google.com> wrote: > Hi Nick, > > That makes sense. I think a main issue here is that the ranges of these PTX > special registers (e.g., threadIdx.x) depend on -target-cpu which is only > visible to clang and llc. Would you mind we specify "target cpu" in the IR > similar

Be careful about oversimplifying signed integer comparisons -- integer arithmetic can easily overflow, so you cannot transform A > B to A - B > 0. The compare instructions in most processors do not simply subtract and test the most significant bit; they compute what the sign of the difference would be in extended precision. On Apr 17, 2010, at 1:00 PM, llvmdev-request at cs.uiuc.edu wrote:

On 06/14/2017 05:05 PM, Connor Abbott wrote: > On Tue, Jun 13, 2017 at 6:13 PM, Tom Stellard <tstellar at redhat.com> wrote: >> On 06/13/2017 07:33 PM, Matt Arsenault wrote: >>> >>>> On Jun 12, 2017, at 17:23, Tom Stellard <tstellar at redhat.com <mailto:tstellar at redhat.com>> wrote: >>>> >>>> On 06/12/2017 08:03 PM, Connor

I don't see any problem mangling it to be honest even though there seems to be only one prototype anyways. We could add restrict in as well. Cheers, Anastasia ________________________________ From: Hongbin Zheng <etherzhhb at gmail.com> Sent: 17 September 2016 05:32:54 To: Liu, Yaxun (Sam) Cc: cfe-dev at lists.llvm.org; llvm-dev; Bader, Alexey (alexey.bader at intel.com); Anastasia

+ Alexey Anastasia According to SPIR spec v1.2 s2.10.3 2.10.3 The printf function The printf function is supported, and is mangled according to its prototype as follows: int printf(constant char * restrict fmt, ... ) Note that the ellipsis formal argument (...) is mangled to argument type specifier z It seems printf should be mangled. Alexey/Anastasia, What do you think? Thanks. Sam From:

Thanks Chris for the comment. Since there is no objection, I attached a new patch which enables response file support and removes the argument for controlling/disabling response file support. The patch also contains a simple test. I did regression check and there are no regressions. + llvmdev Sam From: Chris Lattner [mailto:clattner at apple.com] Sent: Sunday, September 30, 2012 2:20 PM To:

On Thu, Jun 18, 2015 at 10:26 AM, Mehdi Amini <mehdi.amini at apple.com> wrote: > > On Jun 18, 2015, at 9:31 AM, Liu, Yaxun (Sam) <Yaxun.Liu at amd.com> wrote: > > > > *From:* Mehdi Amini [mailto:mehdi.amini at apple.com <mehdi.amini at apple.com>] > > *Sent:* Thursday, June 18, 2015 11:24 AM > *To:* Liu, Yaxun (Sam) > *Cc:* llvmdev at cs.uiuc.edu

Dear Developers, I compiled a OpenCL kernel before (on Nov. last year) like __kernel void g(__global float* array) { array[get_global_id(0)] = 1; } with libclc, which would originally use the instrinsics like llvm.r600.read.local.size.x(). I executed the generated object file with one version of the hsa-runtime [1] provided by Mr. Stellard, when there was more than one workgroup, the output

+ llvm-dev From: Sumner, Brian Sent: Tuesday, May 10, 2016 3:11 PM To: Anastasia Stulova <Anastasia.Stulova at arm.com>; Liu, Yaxun (Sam) <Yaxun.Liu at amd.com>; cfe-dev (cfe-dev at lists.llvm.org) <cfe-dev at lists.llvm.org>; Pan, Xiuli <xiuli.pan at intel.com>; Bader, Alexey (alexey.bader at intel.com) <alexey.bader at intel.com> Cc: Stellard, Thomas

So we have been in discussions within the Khronos SPIR-V work group on our push to get our SPIR-V code into tip LLVM and have drawn the following conclusions; * We absolutely must create a fully fledged backend that uses all the machinery that target backends are expected to use. * We probably have to split out the SPIR-V -> LLVM IR into a separate project from LLVM ala Clang et

On 06/13/2017 07:33 PM, Matt Arsenault wrote: > >> On Jun 12, 2017, at 17:23, Tom Stellard <tstellar at redhat.com <mailto:tstellar at redhat.com>> wrote: >> >> On 06/12/2017 08:03 PM, Connor Abbott wrote: >>> On Mon, Jun 12, 2017 at 4:56 PM, Tom Stellard <tstellar at redhat.com <mailto:tstellar at redhat.com>> wrote: >>>> On

Hey Tom, Really it was at the behest of the replies - we got a lot of feedback from the mailing list that indicated we'd be putting extra workload of people changing features of the IR if we didn't follow the same mechanisms of the other backends (mostly led by Chandler's very astute comments on the subject). Cheers, -Neil. On 07/07/15 14:43, Tom Stellard wrote: > On Tue, Jul