DAS重建程式碼（CUDA實現環陣、matlab實現線陣）
__global__
void
recKernel
（
GPUMemory
gpu_memory
，
RecoParam
reco_param
）
//重建
{
int
i
=
threadIdx
。
x
+
blockIdx
。
x
*
blockDim
。
x
；
int
j
=
threadIdx
。
y
+
blockIdx
。
y
*
blockDim
。
y
；
int
k
=
j
*
512
+
i
；
float
dx
=
0
；
float
dy
=
0
；
float
rr0
=
0
；
float
idx_temp
=
0
；
float
pa_temp
=
0
；
float
new_a
=
0
；
float
Vs
=
reco_param
。
voice_speed
；
float
Fr
=
reco_param
。
sample_fr
；
int
Tp
=
reco_param
。
time_point
；
float
R
=
reco_param
。
radius
=
0。055
；
float
pa
=
0
；
float
pa1
=
0
；
float
pa2
=
0
；
float
pa3
=
0
；
float
pa4
=
0
；
float
p
；
int
int_idx
=
0
；
float
angle_per_step
=
2
*
3。1415926
/
256
；
for
（
int
iStep
=
0
；
iStep
<
reco_param
。
channel
；
iStep
++
）
//迴圈通道，畫素迴圈靠執行緒
{
dx
=
（
k
%
512
+
1
-
512
/
2。0
）
*
reco_param
。
pixel_size
-
cos
（
-
iStep
*
angle_per_step
+
3。141592623
/
2
）
*
R
；
dy
=
（（
512
-
k
/
512
）
-
512
/
2。0
）
*
reco_param
。
pixel_size
-
sin
（
-
iStep
*
angle_per_step
+
3。141592623
/
2
）
*
R
；
rr0
=
sqrt
（
dx
*
dx
+
dy
*
dy
）；
idx_temp
=
rr0
/
Vs
*
Fr
+
74
；
if
（
idx_temp
>
（
Tp
-
1
-
4
））
{
idx_temp
=
Tp
-
1
-
4
；
}
else
if
（
idx_temp
<
1
）
{
idx_temp
=
1
；
}
int_idx
=
int
（
idx_temp
）；
pa1
=
gpu_memory
。
Rawdata
［（
int_idx
）
*
reco_param
。
channel
+
iStep
-
1
］；
pa2
=
gpu_memory
。
Rawdata
［（
int_idx
）
*
reco_param
。
channel
+
iStep
］；
pa3
=
gpu_memory
。
Rawdata
［（
int_idx
）
*
reco_param
。
channel
+
iStep
+
1
］；
p
=
（
idx_temp
-
（
int_idx
））
*
（
idx_temp
-
（
int_idx
+
1
））
*
pa1
/
2
+
（
idx_temp
-
（
int_idx
-
1
））
*
（
idx_temp
-
（
int_idx
+
1
））
*
pa2
/
（
-
1
）
+
（
idx_temp
-
（
int_idx
-
1
））
*
（
idx_temp
-
（
int_idx
））
*
pa3
/
2
；
pa
=
pa
+
p
；
}
gpu_memory
。
ImageArrayR
［
k
］
=
pa
；
new_a
=
exp
（（
reco_param
。
amp
-
50
）
*
0。2
）
*
0。1
；
pa_temp
=
pa
*
new_a
/
200
+
（
reco_param
。
dynamicRange
-
50
）
*
5
；
if
（
pa_temp
>
255
）
pa_temp
=
255
；
else
if
（
pa_temp
<
0
）
pa_temp
=
0
；
gpu_memory
。
ImageArray
［
k
］
=
uchar
（
pa_temp
）；
}
void
CDefaultBuild
：：
GpuDrawing
（
float
*
Rawdata
，
cv
：：
Mat
&
ImgR
，
cv
：：
Mat
&
Img
，
GPUMemory
gpu_memory
，
RecoParam
reco_param
）
{
cudaError_t
cudaStatus
；
dim3
blocks
（
512
/
16
，
512
/
16
）；
dim3
threads
（
16
，
16
）；
cudaStatus
=
cudaMemcpy
（
gpu_memory
。
Rawdata
，
Rawdata
，
reco_param
。
channel
*
reco_param
。
time_point
*
sizeof
（
float
），
cudaMemcpyHostToDevice
）；
//int i = 0；
recKernel
<<
<
blocks
，
threads
>>
>
（
gpu_memory
，
reco_param
）；
cudaStatus
=
cudaGetLastError
（）；
if
（
cudaStatus
！=
cudaSuccess
）
{
fprintf
（
stderr
，
“addKernel launch failed： %s
\n
”
，
cudaGetErrorString
（
cudaStatus
））；
}
cudaStatus
=
cudaDeviceSynchronize
（）；
if
（
cudaStatus
！=
cudaSuccess
）
{
std
：：
cout
<<
“ error”
<<
std
：：
endl
；
}
cudaMemcpy
（
ImgR
。
data
，
gpu_memory
。
ImageArrayR
，
512
*
512
*
sizeof
（
float
），
cudaMemcpyDeviceToHost
）；
cudaMemcpy
（
Img
。
data
，
gpu_memory
。
ImageArray
，
512
*
512
*
sizeof
（
uchar
），
cudaMemcpyDeviceToHost
）；
}；
void
CDefaultBuild
：：
InitGpu
（
RecoParam
reco_param
，
GPUMemory
&
gpu_memory
）
{
float
*
dev_raw_data
，
*
dev_pa_imgR
；
uchar
*
dev_pa_img
；
cudaMalloc
（（
void
**
）
&
dev_raw_data
，
reco_param
。
channel
*
reco_param
。
time_point
*
sizeof
（
float
））；
cudaMalloc
（（
void
**
）
&
dev_pa_img
，
512
*
512
*
sizeof
（
uchar
））；
cudaMalloc
（（
void
**
）
&
dev_pa_imgR
，
512
*
512
*
sizeof
（
float
））；
gpu_memory
。
ImageArray
=
dev_pa_img
，
gpu_memory
。
ImageArrayR
=
dev_pa_imgR
，
gpu_memory
。
Rawdata
=
dev_raw_data
；
std
：：
cout
<<
“init”
<<
std
：：
endl
；
//std：：cout << gpu_memory。ImageArrayR << std：：endl；
//std：：cout << “init” << std：：endl；
//CallBack（）；
}
CDefaultBuild
：：
CDefaultBuild
（）
{
// 由於RecoParam有預設引數，因此在建構函式中可以進行一次Init
// 執行void InitGpu（RecoParam reco_param， GPUMemory & gpu_memory， void（*CallBack）（void））；
// CallBack引數用不到取消掉
// RecoParam reco_param， GPUMemory & gpu_memory直接用成員變數
InitGpu
（
m_RecoParam
，
m_GPUMemory
）；
}
CDefaultBuild
：：~
CDefaultBuild
（）
{
// 解構函式中進行空間釋放
// 執行void FreeGpu（GPUMemory gpu_memory）；
// GPUMemory gpu_memory直接用成員變數
cudaFree
（
m_GPUMemory
。
ImageArray
）；
cudaFree
（
m_GPUMemory
。
ImageArrayR
）；
cudaFree
（
m_GPUMemory
。
Rawdata
）；
}
bool
CDefaultBuild
：：
Build
（
float
*
Rawdata
，
const
RecoParam
&
recoParam
，
cv
：：
Mat
&
ImgR
，
cv
：：
Mat
&
Img
）
{
//0 test
std
：：
cout
<<
“I am CDefaultBuild func Build！！！”
<<
std
：：
endl
；
ImgR
=
cv
：：
Mat
（
512
，
512
，
CV_8UC1
）；
Img
=
cv
：：
Mat
（
512
，
512
，
CV_8UC1
）；
//1 判斷引數是否有變化
if
（
m_RecoParam
！=
recoParam
）
{
// 如果引數變化了，再進一步判斷是否需要InitGpu
// m_RecoParam要賦予新的值
if
（
m_RecoParam
。
channel
！=
recoParam
。
channel
||
m_RecoParam
。
time_point
！=
recoParam
。
time_point
）
{
cudaFree
（
m_GPUMemory
。
ImageArray
）；
cudaFree
（
m_GPUMemory
。
ImageArrayR
）；
cudaFree
（
m_GPUMemory
。
Rawdata
）；
InitGpu
（
recoParam
，
m_GPUMemory
）；
}
m_RecoParam
=
recoParam
；
}
//2 影象構建
// void GpuDrawing（float * Rawdata， cv：：Mat& ImgR， cv：：Mat& Img）
// GPUMemory gpu_memory， RecoParam reco_param直接用成員變數
GpuDrawing
（
Rawdata
，
ImgR
，
Img
，
m_GPUMemory
，
m_RecoParam
）；
return
true
；
}
線陣：
image
=
zeros
（
600
，
600
）；
dx
=
0。1e-3
；
dy
=
0。1e-3
；
vs
=
1。5e3
；
dt
=
20e-9
；
fs
=
1
/
dt
；
pitch
=
0。3e-3
；
for
i = 1：1：600
for
j = 1：1：600
for
channel_i = 1：1：128
l
=
sqrt
（（
j
*
dy
-
0
）
^
2
+
（（
i
-
300
）
*
dx
-
（
channel_i
-
64
）
*
pitch
）
^
2
）；
time
=
floor
（
l
/
vs
*
fs
）
+
1
；
image
（
i
，
j
）
=
image
（
i
，
j
）
+
sensor_data
（
channel_i
，
time
）；
end
end
end
