From 27690e3b6e442d02f78ecc2b3473ec8ce244f16e Mon Sep 17 00:00:00 2001 From: Alexey Spizhevoy Date: Wed, 24 Nov 2010 11:40:14 +0000 Subject: [PATCH] added minMaxLoc function into gpu module --- modules/gpu/include/opencv2/gpu/gpu.hpp | 5 +- modules/gpu/src/arithm.cpp | 52 +++++++ modules/gpu/src/cuda/mathfunc.cu | 248 +++++++++++++++++++++++++++++--- tests/gpu/src/arithm.cpp | 66 +++++++++ 4 files changed, 354 insertions(+), 17 deletions(-) diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index 70f9eda..4dcbd0b 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -422,7 +422,10 @@ namespace cv CV_EXPORTS Scalar sum(const GpuMat& m); //! finds global minimum and maximum array elements and returns their values - CV_EXPORTS void minMax(const GpuMat& src, double* minVal, double* maxVal = 0); + CV_EXPORTS void minMax(const GpuMat& src, double* minVal, double* maxVal=0); + + //! finds global minimum and maximum array elements and returns their values with locations + CV_EXPORTS void minMaxLoc(const GpuMat& src, double* minVal, double* maxVal=0, Point* minLoc=0, Point* maxLoc=0); //! transforms 8-bit unsigned integers using lookup table: dst(i)=lut(src(i)) //! destination array will have the depth type as lut and the same channels number as source diff --git a/modules/gpu/src/arithm.cpp b/modules/gpu/src/arithm.cpp index 7c54719..d74fbb2 100644 --- a/modules/gpu/src/arithm.cpp +++ b/modules/gpu/src/arithm.cpp @@ -66,6 +66,7 @@ double cv::gpu::norm(const GpuMat&, const GpuMat&, int) { throw_nogpu(); return void cv::gpu::flip(const GpuMat&, GpuMat&, int) { throw_nogpu(); } Scalar cv::gpu::sum(const GpuMat&) { throw_nogpu(); return Scalar(); } void cv::gpu::minMax(const GpuMat&, double*, double*) { throw_nogpu(); } +void cv::gpu::minMaxLoc(const GpuMat&, double*, double*, Point*, Point*) { throw_nogpu(); } void cv::gpu::LUT(const GpuMat&, const Mat&, GpuMat&) { throw_nogpu(); } void cv::gpu::exp(const GpuMat&, GpuMat&) { throw_nogpu(); } void cv::gpu::log(const GpuMat&, GpuMat&) { throw_nogpu(); } @@ -530,6 +531,57 @@ void cv::gpu::minMax(const GpuMat& src, double* minVal, double* maxVal) } } + +//////////////////////////////////////////////////////////////////////// +// minMaxLoc + +namespace cv { namespace gpu { namespace mathfunc { + template + void min_max_loc_caller(const DevMem2D src, double* minval, double* maxval, int* minlocx, int* minlocy, + int* maxlocx, int* maxlocy); +}}} + +void cv::gpu::minMaxLoc(const GpuMat& src, double* minVal, double* maxVal, Point* minLoc, Point* maxLoc) +{ + CV_Assert(src.channels() == 1); + + double maxVal_; + if (!maxVal) maxVal = &maxVal_; + + cv::Point minLoc_; + if (!minLoc) minLoc = &minLoc_; + + cv::Point maxLoc_; + if (!maxLoc) maxLoc = &maxLoc_; + + switch (src.type()) + { + case CV_8U: + mathfunc::min_max_loc_caller(src, minVal, maxVal, &minLoc->x, &minLoc->y, &maxLoc->x, &maxLoc->y); + break; + case CV_8S: + mathfunc::min_max_loc_caller(src, minVal, maxVal, &minLoc->x, &minLoc->y, &maxLoc->x, &maxLoc->y); + break; + case CV_16U: + mathfunc::min_max_loc_caller(src, minVal, maxVal, &minLoc->x, &minLoc->y, &maxLoc->x, &maxLoc->y); + break; + case CV_16S: + mathfunc::min_max_loc_caller(src, minVal, maxVal, &minLoc->x, &minLoc->y, &maxLoc->x, &maxLoc->y); + break; + case CV_32S: + mathfunc::min_max_loc_caller(src, minVal, maxVal, &minLoc->x, &minLoc->y, &maxLoc->x, &maxLoc->y); + break; + case CV_32F: + mathfunc::min_max_loc_caller(src, minVal, maxVal, &minLoc->x, &minLoc->y, &maxLoc->x, &maxLoc->y); + break; + case CV_64F: + mathfunc::min_max_loc_caller(src, minVal, maxVal, &minLoc->x, &minLoc->y, &maxLoc->x, &maxLoc->y); + break; + default: + CV_Error(CV_StsBadArg, "Unsupported type"); + } +} + //////////////////////////////////////////////////////////////////////// // LUT diff --git a/modules/gpu/src/cuda/mathfunc.cu b/modules/gpu/src/cuda/mathfunc.cu index d69a32a..bff3a30 100644 --- a/modules/gpu/src/cuda/mathfunc.cu +++ b/modules/gpu/src/cuda/mathfunc.cu @@ -410,10 +410,10 @@ namespace cv { namespace gpu { namespace mathfunc template <> struct MinMaxTypeTraits { typedef float best_type; }; template <> struct MinMaxTypeTraits { typedef double best_type; }; - template struct Cmp {}; + template struct Opt {}; template - struct Cmp + struct Opt { static __device__ void call(unsigned int tid, unsigned int offset, volatile T* optval) { @@ -422,7 +422,7 @@ namespace cv { namespace gpu { namespace mathfunc }; template - struct Cmp + struct Opt { static __device__ void call(unsigned int tid, unsigned int offset, volatile T* optval) { @@ -448,23 +448,22 @@ namespace cv { namespace gpu { namespace mathfunc __syncthreads(); - if (nthreads >= 512) if (tid < 256) { Cmp::call(tid, 256, soptval); __syncthreads(); } - if (nthreads >= 256) if (tid < 128) { Cmp::call(tid, 128, soptval); __syncthreads(); } - if (nthreads >= 128) if (tid < 64) { Cmp::call(tid, 64, soptval); __syncthreads(); } + if (nthreads >= 512) if (tid < 256) { Opt::call(tid, 256, soptval); __syncthreads(); } + if (nthreads >= 256) if (tid < 128) { Opt::call(tid, 128, soptval); __syncthreads(); } + if (nthreads >= 128) if (tid < 64) { Opt::call(tid, 64, soptval); __syncthreads(); } if (tid < 32) { - if (nthreads >= 64) Cmp::call(tid, 32, soptval); - if (nthreads >= 32) Cmp::call(tid, 16, soptval); - if (nthreads >= 16) Cmp::call(tid, 8, soptval); - if (nthreads >= 8) Cmp::call(tid, 4, soptval); - if (nthreads >= 4) Cmp::call(tid, 2, soptval); - if (nthreads >= 2) Cmp::call(tid, 1, soptval); + if (nthreads >= 64) Opt::call(tid, 32, soptval); + if (nthreads >= 32) Opt::call(tid, 16, soptval); + if (nthreads >= 16) Opt::call(tid, 8, soptval); + if (nthreads >= 8) Opt::call(tid, 4, soptval); + if (nthreads >= 4) Opt::call(tid, 2, soptval); + if (nthreads >= 2) Opt::call(tid, 1, soptval); } if (tid == 0) ((T*)optval.ptr(blockIdx.y))[blockIdx.x] = (T)soptval[0]; } - template void min_max_caller(const DevMem2D src, double* minval, double* maxval) @@ -472,17 +471,19 @@ namespace cv { namespace gpu { namespace mathfunc dim3 threads(32, 8); // Allocate memory for aux. buffers - DevMem2D minval_buf[2]; DevMem2D maxval_buf[2]; + DevMem2D minval_buf[2]; minval_buf[0].cols = divUp(src.cols, threads.x); minval_buf[0].rows = divUp(src.rows, threads.y); minval_buf[1].cols = divUp(minval_buf[0].cols, threads.x); minval_buf[1].rows = divUp(minval_buf[0].rows, threads.y); + cudaSafeCall(cudaMallocPitch(&minval_buf[0].data, &minval_buf[0].step, minval_buf[0].cols * sizeof(T), minval_buf[0].rows)); + cudaSafeCall(cudaMallocPitch(&minval_buf[1].data, &minval_buf[1].step, minval_buf[1].cols * sizeof(T), minval_buf[1].rows)); + + DevMem2D maxval_buf[2]; maxval_buf[0].cols = divUp(src.cols, threads.x); maxval_buf[0].rows = divUp(src.rows, threads.y); maxval_buf[1].cols = divUp(maxval_buf[0].cols, threads.x); maxval_buf[1].rows = divUp(maxval_buf[0].rows, threads.y); - cudaSafeCall(cudaMallocPitch(&minval_buf[0].data, &minval_buf[0].step, minval_buf[0].cols * sizeof(T), minval_buf[0].rows)); - cudaSafeCall(cudaMallocPitch(&minval_buf[1].data, &minval_buf[1].step, minval_buf[1].cols * sizeof(T), minval_buf[1].rows)); cudaSafeCall(cudaMallocPitch(&maxval_buf[0].data, &maxval_buf[0].step, maxval_buf[0].cols * sizeof(T), maxval_buf[0].rows)); cudaSafeCall(cudaMallocPitch(&maxval_buf[1].data, &maxval_buf[1].step, maxval_buf[1].cols * sizeof(T), maxval_buf[1].rows)); @@ -528,4 +529,219 @@ namespace cv { namespace gpu { namespace mathfunc template void min_max_caller(const DevMem2D, double*, double*); template void min_max_caller(const DevMem2D, double*, double*); + template struct OptLoc {}; + + template + struct OptLoc + { + static __device__ void call(unsigned int tid, unsigned int offset, volatile T* optval, volatile unsigned int* optloc) + { + T val = optval[tid + offset]; + if (val < optval[tid]) + { + optval[tid] = val; + optloc[tid] = optloc[tid + offset]; + } + } + }; + + template + struct OptLoc + { + static __device__ void call(unsigned int tid, unsigned int offset, volatile T* optval, volatile unsigned int* optloc) + { + T val = optval[tid + offset]; + if (val > optval[tid]) + { + optval[tid] = val; + optloc[tid] = optloc[tid + offset]; + } + } + }; + + template + __global__ void opt_loc_init_kernel(int cols, int rows, const PtrStep src, PtrStep optval, PtrStep optloc) + { + typedef typename MinMaxTypeTraits::best_type best_type; + __shared__ best_type soptval[nthreads]; + __shared__ unsigned int soptloc[nthreads]; + + unsigned int x0 = blockIdx.x * blockDim.x; + unsigned int y0 = blockIdx.y * blockDim.y; + unsigned int tid = threadIdx.y * blockDim.x + threadIdx.x; + + if (x0 + threadIdx.x < cols && y0 + threadIdx.y < rows) + { + soptval[tid] = ((const T*)src.ptr(y0 + threadIdx.y))[x0 + threadIdx.x]; + soptloc[tid] = (y0 + threadIdx.y) * cols + x0 + threadIdx.x; + } + else + { + soptval[tid] = ((const T*)src.ptr(y0))[x0]; + soptloc[tid] = y0 * cols + x0; + } + + __syncthreads(); + + if (nthreads >= 512) if (tid < 256) { OptLoc::call(tid, 256, soptval, soptloc); __syncthreads(); } + if (nthreads >= 256) if (tid < 128) { OptLoc::call(tid, 128, soptval, soptloc); __syncthreads(); } + if (nthreads >= 128) if (tid < 64) { OptLoc::call(tid, 64, soptval, soptloc); __syncthreads(); } + + if (tid < 32) + { + if (nthreads >= 64) OptLoc::call(tid, 32, soptval, soptloc); + if (nthreads >= 32) OptLoc::call(tid, 16, soptval, soptloc); + if (nthreads >= 16) OptLoc::call(tid, 8, soptval, soptloc); + if (nthreads >= 8) OptLoc::call(tid, 4, soptval, soptloc); + if (nthreads >= 4) OptLoc::call(tid, 2, soptval, soptloc); + if (nthreads >= 2) OptLoc::call(tid, 1, soptval, soptloc); + } + + if (tid == 0) + { + ((T*)optval.ptr(blockIdx.y))[blockIdx.x] = (T)soptval[0]; + ((unsigned int*)optloc.ptr(blockIdx.y))[blockIdx.x] = soptloc[0]; + } + } + + template + __global__ void opt_loc_kernel(int cols, int rows, const PtrStep src, const PtrStep loc, PtrStep optval, PtrStep optloc) + { + typedef typename MinMaxTypeTraits::best_type best_type; + __shared__ best_type soptval[nthreads]; + __shared__ unsigned int soptloc[nthreads]; + + unsigned int x0 = blockIdx.x * blockDim.x; + unsigned int y0 = blockIdx.y * blockDim.y; + unsigned int tid = threadIdx.y * blockDim.x + threadIdx.x; + + if (x0 + threadIdx.x < cols && y0 + threadIdx.y < rows) + { + soptval[tid] = ((const T*)src.ptr(y0 + threadIdx.y))[x0 + threadIdx.x]; + soptloc[tid] = ((const unsigned int*)loc.ptr(y0 + threadIdx.y))[x0 + threadIdx.x]; + } + else + { + soptval[tid] = ((const T*)src.ptr(y0))[x0]; + soptloc[tid] = ((const unsigned int*)loc.ptr(y0))[x0]; + } + + __syncthreads(); + + if (nthreads >= 512) if (tid < 256) { OptLoc::call(tid, 256, soptval, soptloc); __syncthreads(); } + if (nthreads >= 256) if (tid < 128) { OptLoc::call(tid, 128, soptval, soptloc); __syncthreads(); } + if (nthreads >= 128) if (tid < 64) { OptLoc::call(tid, 64, soptval, soptloc); __syncthreads(); } + + if (tid < 32) + { + if (nthreads >= 64) OptLoc::call(tid, 32, soptval, soptloc); + if (nthreads >= 32) OptLoc::call(tid, 16, soptval, soptloc); + if (nthreads >= 16) OptLoc::call(tid, 8, soptval, soptloc); + if (nthreads >= 8) OptLoc::call(tid, 4, soptval, soptloc); + if (nthreads >= 4) OptLoc::call(tid, 2, soptval, soptloc); + if (nthreads >= 2) OptLoc::call(tid, 1, soptval, soptloc); + } + + if (tid == 0) + { + ((T*)optval.ptr(blockIdx.y))[blockIdx.x] = (T)soptval[0]; + ((unsigned int*)optloc.ptr(blockIdx.y))[blockIdx.x] = soptloc[0]; + } + } + + template + void min_max_loc_caller(const DevMem2D src, double* minval, double* maxval, int* minlocx, int* minlocy, + int* maxlocx, int* maxlocy) + { + dim3 threads(32, 8); + + // Allocate memory for aux. buffers + + DevMem2D minval_buf[2]; + minval_buf[0].cols = divUp(src.cols, threads.x); + minval_buf[0].rows = divUp(src.rows, threads.y); + minval_buf[1].cols = divUp(minval_buf[0].cols, threads.x); + minval_buf[1].rows = divUp(minval_buf[0].rows, threads.y); + cudaSafeCall(cudaMallocPitch(&minval_buf[0].data, &minval_buf[0].step, minval_buf[0].cols * sizeof(T), minval_buf[0].rows)); + cudaSafeCall(cudaMallocPitch(&minval_buf[1].data, &minval_buf[1].step, minval_buf[1].cols * sizeof(T), minval_buf[1].rows)); + + DevMem2D maxval_buf[2]; + maxval_buf[0].cols = divUp(src.cols, threads.x); + maxval_buf[0].rows = divUp(src.rows, threads.y); + maxval_buf[1].cols = divUp(maxval_buf[0].cols, threads.x); + maxval_buf[1].rows = divUp(maxval_buf[0].rows, threads.y); + cudaSafeCall(cudaMallocPitch(&maxval_buf[0].data, &maxval_buf[0].step, maxval_buf[0].cols * sizeof(T), maxval_buf[0].rows)); + cudaSafeCall(cudaMallocPitch(&maxval_buf[1].data, &maxval_buf[1].step, maxval_buf[1].cols * sizeof(T), maxval_buf[1].rows)); + + DevMem2D minloc_buf[2]; + minloc_buf[0].cols = divUp(src.cols, threads.x); + minloc_buf[0].rows = divUp(src.rows, threads.y); + minloc_buf[1].cols = divUp(minloc_buf[0].cols, threads.x); + minloc_buf[1].rows = divUp(minloc_buf[0].rows, threads.y); + cudaSafeCall(cudaMallocPitch(&minloc_buf[0].data, &minloc_buf[0].step, minloc_buf[0].cols * sizeof(int), minloc_buf[0].rows)); + cudaSafeCall(cudaMallocPitch(&minloc_buf[1].data, &minloc_buf[1].step, minloc_buf[1].cols * sizeof(int), minloc_buf[1].rows)); + + DevMem2D maxloc_buf[2]; + maxloc_buf[0].cols = divUp(src.cols, threads.x); + maxloc_buf[0].rows = divUp(src.rows, threads.y); + maxloc_buf[1].cols = divUp(maxloc_buf[0].cols, threads.x); + maxloc_buf[1].rows = divUp(maxloc_buf[0].rows, threads.y); + cudaSafeCall(cudaMallocPitch(&maxloc_buf[0].data, &maxloc_buf[0].step, maxloc_buf[0].cols * sizeof(int), maxloc_buf[0].rows)); + cudaSafeCall(cudaMallocPitch(&maxloc_buf[1].data, &maxloc_buf[1].step, maxloc_buf[1].cols * sizeof(int), maxloc_buf[1].rows)); + + int curbuf = 0; + dim3 cursize(src.cols, src.rows); + dim3 grid(divUp(cursize.x, threads.x), divUp(cursize.y, threads.y)); + + opt_loc_init_kernel<256, MIN, T><<>>(cursize.x, cursize.y, src, minval_buf[curbuf], minloc_buf[curbuf]); + opt_loc_init_kernel<256, MAX, T><<>>(cursize.x, cursize.y, src, maxval_buf[curbuf], maxloc_buf[curbuf]); + cursize = grid; + + while (cursize.x > 1 || cursize.y > 1) + { + grid.x = divUp(cursize.x, threads.x); + grid.y = divUp(cursize.y, threads.y); + opt_loc_kernel<256, MIN, T><<>>(cursize.x, cursize.y, minval_buf[curbuf], minloc_buf[curbuf], + minval_buf[1 - curbuf], minloc_buf[1 - curbuf]); + opt_loc_kernel<256, MAX, T><<>>(cursize.x, cursize.y, maxval_buf[curbuf], maxloc_buf[curbuf], + maxval_buf[1 - curbuf], maxloc_buf[1 - curbuf]); + curbuf = 1 - curbuf; + cursize = grid; + } + + cudaSafeCall(cudaThreadSynchronize()); + + // Copy results from device to host + + T minval_, maxval_; + cudaSafeCall(cudaMemcpy(&minval_, minval_buf[curbuf].ptr(0), sizeof(T), cudaMemcpyDeviceToHost)); + cudaSafeCall(cudaMemcpy(&maxval_, maxval_buf[curbuf].ptr(0), sizeof(T), cudaMemcpyDeviceToHost)); + *minval = minval_; + *maxval = maxval_; + + unsigned int minloc, maxloc; + cudaSafeCall(cudaMemcpy(&minloc, minloc_buf[curbuf].ptr(0), sizeof(int), cudaMemcpyDeviceToHost)); + cudaSafeCall(cudaMemcpy(&maxloc, maxloc_buf[curbuf].ptr(0), sizeof(int), cudaMemcpyDeviceToHost)); + *minlocy = minloc / src.cols; *minlocx = minloc - *minlocy * src.cols; + *maxlocy = maxloc / src.cols; *maxlocx = maxloc - *maxlocy * src.cols; + + // Release aux. buffers + cudaSafeCall(cudaFree(minval_buf[0].data)); + cudaSafeCall(cudaFree(minval_buf[1].data)); + cudaSafeCall(cudaFree(maxval_buf[0].data)); + cudaSafeCall(cudaFree(maxval_buf[1].data)); + cudaSafeCall(cudaFree(minloc_buf[0].data)); + cudaSafeCall(cudaFree(minloc_buf[1].data)); + cudaSafeCall(cudaFree(maxloc_buf[0].data)); + cudaSafeCall(cudaFree(maxloc_buf[1].data)); + } + + template void min_max_loc_caller(const DevMem2D, double*, double*, int*, int*, int*, int*); + template void min_max_loc_caller(const DevMem2D, double*, double*, int*, int*, int*, int*); + template void min_max_loc_caller(const DevMem2D, double*, double*, int*, int*, int*, int*); + template void min_max_loc_caller(const DevMem2D, double*, double*, int*, int*, int*, int*); + template void min_max_loc_caller(const DevMem2D, double*, double*, int*, int*, int*, int*); + template void min_max_loc_caller(const DevMem2D, double*, double*, int*, int*, int*, int*); + template void min_max_loc_caller(const DevMem2D, double*, double*, int*, int*, int*, int*); + }}} diff --git a/tests/gpu/src/arithm.cpp b/tests/gpu/src/arithm.cpp index f8b65ac..daabc9f 100644 --- a/tests/gpu/src/arithm.cpp +++ b/tests/gpu/src/arithm.cpp @@ -733,6 +733,71 @@ struct CV_GpuMinMaxTest: public CvTest }; +//////////////////////////////////////////////////////////////////////////////// +// Min max loc + +struct CV_GpuMinMaxLocTest: public CvTest +{ + CV_GpuMinMaxLocTest(): CvTest("GPU-MinMaxLocTest", "minMaxLoc") {} + + void run(int) + { + for (int depth = CV_8U; depth <= CV_64F; ++depth) + { + int rows = 1, cols = 3; + test(rows, cols, depth); + for (int i = 0; i < 4; ++i) + { + int rows = 1 + rand() % 1000; + int cols = 1 + rand() % 1000; + test(rows, cols, depth); + } + } + } + + void test(int rows, int cols, int depth) + { + cv::Mat src(rows, cols, depth); + cv::RNG rng; + for (int i = 0; i < src.rows; ++i) + { + Mat row(1, src.cols * src.elemSize(), CV_8U, src.ptr(i)); + rng.fill(row, RNG::UNIFORM, Scalar(0), Scalar(255)); + } + + double minVal, maxVal; + cv::Point minLoc, maxLoc; + + if (depth != CV_8S) + cv::minMaxLoc(src, &minVal, &maxVal, &minLoc, &maxLoc); + else + { + // OpenCV's minMaxLoc doesn't support CV_8S type + minVal = std::numeric_limits::max(); + maxVal = std::numeric_limits::min(); + for (int i = 0; i < src.rows; ++i) + for (int j = 0; j < src.cols; ++j) + { + char val = src.at(i, j); + if (val < minVal) { minVal = val; minLoc = cv::Point(j, i); } + if (val > maxVal) { maxVal = val; maxLoc = cv::Point(j, i); } + } + } + + double minVal_, maxVal_; + cv::Point minLoc_, maxLoc_; + cv::gpu::minMaxLoc(cv::gpu::GpuMat(src), &minVal_, &maxVal_, &minLoc_, &maxLoc_); + + CHECK(minVal == minVal_, CvTS::FAIL_INVALID_OUTPUT); + CHECK(maxVal == maxVal_, CvTS::FAIL_INVALID_OUTPUT); + CHECK(0 == memcmp(src.ptr(minLoc.y) + minLoc.x * src.elemSize(), src.ptr(minLoc_.y) + minLoc_.x * src.elemSize(), src.elemSize()), + CvTS::FAIL_INVALID_OUTPUT); + CHECK(0 == memcmp(src.ptr(maxLoc.y) + maxLoc.x * src.elemSize(), src.ptr(maxLoc_.y) + maxLoc_.x * src.elemSize(), src.elemSize()), + CvTS::FAIL_INVALID_OUTPUT); + } +}; + + ///////////////////////////////////////////////////////////////////////////// /////////////////// tests registration ///////////////////////////////////// ///////////////////////////////////////////////////////////////////////////// @@ -760,3 +825,4 @@ CV_GpuNppImagePhaseTest CV_GpuNppImagePhase_test; CV_GpuNppImageCartToPolarTest CV_GpuNppImageCartToPolar_test; CV_GpuNppImagePolarToCartTest CV_GpuNppImagePolarToCart_test; CV_GpuMinMaxTest CV_GpuMinMaxTest_test; +CV_GpuMinMaxLocTest CV_GpuMinMaxLocTest_test; -- 2.7.4