similar to: BitcodeReader non explicit error

Hi! I would appreciate some feedback from someone with experience in SCEV/SE. D39346 tries to fix an issue in LV (PR34965) that exposes a limitation in SCEV/SE. The best solution to the LV issue might not be a fix at SCEV/SE level but we may want to report/address SCEV/SE limitation as well. For the snippet below, LV expects SE to return a SCEVAddRecExpr for %21. However, SE returns ((4 * (zext

Hi, After a few distractions I'm back to work on the llvm tgsi backend. I've added clang integration and I can now compile a simple opencl program to something which sort of looks like tgsi. You can find my latest work on this here: http://cgit.freedesktop.org/~jwrdegoede/llvm http://cgit.freedesktop.org/~jwrdegoede/clang (the latter may still need to sync) I've a little test

Thanks for the feedback, Sanjoy. > SCEV is fairly conservative around PHI nodes that aren't recurrences and aren't obviously equivalent to a min-max branch-phi idiom. Is that the limitation you're running into here? Yes, that's exactly the problem. The problematic PHI nodes (%bc.resume.val and %bc.resume.val1) aren't either recurrences or related to min-max idioms. I

Hi Tom, Thanks for taking the time to answer this. On 11-01-16 18:10, Tom Stellard wrote: > On Mon, Jan 11, 2016 at 12:07:14PM +0100, Hans de Goede wrote: >> Hi, >> >> After a few distractions I'm back to work on the llvm tgsi backend. I've >> added clang integration and I can now compile a simple opencl program >> to something which sort of looks like

Hi all, According to the SPIR 2.0 spec[1], the name of OpenCL builtins are mangled. However, when I compile OpenCl code with Clang 3.9 with the "spir64-unknown-unknown" target, Clang generates IR without mangling the builtins, e.g. for: __kernel void input_zip_int(__global int *in0) { *in0 = get_global_id(0); } clang generates: define spir_kernel void @input_zip_int(i32

Hello, In the following I tried 3 versions of an example in R help. Only the two first predict command work. After : library(splines) require(stats) 1) fm1 <- lm(weight ~ bs(height, df = 5), data = women) ht1 <- seq(57, 73, len = 200) ph1 <- predict(fm1, data.frame(height=ht1)) # OK plot(women, xlab = "Height (in)", ylab = "Weight (lb)") lines(ht1, ph1) 2)

+ 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 a lot. On Mon, Sep 12, 2016 at 1:42 PM, Liu, Yaxun (Sam) <Yaxun.Liu at amd.com> wrote: > If you use the default header file under clang/lib/Headers/opencl-c.h, > get_global_id will be mangled. > > > > If you want to declare get_global_id in your own header, add > __attribute__((overloadable)), then it will be mangled. > > > > Sam > > > >

<div dir='rtl'><div> <div dir="rtl"> <div dir="ltr">Could be... But the wierd thing here is if I change the array to be of size 256 and the index to be 'unsigned char', seems like there is no way to access the 'size' field throught "y->ptr[index]" (in your example) but clang still performs the re-read: <a

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

Is copy.0 semantically equivalent to copy.1 in the following example? define void @copy.0(i8 addrspace(1)* addrspace(1)* %src, i8 addrspace(1)* addrspace(1)* %dst) { entry: %val = load i8 addrspace(1)* addrspace(1)* %src store i8 addrspace(1)* %val, i8 addrspace(1)* addrspace(1)* %dst ret void } define void @copy.1(i8 addrspace(1)* addrspace(1)* %src, i8 addrspace(1)* addrspace(1)* %dst)

> On Dec 8, 2014, at 5:12 PM, Sanjoy Das <sanjoy at playingwithpointers.com> wrote: > > Partially answering my own question, in general these are not > equivalent because LLVM allows for pointers in different address > spaces to have different sizes. However, are they equivalent if > pointers in addrspace(1) have the same size as pointers in > addrspace(0)? > >

Hi Tom, Matt, I'm running into strange issues with the cos test (piglit generated_tests/cl/builtin/math/builtin-float-cos-1.0.generated.c) I have been seeing random failures (incorrect results) for some time and tried to investigate. the weird part is that the failures are not 100% reproducible, sometimes the tests pass, or partly pass (it's usually float8 and float16 subtests that

I am attempting to get indexing code generation working with my backend. However, it seems that the addresses being calculated is being multiplied by the width of the data type. define void @ test_input_index_constant_int(i32 addrspace(11)* %input, i32 addrspace(11)* %result) { entry: %input.addr = alloca i32 addrspace(11)* ; <i32 addrspace(11)**> [#uses=2]

Hello, With the following reduced test case, cmd "opt -always-inline t.ll" crashes after inlining. Notice that byval argument %a will be remapped to %1 below, and consequently produces an illegal store. %1 = alloca i32, align 4 store i32 * %1, i32 addrspace(1)** %a.addr, align 8 Looks like Inliner assumes that byval arguments are from address space 0. Or this is just a bug in inliner?

OK, I’ve hit one more existing regression test that I’m weary of: define void @test2_addrspacecast() { %A = alloca %T %B = alloca %T %a = addrspacecast %T* %A to i8 addrspace(1)* %b = addrspacecast %T* %B to i8 addrspace(1)* call void @llvm.memcpy.p1i8.p0i8.i64(i8 addrspace(1)* %a, i8* bitcast (%T* @G to i8*), i64 124, i32 4, i1 false) call void

Hello, LLVM devs. I have the following IR: %x = alloca i32, align 4 %p = alloca i32*, align 8 store i32* %x, i32** %p, align 8 Now I change module's data layout and run InferAddressSpacePass. This turns that piece of code into %x = alloca i32, align 4, addrspace(1) %p = alloca i32*, align 8, addrspace(1) store i32 addrspace(1)* %x, i32* addrspace(1)* %p, align 8 But Verifier

It is problematic when byval argument is not from address space 0. When the default alloca address space is 0, should we consider this IR illegal? define internal i32 @bar(i32 addrspace(1)* byval %a) alwaysinline From: Reid Kleckner [mailto:rnk at google.com] Sent: Tuesday, September 25, 2018 2:38 PM To: Pan, Wei <wei.pan at intel.com> Cc: llvm-dev <llvm-dev at lists.llvm.org>

On Fri, Mar 25, 2016 at 02:22:11PM -0400, Jan Vesely wrote: > Hi Tom, Matt, > > I'm working on a project that needs few coherent atomic operations (HSA > mode: load, store, compare-and-swap) for std::atomic_uint in HCC. > > the attached patch implements atomic compare and swap for SI+ > (untested). I tried to stay within what was available, but there are > few issues

We'd like to be able to vary the latency of our load instructions based on what address space is being loaded from. I was thinking I could do this by overriding getOperandLatency in our target, but I'm wondering if the addrspace info is available when instructions are scheduled? For example, I have this in our llvm IR: %0 = load i32 addrspace(4)* @answer, align 4 store i32 %0, i32*