Skip to content

Commit

Permalink
Extent gcc miscompilation workaround for replace.cu (#1461)
Browse files Browse the repository at this point in the history 
* Extent gcc miscompilation workaround for replace.cu

Fixes nvbug4526829

* Extend the general workaround to gcc-11 to be safe
  • Loading branch information
@miscco
miscco authored Feb 29, 2024
1 parent 8b6edca commit 0ebdddf
Show file tree
Hide file tree
Showing 3 changed files with 73 additions and 57 deletions.
4 changes: 2 additions & 2 deletions thrust/testing/functional.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -7,9 +7,9 @@

THRUST_DISABLE_MSVC_POSSIBLE_LOSS_OF_DATA_WARNING_BEGIN

// There is a unfortunate miscompilation of the gcc-12 vectorizer leading to OOB writes
// There is a unfortunate miscompilation of the gcc-11 vectorizer leading to OOB writes
// Adding this attribute suffices that this miscompilation does not appear anymore
#if defined(_CCCL_COMPILER_GCC) && __GNUC__ >= 12
#if defined(_CCCL_COMPILER_GCC) && __GNUC__ >= 11
#define THRUST_DISABLE_BROKEN_GCC_VECTORIZER __attribute__((optimize("no-tree-vectorize")))
#else
#define THRUST_DISABLE_BROKEN_GCC_VECTORIZER
Expand Down
122 changes: 69 additions & 53 deletions thrust/testing/replace.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -3,28 +3,35 @@
#include <thrust/iterator/discard_iterator.h>
#include <thrust/iterator/retag.h>

// There is a unfortunate miscompilation of the gcc-11 vectorizer leading to OOB writes
// Adding this attribute suffices that this miscompilation does not appear anymore
#if defined(_CCCL_COMPILER_GCC) && __GNUC__ >= 11
#define THRUST_DISABLE_BROKEN_GCC_VECTORIZER __attribute__((optimize("no-tree-vectorize")))
#else
#define THRUST_DISABLE_BROKEN_GCC_VECTORIZER
#endif

template <class Vector>
void TestReplaceSimple(void)
{
typedef typename Vector::value_type T;

Vector data(5);
data[0] = 1;
data[1] = 2;
data[0] = 1;
data[1] = 2;
data[2] = 1;
data[3] = 3;
data[4] = 2;
data[3] = 3;
data[4] = 2;

thrust::replace(data.begin(), data.end(), (T) 1, (T) 4);
thrust::replace(data.begin(), data.end(), (T) 2, (T) 5);

Vector result(5);
result[0] = 4;
result[1] = 5;
result[0] = 4;
result[1] = 5;
result[2] = 4;
result[3] = 3;
result[4] = 5;
result[3] = 3;
result[4] = 5;

ASSERT_EQUAL(data, result);
}
Expand Down Expand Up @@ -100,23 +107,23 @@ void TestReplaceCopySimple(void)
typedef typename Vector::value_type T;

Vector data(5);
data[0] = 1;
data[1] = 2;
data[0] = 1;
data[1] = 2;
data[2] = 1;
data[3] = 3;
data[4] = 2;
data[3] = 3;
data[4] = 2;

Vector dest(5);

thrust::replace_copy(data.begin(), data.end(), dest.begin(), (T) 1, (T) 4);
thrust::replace_copy(dest.begin(), dest.end(), dest.begin(), (T) 2, (T) 5);

Vector result(5);
result[0] = 4;
result[1] = 5;
result[0] = 4;
result[1] = 5;
result[2] = 4;
result[3] = 3;
result[4] = 5;
result[3] = 3;
result[4] = 5;

ASSERT_EQUAL(dest, result);
}
Expand Down Expand Up @@ -182,10 +189,10 @@ void TestReplaceCopy(const size_t n)
{
thrust::host_vector<T> h_data = unittest::random_samples<T>(n);
thrust::device_vector<T> d_data = h_data;

T old_value = 0;
T new_value = 1;

thrust::host_vector<T> h_dest(n);
thrust::device_vector<T> d_dest(n);

Expand All @@ -203,7 +210,7 @@ void TestReplaceCopyToDiscardIterator(const size_t n)
{
thrust::host_vector<T> h_data = unittest::random_samples<T>(n);
thrust::device_vector<T> d_data = h_data;

T old_value = 0;
T new_value = 1;

Expand Down Expand Up @@ -234,20 +241,20 @@ void TestReplaceIfSimple(void)
typedef typename Vector::value_type T;

Vector data(5);
data[0] = 1;
data[1] = 3;
data[0] = 1;
data[1] = 3;
data[2] = 4;
data[3] = 6;
data[4] = 5;
data[3] = 6;
data[4] = 5;

thrust::replace_if(data.begin(), data.end(), less_than_five<T>(), (T) 0);

Vector result(5);
result[0] = 0;
result[1] = 0;
result[0] = 0;
result[1] = 0;
result[2] = 0;
result[3] = 6;
result[4] = 5;
result[3] = 6;
result[4] = 5;

ASSERT_EQUAL(data, result);
}
Expand Down Expand Up @@ -303,16 +310,17 @@ DECLARE_UNITTEST(TestReplaceIfDispatchImplicit);


template <class Vector>
THRUST_DISABLE_BROKEN_GCC_VECTORIZER
void TestReplaceIfStencilSimple(void)
{
typedef typename Vector::value_type T;

Vector data(5);
data[0] = 1;
data[1] = 3;
data[0] = 1;
data[1] = 3;
data[2] = 4;
data[3] = 6;
data[4] = 5;
data[3] = 6;
data[4] = 5;

Vector stencil(5);
stencil[0] = 5;
Expand All @@ -324,11 +332,11 @@ void TestReplaceIfStencilSimple(void)
thrust::replace_if(data.begin(), data.end(), stencil.begin(), less_than_five<T>(), (T) 0);

Vector result(5);
result[0] = 1;
result[1] = 0;
result[0] = 1;
result[1] = 0;
result[2] = 4;
result[3] = 0;
result[4] = 5;
result[3] = 0;
result[4] = 5;

ASSERT_EQUAL(data, result);
}
Expand Down Expand Up @@ -388,6 +396,7 @@ DECLARE_UNITTEST(TestReplaceIfStencilDispatchImplicit);


template <typename T>
THRUST_DISABLE_BROKEN_GCC_VECTORIZER
void TestReplaceIf(const size_t n)
{
thrust::host_vector<T> h_data = unittest::random_samples<T>(n);
Expand All @@ -402,6 +411,7 @@ DECLARE_VARIABLE_UNITTEST(TestReplaceIf);


template <typename T>
THRUST_DISABLE_BROKEN_GCC_VECTORIZER
void TestReplaceIfStencil(const size_t n)
{
thrust::host_vector<T> h_data = unittest::random_samples<T>(n);
Expand All @@ -419,27 +429,28 @@ DECLARE_VARIABLE_UNITTEST(TestReplaceIfStencil);


template <class Vector>
THRUST_DISABLE_BROKEN_GCC_VECTORIZER
void TestReplaceCopyIfSimple(void)
{
typedef typename Vector::value_type T;

Vector data(5);
data[0] = 1;
data[1] = 3;
data[0] = 1;
data[1] = 3;
data[2] = 4;
data[3] = 6;
data[4] = 5;
data[3] = 6;
data[4] = 5;

Vector dest(5);

thrust::replace_copy_if(data.begin(), data.end(), dest.begin(), less_than_five<T>(), (T) 0);

Vector result(5);
result[0] = 0;
result[1] = 0;
result[0] = 0;
result[1] = 0;
result[2] = 0;
result[3] = 6;
result[4] = 5;
result[3] = 6;
result[4] = 5;

ASSERT_EQUAL(dest, result);
}
Expand Down Expand Up @@ -501,16 +512,17 @@ DECLARE_UNITTEST(TestReplaceCopyIfDispatchImplicit);


template <class Vector>
THRUST_DISABLE_BROKEN_GCC_VECTORIZER
void TestReplaceCopyIfStencilSimple(void)
{
typedef typename Vector::value_type T;

Vector data(5);
data[0] = 1;
data[1] = 3;
data[0] = 1;
data[1] = 3;
data[2] = 4;
data[3] = 6;
data[4] = 5;
data[3] = 6;
data[4] = 5;

Vector stencil(5);
stencil[0] = 1;
Expand All @@ -524,11 +536,11 @@ void TestReplaceCopyIfStencilSimple(void)
thrust::replace_copy_if(data.begin(), data.end(), stencil.begin(), dest.begin(), less_than_five<T>(), (T) 0);

Vector result(5);
result[0] = 0;
result[1] = 3;
result[0] = 0;
result[1] = 3;
result[2] = 0;
result[3] = 6;
result[4] = 5;
result[3] = 6;
result[4] = 5;

ASSERT_EQUAL(dest, result);
}
Expand Down Expand Up @@ -595,6 +607,7 @@ DECLARE_UNITTEST(TestReplaceCopyIfStencilDispatchImplicit);


template <typename T>
THRUST_DISABLE_BROKEN_GCC_VECTORIZER
void TestReplaceCopyIf(const size_t n)
{
thrust::host_vector<T> h_data = unittest::random_samples<T>(n);
Expand All @@ -613,6 +626,7 @@ DECLARE_VARIABLE_UNITTEST(TestReplaceCopyIf);


template <typename T>
THRUST_DISABLE_BROKEN_GCC_VECTORIZER
void TestReplaceCopyIfToDiscardIterator(const size_t n)
{
thrust::host_vector<T> h_data = unittest::random_samples<T>(n);
Expand All @@ -632,6 +646,7 @@ void TestReplaceCopyIfToDiscardIterator(const size_t n)
DECLARE_VARIABLE_UNITTEST(TestReplaceCopyIfToDiscardIterator);

template <typename T>
THRUST_DISABLE_BROKEN_GCC_VECTORIZER
void TestReplaceCopyIfStencil(const size_t n)
{
thrust::host_vector<T> h_data = unittest::random_samples<T>(n);
Expand All @@ -652,6 +667,7 @@ void TestReplaceCopyIfStencil(const size_t n)
DECLARE_VARIABLE_UNITTEST(TestReplaceCopyIfStencil);

template <typename T>
THRUST_DISABLE_BROKEN_GCC_VECTORIZER
void TestReplaceCopyIfStencilToDiscardIterator(const size_t n)
{
thrust::host_vector<T> h_data = unittest::random_samples<T>(n);
Expand Down
4 changes: 2 additions & 2 deletions thrust/testing/transform.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -7,9 +7,9 @@
#include <thrust/iterator/zip_iterator.h>
#include <thrust/iterator/retag.h>

// There is a unfortunate miscompilation of the gcc-12 vectorizer leading to OOB writes
// There is a unfortunate miscompilation of the gcc-11 vectorizer leading to OOB writes
// Adding this attribute suffices that this miscompilation does not appear anymore
#if defined(_CCCL_COMPILER_GCC) && __GNUC__ >= 12
#if defined(_CCCL_COMPILER_GCC) && __GNUC__ >= 11
#define THRUST_DISABLE_BROKEN_GCC_VECTORIZER __attribute__((optimize("no-tree-vectorize")))
#else
#define THRUST_DISABLE_BROKEN_GCC_VECTORIZER
Expand Down

0 comments on commit 0ebdddf

Please sign in to comment.