complexmat.cu: Remove unused member foreign_data

[hercules2020/kcf.git] / src / complexmat.cu

diff --git a/src/complexmat.cu b/src/complexmat.cu
index 7d4a43dbceece696d6dd03fd923f6f9d035ff6f9..d06e66fa4cc1dcdc7ef90d6c30473026393a9cdf 100644 (file)
--- a/src/complexmat.cu
+++ b/src/complexmat.cu
@@ -27,15 +27,15 @@ __global__ void sqr_norm_kernel(int n, float *out, float *data, float rows, floa
     }
 }
 
-void ComplexMat::sqr_norm(float *result) const
+void ComplexMat::sqr_norm(DynMem &result) const
 {
-    CudaSafeCall(cudaMemsetAsync(result, 0, n_scales * sizeof(float), this->stream));
+    CudaSafeCall(cudaMemsetAsync(result.deviceMem(), 0, n_scales * sizeof(float)));
 
     dim3 threadsPerBlock(rows, cols);
     dim3 numBlocks(n_channels / n_scales, n_scales);
 
-    sqr_norm_kernel<<<numBlocks, threadsPerBlock, rows * cols * sizeof(float), this->stream>>>(
-        n_channels / n_scales, result, this->p_data, rows, cols);
+    sqr_norm_kernel<<<numBlocks, threadsPerBlock, rows * cols * sizeof(float)>>>(
+        n_channels / n_scales, result.deviceMem(), this->p_data, rows, cols);
     CudaCheckError();
 
     return;
@@ -52,11 +52,11 @@ __global__ void sqr_mag_kernel(float *data, float *result)
 
 ComplexMat ComplexMat::sqr_mag() const
 {
-    ComplexMat result(this->rows, this->cols, this->channels(), this->n_scales, this->stream);
+    ComplexMat result(this->rows, this->cols, this->channels(), this->n_scales);
 
     dim3 threadsPerBlock(rows, cols);
     dim3 numBlocks(n_channels / n_scales, n_scales);
-    sqr_mag_kernel<<<numBlocks, threadsPerBlock, 0, this->stream>>>(this->p_data, result.p_data);
+    sqr_mag_kernel<<<numBlocks, threadsPerBlock, 0>>>(this->p_data, result.p_data);
     CudaCheckError();
 
     return result;
@@ -73,11 +73,11 @@ __global__ void conj_kernel(float *data, float *result)
 
 ComplexMat ComplexMat::conj() const
 {
-    ComplexMat result(this->rows, this->cols, this->channels(), this->n_scales, this->stream);
+    ComplexMat result(this->rows, this->cols, this->channels(), this->n_scales);
 
     dim3 threadsPerBlock(rows, cols);
     dim3 numBlocks(n_channels / n_scales, n_scales);
-    conj_kernel<<<numBlocks, threadsPerBlock, 0, this->stream>>>(this->p_data, result.p_data);
+    conj_kernel<<<numBlocks, threadsPerBlock, 0>>>(this->p_data, result.p_data);
     CudaCheckError();
 
     return result;
@@ -86,7 +86,7 @@ ComplexMat ComplexMat::conj() const
 ComplexMat ComplexMat::sum_over_channels() const
 {
     //     assert(p_data.size() > 1);
-    ComplexMat result(this->rows, this->cols, 1, this->stream);
+    ComplexMat result(this->rows, this->cols, 1);
     return result;
 }
 
@@ -109,11 +109,11 @@ ComplexMat ComplexMat::operator*(const ComplexMat &rhs) const
 {
     assert(rhs.n_channels == n_channels && rhs.cols == cols && rhs.rows == rows);
 
-    ComplexMat result(this->rows, this->cols, this->channels(), this->n_scales, this->stream);
+    ComplexMat result(this->rows, this->cols, this->channels(), this->n_scales);
 
     dim3 threadsPerBlock(rows, cols);
     dim3 numBlocks(n_channels / n_scales, n_scales);
-    same_num_channels_mul_kernel<<<numBlocks, threadsPerBlock, 0, this->stream>>>(this->p_data, rhs.p_data,
+    same_num_channels_mul_kernel<<<numBlocks, threadsPerBlock, 0>>>(this->p_data, rhs.p_data,
                                                                                   result.p_data);
     CudaCheckError();
 
@@ -135,11 +135,11 @@ ComplexMat ComplexMat::operator/(const ComplexMat &rhs) const
 {
     assert(rhs.n_channels == n_channels && rhs.cols == cols && rhs.rows == rows);
 
-    ComplexMat result(this->rows, this->cols, this->channels(), this->n_scales, this->stream);
+    ComplexMat result(this->rows, this->cols, this->channels(), this->n_scales);
 
     dim3 threadsPerBlock(rows, cols);
     dim3 numBlocks(n_channels / n_scales, n_scales);
-    same_num_channels_div_kernel<<<numBlocks, threadsPerBlock, 0, this->stream>>>(this->p_data, rhs.p_data,
+    same_num_channels_div_kernel<<<numBlocks, threadsPerBlock, 0>>>(this->p_data, rhs.p_data,
                                                                                   result.p_data);
     CudaCheckError();
 
@@ -159,11 +159,11 @@ ComplexMat ComplexMat::operator+(const ComplexMat &rhs) const
 {
     assert(rhs.n_channels == n_channels && rhs.cols == cols && rhs.rows == rows);
 
-    ComplexMat result(this->rows, this->cols, this->channels(), this->n_scales, this->stream);
+    ComplexMat result(this->rows, this->cols, this->channels(), this->n_scales);
 
     dim3 threadsPerBlock(rows, cols);
     dim3 numBlocks(n_channels / n_scales, n_scales);
-    same_num_channels_add_kernel<<<numBlocks, threadsPerBlock, 0, this->stream>>>(this->p_data, rhs.p_data,
+    same_num_channels_add_kernel<<<numBlocks, threadsPerBlock, 0>>>(this->p_data, rhs.p_data,
                                                                                   result.p_data);
     CudaCheckError();
 
@@ -181,11 +181,11 @@ __global__ void constant_mul_kernel(float *data_l, float constant, float *result
 
 ComplexMat ComplexMat::operator*(const float &rhs) const
 {
-    ComplexMat result(this->rows, this->cols, this->channels(), this->n_scales, this->stream);
+    ComplexMat result(this->rows, this->cols, this->channels(), this->n_scales);
 
     dim3 threadsPerBlock(rows, cols);
     dim3 numBlocks(n_channels / n_scales, n_scales);
-    constant_mul_kernel<<<numBlocks, threadsPerBlock, 0, this->stream>>>(this->p_data, rhs, result.p_data);
+    constant_mul_kernel<<<numBlocks, threadsPerBlock, 0>>>(this->p_data, rhs, result.p_data);
     CudaCheckError();
 
     return result;
@@ -202,11 +202,11 @@ __global__ void constant_add_kernel(float *data_l, float constant, float *result
 
 ComplexMat ComplexMat::operator+(const float &rhs) const
 {
-    ComplexMat result(this->rows, this->cols, this->channels(), this->n_scales, this->stream);
+    ComplexMat result(this->rows, this->cols, this->channels(), this->n_scales);
 
     dim3 threadsPerBlock(rows, cols);
     dim3 numBlocks(n_channels / n_scales, n_scales);
-    constant_add_kernel<<<numBlocks, threadsPerBlock, 0, this->stream>>>(this->p_data, rhs, result.p_data);
+    constant_add_kernel<<<numBlocks, threadsPerBlock, 0>>>(this->p_data, rhs, result.p_data);
     CudaCheckError();
 
     return result;
@@ -227,11 +227,11 @@ ComplexMat ComplexMat::mul(const ComplexMat &rhs) const
 {
     assert(rhs.n_channels == 1 && rhs.cols == cols && rhs.rows == rows);
 
-    ComplexMat result(this->rows, this->cols, this->channels(), this->n_scales, this->stream);
+    ComplexMat result(this->rows, this->cols, this->channels(), this->n_scales);
 
     dim3 threadsPerBlock(rows, cols);
     dim3 numBlocks(n_channels / n_scales, n_scales);
-    one_channel_mul_kernel<<<numBlocks, threadsPerBlock, 0, this->stream>>>(this->p_data, rhs.p_data, result.p_data);
+    one_channel_mul_kernel<<<numBlocks, threadsPerBlock, 0>>>(this->p_data, rhs.p_data, result.p_data);
     CudaCheckError();
 
     return result;
@@ -252,11 +252,11 @@ ComplexMat ComplexMat::mul2(const ComplexMat &rhs) const
 {
     assert(rhs.n_channels == n_channels / n_scales && rhs.cols == cols && rhs.rows == rows);
 
-    ComplexMat result(this->rows, this->cols, this->channels(), this->n_scales, this->stream);
+    ComplexMat result(this->rows, this->cols, this->channels(), this->n_scales);
 
     dim3 threadsPerBlock(rows, cols);
     dim3 numBlocks(n_channels / n_scales, n_scales);
-    scales_channel_mul_kernel<<<numBlocks, threadsPerBlock, 0, this->stream>>>(this->p_data, rhs.p_data, result.p_data);
+    scales_channel_mul_kernel<<<numBlocks, threadsPerBlock, 0>>>(this->p_data, rhs.p_data, result.p_data);
     CudaCheckError();
 
     return result;
@@ -268,8 +268,6 @@ void ComplexMat::operator=(ComplexMat &rhs)
     rows = rhs.rows;
     n_channels = rhs.n_channels;
     n_scales = rhs.n_scales;
-    stream = rhs.stream;
-    foreign_data = true;
 
     p_data = rhs.p_data;
 }
@@ -280,7 +278,6 @@ void ComplexMat::operator=(ComplexMat &&rhs)
     rows = rhs.rows;
     n_channels = rhs.n_channels;
     n_scales = rhs.n_scales;
-    stream = rhs.stream;
 
     p_data = rhs.p_data;