microsoft
/
caffe
зеркало из https://github.com/microsoft/caffe.git
1
0
Форкнуть 0
Код Задачи Пакеты Проекты Релизы Вики Активность

Remove hamming_distance and popcount

This commit is contained in:
Tea 2015-12-02 15:39:19 +08:00
Родитель 196b995efe
Коммит 36bf811574
4 изменённых файлов: 0 добавлений и 115 удалений
Показать статистику Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы

7
include/caffe/util/math_functions.hpp
Просмотреть файл

@ -101,9 +101,6 @@ template <typename Dtype>
Dtype caffe_cpu_strided_dot(const int n, const Dtype* x, const int incx,
const Dtype* y, const int incy);
template <typename Dtype>
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 <typename Dtype>
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 <typename Dtype>
void caffe_gpu_dot(const int n, const Dtype* x, const Dtype* y, Dtype* out);
template <typename Dtype>
uint32_t caffe_gpu_hamming_distance(const int n, const Dtype* x,
const Dtype* y);
template <typename Dtype>
void caffe_gpu_asum(const int n, const Dtype* x, Dtype* y);

41
src/caffe/test/test_math_functions.cpp
Просмотреть файл

@ -39,27 +39,6 @@ class MathFunctionsTest : public MultiDeviceTest<TypeParam> {
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<uint64_t>(x[i]) ^ static_cast<uint64_t>(y[i]);
} else if (sizeof(Dtype) == 4) {
val = static_cast<uint32_t>(x[i]) ^ static_cast<uint32_t>(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<Dtype>* const blob_bottom_;
Blob<Dtype>* 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<TypeParam>(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<GPUDevice<Dtype> > {
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<TypeParam>(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();

22
src/caffe/util/math_functions.cpp
Просмотреть файл

@ -348,28 +348,6 @@ float caffe_cpu_dot<float>(const int n, const float* x, const float* y);
template
double caffe_cpu_dot<double>(const int n, const double* x, const double* y);
template <>
int caffe_cpu_hamming_distance<float>(const int n, const float* x,
const float* y) {
int dist = 0;
for (int i = 0; i < n; ++i) {
dist += __builtin_popcount(static_cast<uint32_t>(x[i]) ^
static_cast<uint32_t>(y[i]));
}
return dist;
}
template <>
int caffe_cpu_hamming_distance<double>(const int n, const double* x,
const double* y) {
int dist = 0;
for (int i = 0; i < n; ++i) {
dist += __builtin_popcountl(static_cast<uint64_t>(x[i]) ^
static_cast<uint64_t>(y[i]));
}
return dist;
}
template <>
float caffe_cpu_asum<float>(const int n, const float* x) {
return cblas_sasum(n, x, 1);

45
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<uint32_t>(a[index]) ^
static_cast<uint32_t>(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<uint64_t>(a[index]) ^
static_cast<uint64_t>(b[index]));
}
}
template <>
uint32_t caffe_gpu_hamming_distance<float>(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<uint8_t> popcounts(n);
// NOLINT_NEXT_LINE(whitespace/operators)
popc_kernel<<<CAFFE_GET_BLOCKS(n), CAFFE_CUDA_NUM_THREADS>>>(
n, x, y, thrust::raw_pointer_cast(popcounts.data()));
return thrust::reduce(popcounts.begin(), popcounts.end(),
(uint32_t) 0, thrust::plus<uint32_t>());
}
template <>
uint32_t caffe_gpu_hamming_distance<double>(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<uint8_t> popcounts(n);
// NOLINT_NEXT_LINE(whitespace/operators)
popcll_kernel<<<CAFFE_GET_BLOCKS(n), CAFFE_CUDA_NUM_THREADS>>>(
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<uint32_t>());
}
void caffe_gpu_rng_uniform(const int n, unsigned int* r) {
CURAND_CHECK(curandGenerate(Caffe::curand_generator(), r, n));
}