cudpp_segmented_scan_cta.cu "Removed dead synchronization intrinsic" advisories

[GoogleCodeExporter]

What steps will reproduce the problem?
1. Build cudpp release or debug

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

Expect no errors or warnings or advisories.  Instead get lots of these:

jS4_PjS9_
src/cta/segmented_scan_cta.cu(868): Advisory: Removed dead synchronization
intrinsic from function
_Z14segmentedScan4If19SegmentedScanTraitsIfL13CUDPPOperator2ELb0ELb0ELb0ELb0ELb1
ELb0EEEvPT_PKS3_PKjjS4_PjS9_

Suggested fix:
I realize that removing this __syncthreads() causes failure.  I believe
though that the compiler is only removing it from some calls to the
function that includes it, not all.  So instead of putting the
syncthreads() inside this function, put it right before the call to the
function, only where it is needed.

Please use labels and text to provide additional information.

Original issue reported on code.google.com by harr...@gmail.com on 17 Jun 2009 at 1:38

[GoogleCodeExporter]

What do you know - CUDA 2.2 seems to have fixed this bug - so I will comment 
the offending syncthreads for 
now (only tested on linux though - I will test on windows as well)

Original comment by shu...@gmail.com on 25 Jun 2009 at 5:18

[GoogleCodeExporter]

Verified on Linux with CUDA 2.2

Original comment by shu...@gmail.com on 25 Jun 2009 at 5:22

• Changed state: Fixed

[GoogleCodeExporter]

Verified on Windows XP with CUDA 2.2

Original comment by shu...@gmail.com on 25 Jun 2009 at 3:21

[GoogleCodeExporter]

Reverted since behavior much more stable with the syncthreads

Original comment by shu...@gmail.com on 25 Jun 2009 at 7:00

• Changed state: WontFix

[GoogleCodeExporter]

FIne, I'll fix it.  I explained how to fix it in the description, I'll try it 
myself.  But we aren't marking this "wont fix".

Original comment by harr...@gmail.com on 25 Jun 2009 at 9:38

• Changed state: Started

[GoogleCodeExporter]

I tried your approach and it didn't work (that is putting the syncthreads 
before calls to warpSegScan)

Original comment by shu...@gmail.com on 25 Jun 2009 at 9:41

[GoogleCodeExporter]

This advisory is fixed in CUDA 2.3.  There is a bug in the compiler where 
moving the
__syncthreads() from the beginning of warpSegScan() to right before each call 
of it
causes the code to fail.  I filed an NVIDIA bug for this.

Closing this for now.

Original comment by harr...@gmail.com on 26 Jun 2009 at 8:03

• Changed state: WontFix

[GoogleCodeExporter]

Original comment by harr...@gmail.com on 26 Jun 2009 at 8:04

• Changed state: Verified

[GoogleCodeExporter]

I fixed this for CUDA 2.2.

Original comment by harr...@gmail.com on 30 Jun 2009 at 5:07