From 57eaf8f3aa53a23e5bba2e83b6a4735a33c99754 Mon Sep 17 00:00:00 2001 From: PiotrKarwat Date: Thu, 18 Nov 2021 20:41:43 +0100 Subject: [PATCH 1/3] iqRaw2Lri - omit initial samples --- api/matlab/arrus/mexcuda/iqRaw2Lri.cu | 38 ++++++++++++++++++++------- 1 file changed, 28 insertions(+), 10 deletions(-) diff --git a/api/matlab/arrus/mexcuda/iqRaw2Lri.cu b/api/matlab/arrus/mexcuda/iqRaw2Lri.cu index 9bfc82ba5..6131d9378 100644 --- a/api/matlab/arrus/mexcuda/iqRaw2Lri.cu +++ b/api/matlab/arrus/mexcuda/iqRaw2Lri.cu @@ -21,7 +21,7 @@ __global__ void iqRaw2Lri( float2 * iqLri, float const * zPix, float const * xP float const * txFoc, float const * txAngZX, float const * txApCentZ, float const * txApCentX, int const * txApFstElem, int const * txApLstElem, - int const * rxApOrigElem, + int const * rxApOrigElem, int const * nSampOmit, float const minRxTang, float const maxRxTang, float const fs, float const fn, float const sos, float const initDel, @@ -114,7 +114,7 @@ __global__ void iqRaw2Lri( float2 * iqLri, float const * zPix, float const * xP time = (txDist + rxDist) * sosInv + initDel; iSamp = time * fs; - if (iSamp<0.f || iSamp>static_cast(nSamp-1)) continue; + if (iSamp(nSampOmit[iTx]) || iSamp>static_cast(nSamp-1)) continue; float2 iqSamp = tex1DLayered(iqRawTex, iSamp + 0.5f, iRx + iTx*nRx); sampRe = iqSamp.x; @@ -184,6 +184,7 @@ void mexFunction(int nlhs, mxArray * plhs[], mxGPUArray const * elemFst; mxGPUArray const * elemLst; mxGPUArray const * rxElemOrig; + mxGPUArray const * nSampOmit; float2 * dev_iqLri; float2 const * dev_iqRaw; @@ -199,6 +200,7 @@ void mexFunction(int nlhs, mxArray * plhs[], int const * dev_elemFst; int const * dev_elemLst; int const * dev_rxElemOrig; + int const * dev_nSampOmit; float minRxTang; float maxRxTang; @@ -222,8 +224,8 @@ void mexFunction(int nlhs, mxArray * plhs[], char const * const invalidOutputMsgId = "iqRaw2Lri:InvalidOutput"; /* Validate mex inputs/outputs */ - if (nrhs!=19) { - mexErrMsgIdAndTxt( invalidInputMsgId, "19 inputs required"); + if (nrhs!=20) { + mexErrMsgIdAndTxt( invalidInputMsgId, "20 inputs required"); } if (nlhs>1) { @@ -251,13 +253,14 @@ void mexFunction(int nlhs, mxArray * plhs[], elemFst = mxGPUCreateFromMxArray(prhs[10]); elemLst = mxGPUCreateFromMxArray(prhs[11]); rxElemOrig = mxGPUCreateFromMxArray(prhs[12]); + nSampOmit = mxGPUCreateFromMxArray(prhs[13]); - minRxTang = mxGetScalar(prhs[13]); - maxRxTang = mxGetScalar(prhs[14]); - fs = mxGetScalar(prhs[15]); - fn = mxGetScalar(prhs[16]); - sos = mxGetScalar(prhs[17]); - initDel = mxGetScalar(prhs[18]); + minRxTang = mxGetScalar(prhs[14]); + maxRxTang = mxGetScalar(prhs[15]); + fs = mxGetScalar(prhs[16]); + fn = mxGetScalar(prhs[17]); + sos = mxGetScalar(prhs[18]); + initDel = mxGetScalar(prhs[19]); /* Validate inputs */ checkData(iqRaw, "iqRaw", true, 3, invalidInputMsgId); @@ -273,6 +276,7 @@ void mexFunction(int nlhs, mxArray * plhs[], // checkData(elemFst, "elemFst", false, 1, invalidInputMsgId); //int // checkData(elemLst, "elemLst", false, 1, invalidInputMsgId); //int // checkData(rxElemOrig,"rxElemOrig",false, 1, invalidInputMsgId); //int +// checkData(nSampOmit,"nSampOmit",false, 1, invalidInputMsgId); //int /* Get some additional information */ nSamp = mxGPUGetDimensions(iqRaw)[0]; @@ -316,6 +320,7 @@ void mexFunction(int nlhs, mxArray * plhs[], dev_elemFst = static_cast(mxGPUGetDataReadOnly(elemFst)); dev_elemLst = static_cast(mxGPUGetDataReadOnly(elemLst)); dev_rxElemOrig = static_cast(mxGPUGetDataReadOnly(rxElemOrig)); + dev_nSampOmit = static_cast(mxGPUGetDataReadOnly(nSampOmit)); /* set constant memory */ if(nElem > 256) { @@ -359,6 +364,17 @@ void mexFunction(int nlhs, mxArray * plhs[], // dev_centX + iPart*nTxPerPart, // minRxTang, maxRxTang, fs, fn, sos, initDel, // nZPix, nXPix, nSamp, nElem, nTxInThisPart); +// iqRaw2Lri<<>>(dev_iqLri + iPart*nZPix*nXPix*nTxPerPart, +// dev_zPix, dev_xPix, +// dev_foc + iPart*nTxPerPart, +// dev_ang + iPart*nTxPerPart, +// dev_centZ + iPart*nTxPerPart, +// dev_centX + iPart*nTxPerPart, +// dev_elemFst + iPart*nTxPerPart, +// dev_elemLst + iPart*nTxPerPart, +// dev_rxElemOrig + iPart*nTxPerPart, +// minRxTang, maxRxTang, fs, fn, sos, initDel, +// nZPix, nXPix, nSamp, nElem, nRx, nTxInThisPart); iqRaw2Lri<<>>(dev_iqLri + iPart*nZPix*nXPix*nTxPerPart, dev_zPix, dev_xPix, dev_foc + iPart*nTxPerPart, @@ -368,6 +384,7 @@ void mexFunction(int nlhs, mxArray * plhs[], dev_elemFst + iPart*nTxPerPart, dev_elemLst + iPart*nTxPerPart, dev_rxElemOrig + iPart*nTxPerPart, + dev_nSampOmit + iPart*nTxPerPart, minRxTang, maxRxTang, fs, fn, sos, initDel, nZPix, nXPix, nSamp, nElem, nRx, nTxInThisPart); @@ -394,6 +411,7 @@ void mexFunction(int nlhs, mxArray * plhs[], mxGPUDestroyGPUArray(elemFst); mxGPUDestroyGPUArray(elemLst); mxGPUDestroyGPUArray(rxElemOrig); + mxGPUDestroyGPUArray(nSampOmit); //cudaDeviceReset(); } From 3cdd93ccbc1164a1bf762f5650b31e080041e6b6 Mon Sep 17 00:00:00 2001 From: PiotrKarwat Date: Thu, 18 Nov 2021 20:42:50 +0100 Subject: [PATCH 2/3] Us4R - pass the nSampOmit to the reconstruction kernel --- api/matlab/arrus/Us4R.m | 2 ++ 1 file changed, 2 insertions(+) diff --git a/api/matlab/arrus/Us4R.m b/api/matlab/arrus/Us4R.m index e2dcfa87a..cd30649f0 100644 --- a/api/matlab/arrus/Us4R.m +++ b/api/matlab/arrus/Us4R.m @@ -606,6 +606,7 @@ function setRecParams(obj,varargin) obj.seq.txApFstElem = gpuArray( int32(obj.seq.txApFstElem - 1)); obj.seq.txApLstElem = gpuArray( int32(obj.seq.txApLstElem - 1)); obj.seq.rxApFstElem = gpuArray( int32(obj.seq.rxApOrig - 1)); % rxApOrig remains unchanged as it is used in data reorganization + obj.seq.nSampOmit = gpuArray( int32(((max(obj.seq.txDel) + obj.seq.txNPer/obj.seq.txFreq) * obj.seq.rxSampFreq + 50) / obj.rec.dec)); obj.rec.bmodeRxTangLim = single(obj.rec.bmodeRxTangLim); obj.rec.colorRxTangLim = single(obj.rec.colorRxTangLim); obj.rec.vect0RxTangLim = single(obj.rec.vect0RxTangLim); @@ -1171,6 +1172,7 @@ function closeSequence(obj) obj.seq.txApFstElem(selFrames), ... obj.seq.txApLstElem(selFrames), ... obj.seq.rxApFstElem(selFrames), ... + obj.seq.nSampOmit(selFrames), ... rxTangLim(1), ... rxTangLim(2), ... obj.seq.rxSampFreq / obj.rec.dec, ... From 7da7811bd086030b0442ceccd81bb7361b4acf77 Mon Sep 17 00:00:00 2001 From: PiotrKarwat Date: Fri, 19 Nov 2021 20:38:57 +0100 Subject: [PATCH 3/3] iqRaw2Lri - remove commented piece of code --- api/matlab/arrus/mexcuda/iqRaw2Lri.cu | 11 ----------- 1 file changed, 11 deletions(-) diff --git a/api/matlab/arrus/mexcuda/iqRaw2Lri.cu b/api/matlab/arrus/mexcuda/iqRaw2Lri.cu index 6131d9378..373aced10 100644 --- a/api/matlab/arrus/mexcuda/iqRaw2Lri.cu +++ b/api/matlab/arrus/mexcuda/iqRaw2Lri.cu @@ -364,17 +364,6 @@ void mexFunction(int nlhs, mxArray * plhs[], // dev_centX + iPart*nTxPerPart, // minRxTang, maxRxTang, fs, fn, sos, initDel, // nZPix, nXPix, nSamp, nElem, nTxInThisPart); -// iqRaw2Lri<<>>(dev_iqLri + iPart*nZPix*nXPix*nTxPerPart, -// dev_zPix, dev_xPix, -// dev_foc + iPart*nTxPerPart, -// dev_ang + iPart*nTxPerPart, -// dev_centZ + iPart*nTxPerPart, -// dev_centX + iPart*nTxPerPart, -// dev_elemFst + iPart*nTxPerPart, -// dev_elemLst + iPart*nTxPerPart, -// dev_rxElemOrig + iPart*nTxPerPart, -// minRxTang, maxRxTang, fs, fn, sos, initDel, -// nZPix, nXPix, nSamp, nElem, nRx, nTxInThisPart); iqRaw2Lri<<>>(dev_iqLri + iPart*nZPix*nXPix*nTxPerPart, dev_zPix, dev_xPix, dev_foc + iPart*nTxPerPart,