memoryChecker reports bad global memory access

[GoogleCodeExporter]

What steps will reproduce the problem?
1. allocate device memory
2. attempt to access the allocated memory within a kernel

What is the expected output? What do you see instead?

I expect the program to succeed with no output. Instead, the memory checker 
claims one of my kernels is accessing memory that is not allocated or mapped. 
However, it clearly is allocated, as its own list of device allocations shows 
it is within the fifth allocation:

terminate called after throwing an instance of 'hydrazine::Exception'
  what():  [PC 26] [thread 0] [cta 0] ld.global.u8 %r24, [%r23 + 0] - Global memory access 0xb79d0f is not within any allocated or mapped range.

Nearby Device Allocations
[0xa66fc0] - [0xa67020] (96 bytes)
[0xa670a0] - [0xa674d8] (1080 bytes)
[0xa67620] - [0xa67680] (96 bytes)
[0xa67700] - [0xa67b38] (1080 bytes)
[0xb79d00] - [0xb7bd00] (8192 bytes)
[0xb7be20] - [0xb7de20] (8192 bytes)
[0xb7df40] - [0xb7e740] (2048 bytes)
[0xb7e860] - [0xb7f060] (2048 bytes)
[0xb7f180] - [0xb7f980] (2048 bytes)
[0xb7faa0] - [0xb802a0] (2048 bytes)
[0xb803c0] - [0xb80bc0] (2048 bytes)
[0xb80ce0] - [0xb814e0] (2048 bytes)
[0xb83520] - [0xb85520] (8192 bytes)
[0xb87660] - [0xb89660] (8192 bytes)
[0xb89780] - [0xb8e780] (20480 bytes)
[0xb8e8a0] - [0xb938a0] (20480 bytes)
[0xb939c0] - [0xb989c0] (20480 bytes)
[0xb98ae0] - [0xb9dae0] (20480 bytes)
[0xb9dc00] - [0xba2c00] (20480 bytes)
[0xba2d20] - [0xba7d20] (20480 bytes)
[0xba7e40] - [0xbace40] (20480 bytes)
[0xbacf60] - [0xbb1f60] (20480 bytes)
[0xbb2080] - [0xbb2880] (2048 bytes)
[0xbb29a0] - [0xbb31a0] (2048 bytes)
[0xbb32c0] - [0xbb3ac0] (2048 bytes)
[0xbb3be0] - [0xbb43e0] (2048 bytes)

What version of the product are you using? On what operating system?

-Ocelot version: SVN r634
-compiled with gcc 4.5 with the macro to disable c++0x feature in 
KernelEntry.cpp as described in 
http://groups.google.com/group/gpuocelot/msg/55339e218bc5bdaa?pli=1
-modifications to type checker described in 
http://groups.google.com/group/gpuocelot/browse_thread/thread/186dcb0bee10ed8b

Cuda version: 2.3

Please provide any additional information below.

Original issue reported on code.google.com by meaner...@gmail.com on 4 Aug 2010 at 4:21

[GoogleCodeExporter]

Does your kernel happen to use memory that is shared between the host and 
device.  There may have been a bug that crept in recently related to this.

Original comment by gregory....@gatech.edu on 26 Aug 2010 at 12:36

[GoogleCodeExporter]

Yes, it does use zero-copy memory. Is there any way I can confirm that is the 
issue?

Original comment by meaner...@gmail.com on 26 Aug 2010 at 4:51

[GoogleCodeExporter]

I added a potential fix in r664 that allowed several zero copy examples to pass 
again on my machine.  Let me know if this fixes your problem.

Original comment by gregory....@gatech.edu on 27 Aug 2010 at 2:53

• Changed state: Accepted

[GoogleCodeExporter]

Original comment by gregory....@gatech.edu on 18 Oct 2010 at 6:47

• Changed state: Fixed