fix address calculation for ldg

[vchuravy]

Should fix #548

The issue was that we forwarded the address calculation to unsafe_cached_load which should operate in bytes.

[vchuravy]

Before:

	ld.param.u64 	%rd2, [%rd3+8];
	ld.local.u64 	%rd9, [%rd1];
	cvta.to.global.u64 	%rd10, %rd9;
	cvt.u64.u32 	%rd11, %r1;
	add.s64 	%rd12, %rd10, %rd11;
	add.s64 	%rd13, %rd12, 1;
	cvta.global.u64 	%rd14, %rd13;
	ld.global.nc.f64 	%fd1, [%rd14+-1];
	mul.wide.u32 	%rd15, %r1, 8;
	add.s64 	%rd16, %rd15, %rd2;
	st.global.f64 	[%rd16], %fd1;

After:

	ld.param.u64 	%rd2, [%rd3+8];
	shl.b32 	%r3, %r2, 3;
	cvt.u64.u32 	%rd9, %r3;
	ld.local.u64 	%rd10, [%rd1];
	add.s64 	%rd11, %rd10, %rd9;
	ld.global.nc.f64 	%fd1, [%rd11+-8];
	mul.wide.u32 	%rd12, %r1, 8;
	add.s64 	%rd13, %rd12, %rd2;
	st.global.f64 	[%rd13], %fd1;
```

[codecov]

Codecov Report

Merging #549 (f381564) into master (1139ffd) will not change coverage.
The diff coverage is n/a.

@@           Coverage Diff           @@
##           master     #549   +/-   ##
=======================================
  Coverage   77.95%   77.95%           
=======================================
  Files         116      116           
  Lines        6422     6422           
=======================================
  Hits         5006     5006           
  Misses       1416     1416           

---

Continue to review full report at Codecov.

Legend - Click here to learn more
Δ = absolute <relative> (impact), ø = not affected, ? = missing data
Powered by Codecov. Last update 8a92e95...f381564. Read the comment docs.

[maleadt]

Nice, thanks! I had a feeling Const was broken but thought it was due to bounds checking missing something.