From 231f8254e6c04e21a7bb84c7cde704c3e8eea9ea Mon Sep 17 00:00:00 2001 From: PiotrKarwat Date: Mon, 22 Nov 2021 08:56:59 +0100 Subject: [PATCH 1/6] iqRaw2Lri - custom apodization (uses texture) Replaces previous (fixed gaussian) apodization --- api/matlab/arrus/mexcuda/iqRaw2Lri.cu | 61 ++++++++++++++++++--------- 1 file changed, 40 insertions(+), 21 deletions(-) diff --git a/api/matlab/arrus/mexcuda/iqRaw2Lri.cu b/api/matlab/arrus/mexcuda/iqRaw2Lri.cu index 373aced10..1a2edf818 100644 --- a/api/matlab/arrus/mexcuda/iqRaw2Lri.cu +++ b/api/matlab/arrus/mexcuda/iqRaw2Lri.cu @@ -9,7 +9,7 @@ __constant__ float xElemConst[256]; __constant__ float tangElemConst[256]; texture iqRawTex; - +texture rxApodTex; __forceinline__ __device__ float ownHypotf(float x, float y) { @@ -43,10 +43,7 @@ __global__ void iqRaw2Lri( float2 * iqLri, float const * zPix, float const * xP float const omega = 2 * M_PI * fn; float const sosInv = 1 / sos; // float const zDistInv = 1 / zPix[z]; - float const nSigma = 3; // number of sigmas in half of the apodization Gaussian curve - float const twoSigSqrInv = nSigma * nSigma * 0.5f; - float const rngRxTangInv = 2 / (maxRxTang - minRxTang); // inverted half range - float const centRxTang = (maxRxTang + minRxTang) * 0.5f; + float const rngRxTangInv = 1 / (maxRxTang - minRxTang); // inverted tangent range for (int iTx=0; iTx maxRxTang) continue; - rxApod = (rxTang-centRxTang)*rngRxTangInv; - rxApod = __expf(-rxApod*rxApod*twoSigSqrInv); + rxApod = (rxTang-minRxTang)*rngRxTangInv; // <0,1>, needs normalized texture fetching, errors at aperture sided + rxApod = tex1D(rxApodTex, rxApod); time = (txDist + rxDist) * sosInv + initDel; iSamp = time * fs; @@ -181,6 +178,7 @@ void mexFunction(int nlhs, mxArray * plhs[], mxGPUArray const * ang; mxGPUArray const * centZ; mxGPUArray const * centX; + mxGPUArray const * rxApod; mxGPUArray const * elemFst; mxGPUArray const * elemLst; mxGPUArray const * rxElemOrig; @@ -197,6 +195,7 @@ void mexFunction(int nlhs, mxArray * plhs[], float const * dev_ang; float const * dev_centZ; float const * dev_centX; + float const * dev_rxApod; int const * dev_elemFst; int const * dev_elemLst; int const * dev_rxElemOrig; @@ -215,6 +214,7 @@ void mexFunction(int nlhs, mxArray * plhs[], int nXPix; int nRx; int nTx; + int nRxApodSamp; dim3 const threadsPerBlock = {16, 16, 1}; dim3 blocksPerGrid; @@ -224,15 +224,15 @@ void mexFunction(int nlhs, mxArray * plhs[], char const * const invalidOutputMsgId = "iqRaw2Lri:InvalidOutput"; /* Validate mex inputs/outputs */ - if (nrhs!=20) { - mexErrMsgIdAndTxt( invalidInputMsgId, "20 inputs required"); + if (nrhs!=21) { + mexErrMsgIdAndTxt( invalidInputMsgId, "21 inputs required"); } if (nlhs>1) { mexErrMsgIdAndTxt( invalidOutputMsgId, "One output allowed"); } -// for (int i=13; i<19; i++) { +// for (int i=15; i<20; i++) { // if (!mxIsSingle(prhs[i]) || mxIsComplex(prhs[i]) || mxGetNumberOfElements(prhs[i]) != 1) { // mexErrMsgIdAndTxt( invalidInputMsgId, "Last 6 inputs must be single, real scalars"); // } @@ -250,17 +250,18 @@ void mexFunction(int nlhs, mxArray * plhs[], ang = mxGPUCreateFromMxArray(prhs[7]); centZ = mxGPUCreateFromMxArray(prhs[8]); centX = mxGPUCreateFromMxArray(prhs[9]); - elemFst = mxGPUCreateFromMxArray(prhs[10]); - elemLst = mxGPUCreateFromMxArray(prhs[11]); - rxElemOrig = mxGPUCreateFromMxArray(prhs[12]); - nSampOmit = mxGPUCreateFromMxArray(prhs[13]); - - minRxTang = mxGetScalar(prhs[14]); - maxRxTang = mxGetScalar(prhs[15]); - fs = mxGetScalar(prhs[16]); - fn = mxGetScalar(prhs[17]); - sos = mxGetScalar(prhs[18]); - initDel = mxGetScalar(prhs[19]); + rxApod = mxGPUCreateFromMxArray(prhs[10]); + elemFst = mxGPUCreateFromMxArray(prhs[11]); + elemLst = mxGPUCreateFromMxArray(prhs[12]); + rxElemOrig= mxGPUCreateFromMxArray(prhs[13]); + nSampOmit = mxGPUCreateFromMxArray(prhs[14]); + + minRxTang = mxGetScalar(prhs[15]); + maxRxTang = mxGetScalar(prhs[16]); + fs = mxGetScalar(prhs[17]); + fn = mxGetScalar(prhs[18]); + sos = mxGetScalar(prhs[19]); + initDel = mxGetScalar(prhs[20]); /* Validate inputs */ checkData(iqRaw, "iqRaw", true, 3, invalidInputMsgId); @@ -273,6 +274,7 @@ void mexFunction(int nlhs, mxArray * plhs[], checkData(ang, "ang", false, 1, invalidInputMsgId); checkData(centZ, "centZ", false, 1, invalidInputMsgId); checkData(centX, "centX", false, 1, invalidInputMsgId); + checkData(rxApod, "rxApod", false, 1, invalidInputMsgId); // checkData(elemFst, "elemFst", false, 1, invalidInputMsgId); //int // checkData(elemLst, "elemLst", false, 1, invalidInputMsgId); //int // checkData(rxElemOrig,"rxElemOrig",false, 1, invalidInputMsgId); //int @@ -284,6 +286,7 @@ void mexFunction(int nlhs, mxArray * plhs[], nElem = mxGPUGetNumberOfElements(xElem); nZPix = mxGPUGetNumberOfElements(zPix); nXPix = mxGPUGetNumberOfElements(xPix); + nRxApodSamp = mxGPUGetNumberOfElements(rxApod); if (mxGPUGetNumberOfDimensions(iqRaw)<3) { nTx = 1; } @@ -317,6 +320,7 @@ void mexFunction(int nlhs, mxArray * plhs[], dev_ang = static_cast(mxGPUGetDataReadOnly(ang)); dev_centZ = static_cast(mxGPUGetDataReadOnly(centZ)); dev_centX = static_cast(mxGPUGetDataReadOnly(centX)); + dev_rxApod = static_cast(mxGPUGetDataReadOnly(rxApod)); dev_elemFst = static_cast(mxGPUGetDataReadOnly(elemFst)); dev_elemLst = static_cast(mxGPUGetDataReadOnly(elemLst)); dev_rxElemOrig = static_cast(mxGPUGetDataReadOnly(rxElemOrig)); @@ -330,6 +334,17 @@ void mexFunction(int nlhs, mxArray * plhs[], cudaMemcpyToSymbol( xElemConst, dev_xElem, nElem*sizeof(float), 0, cudaMemcpyDeviceToDevice); cudaMemcpyToSymbol(tangElemConst, dev_tangElem, nElem*sizeof(float), 0, cudaMemcpyDeviceToDevice); + /* configure texture reference (apodization) */ + rxApodTex.normalized = true; + rxApodTex.addressMode[0] = cudaAddressModeBorder; + rxApodTex.filterMode = cudaFilterModeLinear; + + cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat); + cudaArray* cuArrayApod; + cudaMallocArray(&cuArrayApod, &channelDesc, nRxApodSamp, 0); + cudaMemcpyToArray(cuArrayApod, 0, 0, dev_rxApod, nRxApodSamp*sizeof(float), cudaMemcpyDeviceToDevice); + cudaBindTextureToArray(rxApodTex, cuArrayApod, channelDesc); + /* configure texture reference */ iqRawTex.normalized = false; iqRawTex.addressMode[0] = cudaAddressModeBorder; @@ -386,6 +401,9 @@ void mexFunction(int nlhs, mxArray * plhs[], cudaUnbindTexture(iqRawTex); cudaFreeArray(cuArray); + cudaUnbindTexture(rxApodTex); + cudaFreeArray(cuArrayApod); + mxGPUDestroyGPUArray(iqLri); mxGPUDestroyGPUArray(iqRaw); mxGPUDestroyGPUArray(zElem); @@ -397,6 +415,7 @@ void mexFunction(int nlhs, mxArray * plhs[], mxGPUDestroyGPUArray(ang); mxGPUDestroyGPUArray(centZ); mxGPUDestroyGPUArray(centX); + mxGPUDestroyGPUArray(rxApod); mxGPUDestroyGPUArray(elemFst); mxGPUDestroyGPUArray(elemLst); mxGPUDestroyGPUArray(rxElemOrig); From eff0bcf1ac19dee5b18ada9abf6b4825f5e2a043 Mon Sep 17 00:00:00 2001 From: PiotrKarwat Date: Mon, 22 Nov 2021 10:05:07 +0100 Subject: [PATCH 2/6] iqRaw2Lri - change rxApod position among other parameters --- api/matlab/arrus/mexcuda/iqRaw2Lri.cu | 20 ++++++++++---------- 1 file changed, 10 insertions(+), 10 deletions(-) diff --git a/api/matlab/arrus/mexcuda/iqRaw2Lri.cu b/api/matlab/arrus/mexcuda/iqRaw2Lri.cu index 1a2edf818..01304a7b4 100644 --- a/api/matlab/arrus/mexcuda/iqRaw2Lri.cu +++ b/api/matlab/arrus/mexcuda/iqRaw2Lri.cu @@ -174,11 +174,11 @@ void mexFunction(int nlhs, mxArray * plhs[], mxGPUArray const * tangElem; mxGPUArray const * zPix; mxGPUArray const * xPix; + mxGPUArray const * rxApod; mxGPUArray const * foc; mxGPUArray const * ang; mxGPUArray const * centZ; mxGPUArray const * centX; - mxGPUArray const * rxApod; mxGPUArray const * elemFst; mxGPUArray const * elemLst; mxGPUArray const * rxElemOrig; @@ -191,11 +191,11 @@ void mexFunction(int nlhs, mxArray * plhs[], float const * dev_tangElem; float const * dev_zPix; float const * dev_xPix; + float const * dev_rxApod; float const * dev_foc; float const * dev_ang; float const * dev_centZ; float const * dev_centX; - float const * dev_rxApod; int const * dev_elemFst; int const * dev_elemLst; int const * dev_rxElemOrig; @@ -246,11 +246,11 @@ void mexFunction(int nlhs, mxArray * plhs[], tangElem = mxGPUCreateFromMxArray(prhs[3]); zPix = mxGPUCreateFromMxArray(prhs[4]); xPix = mxGPUCreateFromMxArray(prhs[5]); - foc = mxGPUCreateFromMxArray(prhs[6]); - ang = mxGPUCreateFromMxArray(prhs[7]); - centZ = mxGPUCreateFromMxArray(prhs[8]); - centX = mxGPUCreateFromMxArray(prhs[9]); - rxApod = mxGPUCreateFromMxArray(prhs[10]); + rxApod = mxGPUCreateFromMxArray(prhs[6]); + foc = mxGPUCreateFromMxArray(prhs[7]); + ang = mxGPUCreateFromMxArray(prhs[8]); + centZ = mxGPUCreateFromMxArray(prhs[9]); + centX = mxGPUCreateFromMxArray(prhs[10]); elemFst = mxGPUCreateFromMxArray(prhs[11]); elemLst = mxGPUCreateFromMxArray(prhs[12]); rxElemOrig= mxGPUCreateFromMxArray(prhs[13]); @@ -270,11 +270,11 @@ void mexFunction(int nlhs, mxArray * plhs[], checkData(tangElem, "tangElem", false, 1, invalidInputMsgId); checkData(zPix, "zPix", false, 1, invalidInputMsgId); checkData(xPix, "xPix", false, 1, invalidInputMsgId); + checkData(rxApod, "rxApod", false, 1, invalidInputMsgId); checkData(foc, "foc", false, 1, invalidInputMsgId); checkData(ang, "ang", false, 1, invalidInputMsgId); checkData(centZ, "centZ", false, 1, invalidInputMsgId); checkData(centX, "centX", false, 1, invalidInputMsgId); - checkData(rxApod, "rxApod", false, 1, invalidInputMsgId); // checkData(elemFst, "elemFst", false, 1, invalidInputMsgId); //int // checkData(elemLst, "elemLst", false, 1, invalidInputMsgId); //int // checkData(rxElemOrig,"rxElemOrig",false, 1, invalidInputMsgId); //int @@ -316,11 +316,11 @@ void mexFunction(int nlhs, mxArray * plhs[], dev_tangElem = static_cast(mxGPUGetDataReadOnly(tangElem)); dev_zPix = static_cast(mxGPUGetDataReadOnly(zPix)); dev_xPix = static_cast(mxGPUGetDataReadOnly(xPix)); + dev_rxApod = static_cast(mxGPUGetDataReadOnly(rxApod)); dev_foc = static_cast(mxGPUGetDataReadOnly(foc)); dev_ang = static_cast(mxGPUGetDataReadOnly(ang)); dev_centZ = static_cast(mxGPUGetDataReadOnly(centZ)); dev_centX = static_cast(mxGPUGetDataReadOnly(centX)); - dev_rxApod = static_cast(mxGPUGetDataReadOnly(rxApod)); dev_elemFst = static_cast(mxGPUGetDataReadOnly(elemFst)); dev_elemLst = static_cast(mxGPUGetDataReadOnly(elemLst)); dev_rxElemOrig = static_cast(mxGPUGetDataReadOnly(rxElemOrig)); @@ -411,11 +411,11 @@ void mexFunction(int nlhs, mxArray * plhs[], mxGPUDestroyGPUArray(tangElem); mxGPUDestroyGPUArray(zPix); mxGPUDestroyGPUArray(xPix); + mxGPUDestroyGPUArray(rxApod); mxGPUDestroyGPUArray(foc); mxGPUDestroyGPUArray(ang); mxGPUDestroyGPUArray(centZ); mxGPUDestroyGPUArray(centX); - mxGPUDestroyGPUArray(rxApod); mxGPUDestroyGPUArray(elemFst); mxGPUDestroyGPUArray(elemLst); mxGPUDestroyGPUArray(rxElemOrig); From a52383c3869c794456ec74caa604fc01da40a302 Mon Sep 17 00:00:00 2001 From: PiotrKarwat Date: Mon, 22 Nov 2021 10:43:45 +0100 Subject: [PATCH 3/6] Us4R & Reconstruction - include rxApod --- api/matlab/arrus/Reconstruction.m | 1 + api/matlab/arrus/Us4R.m | 4 ++++ 2 files changed, 5 insertions(+) diff --git a/api/matlab/arrus/Reconstruction.m b/api/matlab/arrus/Reconstruction.m index b68f30d31..27dde8246 100644 --- a/api/matlab/arrus/Reconstruction.m +++ b/api/matlab/arrus/Reconstruction.m @@ -35,6 +35,7 @@ decimation xGrid zGrid + rxApod bmodeEnable = true colorEnable = false vectorEnable = false diff --git a/api/matlab/arrus/Us4R.m b/api/matlab/arrus/Us4R.m index cd30649f0..7a985377a 100644 --- a/api/matlab/arrus/Us4R.m +++ b/api/matlab/arrus/Us4R.m @@ -203,6 +203,7 @@ function upload(obj, sequenceOperation, reconstructOperation) 'decimation', reconstructOperation.decimation, ... 'xGrid', reconstructOperation.xGrid, ... 'zGrid', reconstructOperation.zGrid, ... + 'rxApod', reconstructOperation.rxApod, ... 'bmodeEnable', reconstructOperation.bmodeEnable, ... 'colorEnable', reconstructOperation.colorEnable, ... 'vectorEnable', reconstructOperation.vectorEnable, ... @@ -548,6 +549,7 @@ function setRecParams(obj,varargin) 'decimation', 'dec'; ... 'xGrid', 'xGrid'; ... 'zGrid', 'zGrid'; ... + 'rxApod', 'rxApod'; ... 'bmodeEnable', 'bmodeEnable'; ... 'colorEnable', 'colorEnable'; ... 'vectorEnable', 'vectorEnable'; ... @@ -599,6 +601,7 @@ function setRecParams(obj,varargin) obj.sys.tangElem = gpuArray(single(obj.sys.tangElem)); obj.rec.zGrid = gpuArray(single(obj.rec.zGrid)); obj.rec.xGrid = gpuArray(single(obj.rec.xGrid)); + obj.rec.rxApod = gpuArray(single(obj.rec.rxApod)); obj.seq.txFoc = gpuArray(single(obj.seq.txFoc)); obj.seq.txAngZX = gpuArray(single(obj.seq.txAngZX)); obj.seq.txApCentZ = gpuArray(single(obj.seq.txApCentZ)); @@ -1165,6 +1168,7 @@ function closeSequence(obj) obj.sys.tangElem, ... obj.rec.zGrid, ... obj.rec.xGrid, ... + obj.rec.rxApod, ... obj.seq.txFoc(selFrames), ... obj.seq.txAngZX(selFrames), ... obj.seq.txApCentZ(selFrames), ... From cbd1bf521b8ec3770c1f2595d1c4c7aa258103a2 Mon Sep 17 00:00:00 2001 From: PiotrKarwat Date: Mon, 22 Nov 2021 10:46:20 +0100 Subject: [PATCH 4/6] Control scripts - include rxApod --- api/matlab/examples/Us4R_ATL_control.m | 3 ++- api/matlab/examples/Us4R_ATL_control_batches.m | 3 ++- api/matlab/examples/Us4R_Olympus_control.m | 3 ++- api/matlab/examples/Us4R_Ultrasonix_control.m | 3 ++- api/matlab/examples/Us4R_Vermon_control.m | 3 ++- api/matlab/examples/Us4R_control.m | 3 ++- api/matlab/examples/Us4R_duplex.m | 1 + api/matlab/examples/Us4R_maxSequence.m | 3 ++- 8 files changed, 15 insertions(+), 7 deletions(-) diff --git a/api/matlab/examples/Us4R_ATL_control.m b/api/matlab/examples/Us4R_ATL_control.m index 0a50a2b30..56ff76356 100644 --- a/api/matlab/examples/Us4R_ATL_control.m +++ b/api/matlab/examples/Us4R_ATL_control.m @@ -59,7 +59,8 @@ 'cicOrder', 2, ... 'decimation', 4, ... 'xGrid', (-20:0.10:20)*1e-3, ... - 'zGrid', ( 0:0.10:100)*1e-3); + 'zGrid', ( 0:0.10:100)*1e-3, ... + 'rxApod', hamming(20).'); us.upload(seqPWI,rec); % us.upload(seqSTA,rec); diff --git a/api/matlab/examples/Us4R_ATL_control_batches.m b/api/matlab/examples/Us4R_ATL_control_batches.m index d4e5bdb1e..dad3ebbf1 100644 --- a/api/matlab/examples/Us4R_ATL_control_batches.m +++ b/api/matlab/examples/Us4R_ATL_control_batches.m @@ -92,7 +92,8 @@ 'cicOrder', 1, ... 'decimation', 1, ... 'xGrid', xGrid, ... - 'zGrid', zGrid); + 'zGrid', zGrid, ... + 'rxApod', hamming(20).'); end %% Prepare the acquisition (& reconstruction) diff --git a/api/matlab/examples/Us4R_Olympus_control.m b/api/matlab/examples/Us4R_Olympus_control.m index f14112708..8559deeb2 100644 --- a/api/matlab/examples/Us4R_Olympus_control.m +++ b/api/matlab/examples/Us4R_Olympus_control.m @@ -70,7 +70,8 @@ 'cicOrder', 2, ... 'decimation', 4, ... 'xGrid', (-40:0.20:40)*1e-3, ... - 'zGrid', ( 0:0.20:110)*1e-3); + 'zGrid', ( 0:0.20:110)*1e-3, ... + 'rxApod', hamming(20).'); % us.upload(seqSTA, rec); us.upload(seqPWI, rec); diff --git a/api/matlab/examples/Us4R_Ultrasonix_control.m b/api/matlab/examples/Us4R_Ultrasonix_control.m index 7ac4297fc..291e97dda 100644 --- a/api/matlab/examples/Us4R_Ultrasonix_control.m +++ b/api/matlab/examples/Us4R_Ultrasonix_control.m @@ -59,7 +59,8 @@ 'cicOrder', 2, ... 'decimation', 4, ... 'xGrid', (-20:0.10:20)*1e-3, ... - 'zGrid', ( 0:0.10:50)*1e-3); + 'zGrid', ( 0:0.10:50)*1e-3, ... + 'rxApod', hamming(20).'); us.upload(seqPWI,rec); % us.upload(seqSTA,rec); diff --git a/api/matlab/examples/Us4R_Vermon_control.m b/api/matlab/examples/Us4R_Vermon_control.m index 20adb4d79..3abed25d8 100644 --- a/api/matlab/examples/Us4R_Vermon_control.m +++ b/api/matlab/examples/Us4R_Vermon_control.m @@ -71,7 +71,8 @@ 'cicOrder', 2, ... 'decimation', 1, ... 'xGrid', (-5:0.05:5)*1e-3, ... - 'zGrid', ( 0:0.05:10)*1e-3); + 'zGrid', ( 0:0.05:10)*1e-3, ... + 'rxApod', hamming(20).'); % us.upload(seqSTA, rec); us.upload(seqPWI, rec); diff --git a/api/matlab/examples/Us4R_control.m b/api/matlab/examples/Us4R_control.m index 5f16f79e1..744acc160 100644 --- a/api/matlab/examples/Us4R_control.m +++ b/api/matlab/examples/Us4R_control.m @@ -72,7 +72,8 @@ 'cicOrder', 2, ... 'decimation', 4, ... 'xGrid', (-20:0.10:20)*1e-3, ... - 'zGrid', ( 0:0.10:50)*1e-3); + 'zGrid', ( 0:0.10:50)*1e-3, ... + 'rxApod', hamming(20).'); % us.upload(seqSTA, rec); us.upload(seqPWI, rec); diff --git a/api/matlab/examples/Us4R_duplex.m b/api/matlab/examples/Us4R_duplex.m index 1cd61d31c..f30b2ded1 100644 --- a/api/matlab/examples/Us4R_duplex.m +++ b/api/matlab/examples/Us4R_duplex.m @@ -74,6 +74,7 @@ 'decimation', 2, ... 'xGrid', (-7:0.025:7)*1e-3, ... 'zGrid', ( 0:0.025:7)*1e-3, ... + 'rxApod', hamming(20).', ... 'colorEnable', true, ... 'vectorEnable', false, ... 'bmodeFrames', 1:9, ... diff --git a/api/matlab/examples/Us4R_maxSequence.m b/api/matlab/examples/Us4R_maxSequence.m index d8db4dafe..3b95b9e34 100644 --- a/api/matlab/examples/Us4R_maxSequence.m +++ b/api/matlab/examples/Us4R_maxSequence.m @@ -70,7 +70,8 @@ 'cicOrder', 2, ... 'decimation', 4, ... 'xGrid', (-20:0.10:20)*1e-3, ... - 'zGrid', ( 0:0.10:50)*1e-3); + 'zGrid', ( 0:0.10:50)*1e-3, ... + 'rxApod', hamming(20).'); us.upload(seqSTA,rec); % us.upload(seqPWI,rec); From 824c1e86e7fc9ef72ddd38d32976562880e6d550 Mon Sep 17 00:00:00 2001 From: PiotrKarwat Date: Tue, 23 Nov 2021 10:28:29 +0100 Subject: [PATCH 5/6] iqRaw2Lri - checkData - validation of int32 arrays added --- api/matlab/arrus/mexcuda/iqRaw2Lri.cu | 55 ++++++++++++++++----------- 1 file changed, 32 insertions(+), 23 deletions(-) diff --git a/api/matlab/arrus/mexcuda/iqRaw2Lri.cu b/api/matlab/arrus/mexcuda/iqRaw2Lri.cu index 01304a7b4..9b1b709a5 100644 --- a/api/matlab/arrus/mexcuda/iqRaw2Lri.cu +++ b/api/matlab/arrus/mexcuda/iqRaw2Lri.cu @@ -130,27 +130,36 @@ __global__ void iqRaw2Lri( float2 * iqLri, float const * zPix, float const * xP } } -__host__ void checkData(mxGPUArray const * const data, char const * const name, bool const isComplex, int const nDims, char const * const invalidInputMsgId) +__host__ void checkData(mxGPUArray const * const data, + char const * const name, + bool const mustBeInt, + bool const mustBeComplex, + int const mustBeNDim, + char const * const invalidInputMsgId) { std::string invalidInputMsgTxt(name); - if (mxGPUGetClassID(data) != mxSINGLE_CLASS) + if (mustBeInt && mxGPUGetClassID(data) != mxINT32_CLASS) + invalidInputMsgTxt += std::string(" must be int32."); + + else if (!mustBeInt && mxGPUGetClassID(data) != mxSINGLE_CLASS) invalidInputMsgTxt += std::string(" must be single."); - else if (!isComplex && mxGPUGetComplexity(data)) + else if (!mustBeComplex && mxGPUGetComplexity(data)) invalidInputMsgTxt += std::string(" must be real."); - else if (isComplex && !mxGPUGetComplexity(data)) + else if (mustBeComplex && !mxGPUGetComplexity(data)) invalidInputMsgTxt += std::string(" must be complex."); - else if (nDims==1 && !( mxGPUGetNumberOfDimensions(data) == 1 || - (mxGPUGetNumberOfDimensions(data) == 2 && mxGPUGetDimensions(data)[0] == 1))) + else if (mustBeNDim==1 && !( mxGPUGetNumberOfDimensions(data) == 1 || + (mxGPUGetNumberOfDimensions(data) == 2 && + mxGPUGetDimensions(data)[0] == 1))) invalidInputMsgTxt += std::string(" must be at most 1D vector."); - else if (nDims==2 && !(mxGPUGetNumberOfDimensions(data) <= 2)) + else if (mustBeNDim==2 && !(mxGPUGetNumberOfDimensions(data) <= 2)) invalidInputMsgTxt += std::string(" must be at most 2D array."); - else if (nDims==3 && !(mxGPUGetNumberOfDimensions(data) <= 3)) + else if (mustBeNDim==3 && !(mxGPUGetNumberOfDimensions(data) <= 3)) invalidInputMsgTxt += std::string(" must be at most 3D array."); else @@ -264,21 +273,21 @@ void mexFunction(int nlhs, mxArray * plhs[], initDel = mxGetScalar(prhs[20]); /* Validate inputs */ - checkData(iqRaw, "iqRaw", true, 3, invalidInputMsgId); - checkData(zElem, "zElem", false, 1, invalidInputMsgId); - checkData(xElem, "xElem", false, 1, invalidInputMsgId); - checkData(tangElem, "tangElem", false, 1, invalidInputMsgId); - checkData(zPix, "zPix", false, 1, invalidInputMsgId); - checkData(xPix, "xPix", false, 1, invalidInputMsgId); - checkData(rxApod, "rxApod", false, 1, invalidInputMsgId); - checkData(foc, "foc", false, 1, invalidInputMsgId); - checkData(ang, "ang", false, 1, invalidInputMsgId); - checkData(centZ, "centZ", false, 1, invalidInputMsgId); - checkData(centX, "centX", false, 1, invalidInputMsgId); -// 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 + checkData(iqRaw, "iqRaw", false, true, 3, invalidInputMsgId); + checkData(zElem, "zElem", false, false, 1, invalidInputMsgId); + checkData(xElem, "xElem", false, false, 1, invalidInputMsgId); + checkData(tangElem, "tangElem", false, false, 1, invalidInputMsgId); + checkData(zPix, "zPix", false, false, 1, invalidInputMsgId); + checkData(xPix, "xPix", false, false, 1, invalidInputMsgId); + checkData(rxApod, "rxApod", false, false, 1, invalidInputMsgId); + checkData(foc, "foc", false, false, 1, invalidInputMsgId); + checkData(ang, "ang", false, false, 1, invalidInputMsgId); + checkData(centZ, "centZ", false, false, 1, invalidInputMsgId); + checkData(centX, "centX", false, false, 1, invalidInputMsgId); + checkData(elemFst, "elemFst", true, false, 1, invalidInputMsgId); + checkData(elemLst, "elemLst", true, false, 1, invalidInputMsgId); + checkData(rxElemOrig,"rxElemOrig",true, false, 1, invalidInputMsgId); + checkData(nSampOmit, "nSampOmit", true, false, 1, invalidInputMsgId); /* Get some additional information */ nSamp = mxGPUGetDimensions(iqRaw)[0]; From 05e43519ed4bf7aa3c918c13707f18365ded73ec Mon Sep 17 00:00:00 2001 From: PiotrKarwat Date: Tue, 23 Nov 2021 11:10:32 +0100 Subject: [PATCH 6/6] Reconstruction - default rxApod (uniform) --- api/matlab/arrus/Reconstruction.m | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/api/matlab/arrus/Reconstruction.m b/api/matlab/arrus/Reconstruction.m index 27dde8246..27faf0a58 100644 --- a/api/matlab/arrus/Reconstruction.m +++ b/api/matlab/arrus/Reconstruction.m @@ -35,7 +35,7 @@ decimation xGrid zGrid - rxApod + rxApod = [1 1] bmodeEnable = true colorEnable = false vectorEnable = false