device_vector<float>::assign(size_type n, const T &x) is broken for float x != 0

[GoogleCodeExporter]

#include <thrust/device_vector.h>
#include <thrust/host_vector.h>

#include <cassert>

int main(int argc, char* argv[])
{
    thrust::device_vector<float>    device_some_vector  (10, 0.f);              //  [0, 0, 0, 0, 0, 0, 0, 0, 0, 0]
    thrust::host_vector<float>      host_some_vector    (10, 0.f);              //  [0, 0, 0, 0, 0, 0, 0, 0, 0, 0]
    thrust::host_vector<float>      test_some_vector    (device_some_vector);   //  [0, 0, 0, 0, 0, 0, 0, 0, 0, 0]

    assert(host_some_vector == test_some_vector);   //  Should succeed

    device_some_vector  .assign(10, 5.0f);  //  [5, 5, 5, 5, 5, 5, 5, 5, 5, 5]
    host_some_vector    .assign(10, 5.0f);  //  [5, 5, 5, 5, 5, 5, 5, 5, 5, 5]
    test_some_vector = device_some_vector;  //  [5, 5, 5, 5, 5, 5, 5, 5, 5, 5] expected, but instead we get [0, 0, 0, 0, 0, 0, 0, 0, 0, 0](?!)

    assert(host_some_vector == test_some_vector);   //  Should succeed but fails, 

    return 0;
}

Expected output: in the debugger I expect the second assert to succeed, however 
it fails, the container is filled with zeroes instead of the value 5.0f

Thrust version: 1.30.0

nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2010 NVIDIA Corporation
Built on Sat_Nov__6_00:59:51_PDT_2010
Cuda compilation tools, release 3.2, V0.2.1221

Host compiler: (vc9 32-bit, service pack 1)

Microsoft (R) 32-bit C/C++ Optimizing Compiler Version 15.00.30729.01 for 80x86
Copyright (C) Microsoft Corporation.  All rights reserved.

Operating system:

Windows 7 Ultimate x64, version 6.1 build 7600

GPU:

GeForce 8800M GTX

Original issue reported on code.google.com by george.v...@organicvectory.com on 19 Apr 2011 at 8:42

Attachments:

• assign_bug.cpp

[GoogleCodeExporter]

I think I've tracked down the issue to this line:

thrust/detail/device/cuda/fill.inl(196):

  if(std::strncmp(reinterpret_cast<const char*>(&value), reinterpret_cast<const char*>(&zero), sizeof(T)) == 0)

I think the strncmp statement stops at the first zero-char encountered and does 
not check all the bytes in a float. Perhaps using memcmp() would be the correct 
thing to do here.

Original comment by george.v...@organicvectory.com on 19 Apr 2011 at 9:02

[GoogleCodeExporter]

Hi George, thanks for reporting and tracking down the bug.  I think it's been 
fixed [1] since version 1.3.

Could you confirm that your code works correctly with version 1.4?  You can get 
it with Mercurial:

$ hg clone http://thrust.googlecode.com/hg -r 1.4.0 thrust-1.4

[1] 
http://code.google.com/p/thrust/source/detail?r=a2e913c7ed831c9cf155c4cac43397eb
dd131ce6&path=/thrust/detail/backend/cuda/fill.inl

Original comment by jaredhoberock on 19 Apr 2011 at 11:40

[GoogleCodeExporter]

I have tested the code with Cuda 4.0 and thrust 1.4.0. The only change I had to 
make was to compile with nvcc as the assign now leads to a kernel launch, 
whereas the original version could run without requiring a kernel launch.

So the fix is confirmed, thanks for your quick response.

Original comment by george.v...@organicvectory.com on 21 Apr 2011 at 9:40

[GoogleCodeExporter]

Original comment by jaredhoberock on 21 Apr 2011 at 6:24

• Changed state: Invalid