Skip to content
GitLab
Projects
Groups
Snippets
Help
Loading...
Help
What's new
10
Help
Support
Community forum
Keyboard shortcuts
?
Submit feedback
Contribute to GitLab
Sign in
Toggle navigation
Open sidebar
Kai Karius
Fitter
Commits
99d84a49
Commit
99d84a49
authored
Nov 25, 2020
by
karius
Browse files
Options
Browse Files
Download
Email Patches
Plain Diff
less bad test
parent
e059d41d
Changes
3
Hide whitespace changes
Inline
Side-by-side
Showing
3 changed files
with
39 additions
and
14 deletions
+39
-14
TransformationGrid.cu
TransformationGrid.cu
+37
-14
TransformationGrid.h
TransformationGrid.h
+2
-0
TransformationGrid.o
TransformationGrid.o
+0
-0
No files found.
TransformationGrid.cu
View file @
99d84a49
...
...
@@ -7,6 +7,24 @@
#include <TransformationGrid.h>
__device__
float
identity
(
float
value
)
{
return
value
;
}
__device__
float
entropy
(
float
value
)
{
return
value
;
}
__device__
pointFloatToFloat_t
p_identity
=
identity
;
//the host-side function pointer to your __device__ function
//in host code: copy the function pointers to their host equivalent
void
threeaxisrot
(
double
r11
,
double
r12
,
double
r21
,
double
r31
,
double
r32
,
float3
*
const
&
res
){
res
->
z
=
(
float
)
atan2
(
r31
,
r32
)
+
M_PI
;
res
->
y
=
(
float
)
asin
(
r21
)
+
M_PI_2
;
...
...
@@ -173,6 +191,9 @@ TransformationGrid::TransformationGrid(const std::pair<float3,float3> bounding_b
cudaMalloc
((
void
**
)
&
d_scores
,
h_num_transformations
[
0
]
*
sizeof
(
*
d_scores
));
cudaMalloc
((
void
**
)
&
d_tmp_rot_quadrature_scores
,
h_num_translations
[
0
]
*
sizeof
(
*
d_scores
));
CudaCheckError
();
cudaMemcpyFromSymbol
(
&
h_identity
,
p_identity
,
sizeof
(
pointFloatToFloat_t
));
CudaCheckError
();
}
__host__
void
TransformationGrid
::
quadrature_test_function
(
void
){
...
...
@@ -181,13 +202,14 @@ __host__ void TransformationGrid::quadrature_test_function(void){
//being 8*pi^2 for the euler parametrisation of SO(3)
float
rotation_normalization
=
2
/
M_PI
;
//normalization for 6D space
float
six_d_normalization
=
translation_normalization
*
rotation_normalization
;
//
float six_d_normalization
_inv = rot
ation_normalization;
//
float six_d_normalization = translation_normalization*rotation_normalization;
float
six_d_normalization
=
transl
ation_normalization
;
int
rot_index
;
int
trans_index
;
for
(
uint
i
=
0
;
i
<
h_num_transformations
[
0
];
i
++
){
rot_index
=
i
/
h_num_translations
[
0
];
h_scores
[
i
]
=
six_d_normalization
/
sinf
(
h_rotations_euler
[
rot_index
].
y
);
// h_scores[i] = six_d_normalization/sinf(h_rotations_euler[rot_index].y);
h_scores
[
i
]
=
six_d_normalization
;
}
cudaMemcpy
(
d_scores
,
h_scores
,
h_num_transformations
[
0
]
*
sizeof
(
*
h_scores
),
cudaMemcpyHostToDevice
);
CudaCheckError
();
...
...
@@ -392,7 +414,7 @@ void translation_quadrature(TransformationGrid grid, float * d_tmp_quadrature){
template
<
uint
blockSize
>
__global__
void
rotation_quadrature
(
TransformationGrid
grid
){
void
rotation_quadrature
(
TransformationGrid
grid
,
pointFloatToFloat_t
function
){
__shared__
float
reduce
[
blockSize
];
uint
num_translations
=
grid
.
d_num_translations
[
0
];
uint
num_rotations
=
grid
.
d_num_rotations
[
0
];
...
...
@@ -413,7 +435,8 @@ void rotation_quadrature(TransformationGrid grid){
// . .
// . .
// (r_{N_r-1},t0);(r_{N_r-1},t1);(r_{N_r-1},t2); ... ;(r_{N_r-1},t_{N_t-1});
reduce
[
threadIdx
.
x
]
+=
d_scores
[
t
*
num_translations
+
b
];
//t-Design! - weights implicit at the end
reduce
[
threadIdx
.
x
]
+=
function
(
d_scores
[
t
*
num_translations
+
b
]);
// if (b ==1 && threadIdx.x ==0){
// printf("%i,%i,%f\n",t,t*num_translations,d_scores[t*num_translations]);
// }
...
...
@@ -445,7 +468,7 @@ void TransformationGrid::quadrature(float* const &result){
// template <uint blockSize>
// __global__
// void rotation_quadrature(TransformationGrid grid){
rotation_quadrature
<
block_size
><<<
rot_quat_grid_dim
,
block_size
>>>
(
*
this
);
rotation_quadrature
<
block_size
><<<
rot_quat_grid_dim
,
block_size
>>>
(
*
this
,
h_identity
);
cudaDeviceSynchronize
();
uint
secondary_quadrature_grid_dim
=
h_num_translations
[
0
]
/
block_size
+
1
;
...
...
@@ -770,10 +793,10 @@ void TransformationGrid::shift_and_scale_chamfer_score(float * d_score, float *
cudaMalloc
((
void
**
)
&
d_sum
,
sizeof
(
*
d_sum
));
cudaMemcpy
(
d_sum
,
h_sum
,
sizeof
(
*
h_sum
),
cudaMemcpyHostToDevice
);
CudaCheckError
();
thrust
::
for_each
(
thrust
::
device_pointer_cast
(
&
d_normalized_score
[
0
]),
thrust
::
device_pointer_cast
(
&
d_normalized_score
[
num_scores
-
1
]),
thrust
::
placeholders
::
_1
/=
sum
);
thrust
::
for_each
(
thrust
::
device_pointer_cast
(
&
d_normalized_score
[
0
]),
thrust
::
device_pointer_cast
(
&
d_normalized_score
[
num_scores
-
1
]),
entropy
());
float
entr
=
-
1.0
*
thrust
::
reduce
(
thrust
::
device_pointer_cast
(
&
d_normalized_score
[
0
]),
thrust
::
device_pointer_cast
(
&
d_normalized_score
[
num_scores
-
1
]));
printf
(
"Entropy for envelope score: %f
\n
"
,
entr
);
//
thrust::for_each(thrust::device_pointer_cast(&d_normalized_score[0]),thrust::device_pointer_cast(&d_normalized_score[num_scores-1]), thrust::placeholders::_1 /= sum);
//
thrust::for_each(thrust::device_pointer_cast(&d_normalized_score[0]),thrust::device_pointer_cast(&d_normalized_score[num_scores-1]), entropy());
//
float entr = -1.0*thrust::reduce(thrust::device_pointer_cast(&d_normalized_score[0]),thrust::device_pointer_cast(&d_normalized_score[num_scores-1]));
//
printf("Entropy for envelope score: %f\n",entr);
}
...
...
@@ -792,10 +815,10 @@ void TransformationGrid::shift_and_scale_envelope_score(float * d_score, float *
cudaMalloc
((
void
**
)
&
d_sum
,
sizeof
(
*
d_sum
));
cudaMemcpy
(
d_sum
,
h_sum
,
sizeof
(
*
h_sum
),
cudaMemcpyHostToDevice
);
CudaCheckError
();
thrust
::
for_each
(
thrust
::
device_pointer_cast
(
&
d_normalized_score
[
0
]),
thrust
::
device_pointer_cast
(
&
d_normalized_score
[
num_scores
-
1
]),
thrust
::
placeholders
::
_1
/=
sum
);
thrust
::
for_each
(
thrust
::
device_pointer_cast
(
&
d_normalized_score
[
0
]),
thrust
::
device_pointer_cast
(
&
d_normalized_score
[
num_scores
-
1
]),
entropy
());
float
entr
=
-
1.0
*
thrust
::
reduce
(
thrust
::
device_pointer_cast
(
&
d_normalized_score
[
0
]),
thrust
::
device_pointer_cast
(
&
d_normalized_score
[
num_scores
-
1
]));
printf
(
"Entropy for envelope score: %f
\n
"
,
entr
);
//
thrust::for_each(thrust::device_pointer_cast(&d_normalized_score[0]),thrust::device_pointer_cast(&d_normalized_score[num_scores-1]), thrust::placeholders::_1 /= sum);
//
thrust::for_each(thrust::device_pointer_cast(&d_normalized_score[0]),thrust::device_pointer_cast(&d_normalized_score[num_scores-1]), entropy());
//
float entr = -1.0*thrust::reduce(thrust::device_pointer_cast(&d_normalized_score[0]),thrust::device_pointer_cast(&d_normalized_score[num_scores-1]));
//
printf("Entropy for envelope score: %f\n",entr);
}
TransformationGrid
::~
TransformationGrid
()
{
...
...
TransformationGrid.h
View file @
99d84a49
...
...
@@ -47,6 +47,7 @@
//#include "DelaunayChecker.h"
//
typedef
float
(
*
pointFloatToFloat_t
)(
float
);
class
TransformationGrid
{
public:
...
...
@@ -85,6 +86,7 @@ public:
float
*
d_scores
;
uint
*
h_num_transformations
;
uint
*
d_num_transformations
;
pointFloatToFloat_t
h_identity
;
__device__
int4
d_transformation_index_from_linear
(
const
uint
&
linear_index
);
__host__
int4
h_transformation_index_from_linear
(
const
uint
&
linear_index
);
__device__
void
d_transformation_to_memory
(
const
uint
&
i
,
float4
*
const
&
rotation
,
float4
*
const
&
translation
);
...
...
TransformationGrid.o
View file @
99d84a49
No preview for this file type
Write
Preview
Markdown
is supported
0%
Try again
or
attach a new file
.
Attach a file
Cancel
You are about to add
0
people
to the discussion. Proceed with caution.
Finish editing this message first!
Cancel
Please
register
or
sign in
to comment