From 36bf811574a0787910d80d12fd9ea186481d7939 Mon Sep 17 00:00:00 2001 From: Tea Date: Wed, 2 Dec 2015 15:39:19 +0800 Subject: [PATCH] Remove hamming_distance and popcount --- include/caffe/util/math_functions.hpp | 7 ---- src/caffe/test/test_math_functions.cpp | 41 ----------------------- src/caffe/util/math_functions.cpp | 22 ------------- src/caffe/util/math_functions.cu | 45 -------------------------- 4 files changed, 115 deletions(-) diff --git a/include/caffe/util/math_functions.hpp b/include/caffe/util/math_functions.hpp index 2cacd8e72cd..6f6d3feeae2 100644 --- a/include/caffe/util/math_functions.hpp +++ b/include/caffe/util/math_functions.hpp @@ -101,9 +101,6 @@ template Dtype caffe_cpu_strided_dot(const int n, const Dtype* x, const int incx, const Dtype* y, const int incy); -template -int caffe_cpu_hamming_distance(const int n, const Dtype* x, const Dtype* y); - // Returns the sum of the absolute values of the elements of vector x template Dtype caffe_cpu_asum(const int n, const Dtype* x); @@ -234,10 +231,6 @@ void caffe_gpu_rng_bernoulli(const int n, const Dtype p, int* r); template void caffe_gpu_dot(const int n, const Dtype* x, const Dtype* y, Dtype* out); -template -uint32_t caffe_gpu_hamming_distance(const int n, const Dtype* x, - const Dtype* y); - template void caffe_gpu_asum(const int n, const Dtype* x, Dtype* y); diff --git a/src/caffe/test/test_math_functions.cpp b/src/caffe/test/test_math_functions.cpp index fbee3f9c32d..efc5a2784eb 100644 --- a/src/caffe/test/test_math_functions.cpp +++ b/src/caffe/test/test_math_functions.cpp @@ -39,27 +39,6 @@ class MathFunctionsTest : public MultiDeviceTest { delete blob_top_; } - // http://en.wikipedia.org/wiki/Hamming_distance - int ReferenceHammingDistance(const int n, const Dtype* x, const Dtype* y) { - int dist = 0; - uint64_t val; - for (int i = 0; i < n; ++i) { - if (sizeof(Dtype) == 8) { - val = static_cast(x[i]) ^ static_cast(y[i]); - } else if (sizeof(Dtype) == 4) { - val = static_cast(x[i]) ^ static_cast(y[i]); - } else { - LOG(FATAL) << "Unrecognized Dtype size: " << sizeof(Dtype); - } - // Count the number of set bits - while (val) { - ++dist; - val &= val - 1; - } - } - return dist; - } - Blob* const blob_bottom_; Blob* const blob_top_; }; @@ -76,14 +55,6 @@ TYPED_TEST(CPUMathFunctionsTest, TestNothing) { // due to the set up overhead. } -TYPED_TEST(CPUMathFunctionsTest, TestHammingDistance) { - int n = this->blob_bottom_->count(); - const TypeParam* x = this->blob_bottom_->cpu_data(); - const TypeParam* y = this->blob_top_->cpu_data(); - EXPECT_EQ(this->ReferenceHammingDistance(n, x, y), - caffe_cpu_hamming_distance(n, x, y)); -} - TYPED_TEST(CPUMathFunctionsTest, TestAsum) { int n = this->blob_bottom_->count(); const TypeParam* x = this->blob_bottom_->cpu_data(); @@ -156,18 +127,6 @@ class GPUMathFunctionsTest : public MathFunctionsTest > { TYPED_TEST_CASE(GPUMathFunctionsTest, TestDtypes); -// TODO: Fix caffe_gpu_hamming_distance and re-enable this test. -TYPED_TEST(GPUMathFunctionsTest, DISABLED_TestHammingDistance) { - int n = this->blob_bottom_->count(); - const TypeParam* x = this->blob_bottom_->cpu_data(); - const TypeParam* y = this->blob_top_->cpu_data(); - int reference_distance = this->ReferenceHammingDistance(n, x, y); - x = this->blob_bottom_->gpu_data(); - y = this->blob_top_->gpu_data(); - int computed_distance = caffe_gpu_hamming_distance(n, x, y); - EXPECT_EQ(reference_distance, computed_distance); -} - TYPED_TEST(GPUMathFunctionsTest, TestAsum) { int n = this->blob_bottom_->count(); const TypeParam* x = this->blob_bottom_->cpu_data(); diff --git a/src/caffe/util/math_functions.cpp b/src/caffe/util/math_functions.cpp index 0aab6b17b85..71c02274a75 100644 --- a/src/caffe/util/math_functions.cpp +++ b/src/caffe/util/math_functions.cpp @@ -348,28 +348,6 @@ float caffe_cpu_dot(const int n, const float* x, const float* y); template double caffe_cpu_dot(const int n, const double* x, const double* y); -template <> -int caffe_cpu_hamming_distance(const int n, const float* x, - const float* y) { - int dist = 0; - for (int i = 0; i < n; ++i) { - dist += __builtin_popcount(static_cast(x[i]) ^ - static_cast(y[i])); - } - return dist; -} - -template <> -int caffe_cpu_hamming_distance(const int n, const double* x, - const double* y) { - int dist = 0; - for (int i = 0; i < n; ++i) { - dist += __builtin_popcountl(static_cast(x[i]) ^ - static_cast(y[i])); - } - return dist; -} - template <> float caffe_cpu_asum(const int n, const float* x) { return cblas_sasum(n, x, 1); diff --git a/src/caffe/util/math_functions.cu b/src/caffe/util/math_functions.cu index e4d0c4b04bb..4c587537435 100644 --- a/src/caffe/util/math_functions.cu +++ b/src/caffe/util/math_functions.cu @@ -371,51 +371,6 @@ DEFINE_AND_INSTANTIATE_GPU_UNARY_FUNC(sign, y[index] = (Dtype(0) < x[index]) - (x[index] < Dtype(0))); DEFINE_AND_INSTANTIATE_GPU_UNARY_FUNC(sgnbit, y[index] = signbit(x[index])); -__global__ void popc_kernel(const int n, const float* a, - const float* b, uint8_t* y) { - CUDA_KERNEL_LOOP(index, n) { - y[index] = __popc(static_cast(a[index]) ^ - static_cast(b[index])); - } -} - -__global__ void popcll_kernel(const int n, const double* a, - const double* b, uint8_t* y) { - CUDA_KERNEL_LOOP(index, n) { - y[index] = __popcll(static_cast(a[index]) ^ - static_cast(b[index])); - } -} - -template <> -uint32_t caffe_gpu_hamming_distance(const int n, const float* x, - const float* y) { - // TODO: Fix caffe_gpu_hamming_distance (see failing unit test - // TestHammingDistanceGPU in test_math_functions.cpp). - NOT_IMPLEMENTED; - thrust::device_vector popcounts(n); - // NOLINT_NEXT_LINE(whitespace/operators) - popc_kernel<<>>( - n, x, y, thrust::raw_pointer_cast(popcounts.data())); - return thrust::reduce(popcounts.begin(), popcounts.end(), - (uint32_t) 0, thrust::plus()); -} - -template <> -uint32_t caffe_gpu_hamming_distance(const int n, const double* x, - const double* y) { - // TODO: Fix caffe_gpu_hamming_distance (see failing unit test - // TestHammingDistanceGPU in test_math_functions.cpp). - NOT_IMPLEMENTED; - thrust::device_vector popcounts(n); - // NOLINT_NEXT_LINE(whitespace/operators) - popcll_kernel<<>>( - n, x, y, thrust::raw_pointer_cast(popcounts.data())); - return thrust::reduce(popcounts.begin(), popcounts.end(), - /* NOLINT_NEXT_LINE(build/include_what_you_use) */ - (uint32_t) 0, thrust::plus()); -} - void caffe_gpu_rng_uniform(const int n, unsigned int* r) { CURAND_CHECK(curandGenerate(Caffe::curand_generator(), r, n)); }