search for: maxstorespermemcpy

I had already adjusted MaxStoresPerMemcpy to my preferred value, and this works great but for the cases where load/stores are used on non-size aligned structs the odd behaviour still happens. For my 3-char, 3-byte struct test, the memcpy replacement appears to consist on a single byte load and store of the last char (this is correct), foll...

Hi Tim, Thanks for your reply. That’s what I was afraid of. I will try on the cfe-list as you suggest though. The reason I want structs to be aligned/padded to 2 bytes is because my architecture only has 16 bit operations. I can read (sign and zero extended) and write (truncated) 8 bit data from/to memory, but all intermediate operations in registers are performed in 16 bit registers. This

Similar to how `operator new(size)` can be overridden during linking, to capture and customize all memory allocation operations, I'd like to be able to capture all scenarios where memory is copied or moved. I've tried to use `--wrap memcpy` with some success, however I've noticed that in some instances (e.g. copying a small struct) the memcpy instruction in the IR is replaced with a

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

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

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

You'll need to do the work then. I'd also question why? On most platforms a decent memcpy exists. -eric On Mar 5, 2012, at 12:28 PM, Ryan Taylor <ryta1203 at gmail.com> wrote: > I would like it to always be lowered, I don't want it. > > On Mon, Mar 5, 2012 at 12:27 PM, Eric Christopher <echristo at apple.com> wrote: > You don't have memcpy or want it to

I would like it to always be lowered, I don't want it. On Mon, Mar 5, 2012 at 12:27 PM, Eric Christopher <echristo at apple.com>wrote: > You don't have memcpy or want it to always lower it? > > -eric > > On Mar 5, 2012, at 11:56 AM, Ryan Taylor <ryta1203 at gmail.com> wrote: > > > Christoph, > > > > Yes, you are correct on the lifetime