Skip to content
Projects
Groups
Snippets
Help
Loading...
Help
Submit feedback
Contribute to GitLab
Sign in
Toggle navigation
T
tile_processor_gpu
Project
Project
Details
Activity
Releases
Cycle Analytics
Repository
Repository
Files
Commits
Branches
Tags
Contributors
Graph
Compare
Charts
Issues
0
Issues
0
List
Board
Labels
Milestones
Merge Requests
0
Merge Requests
0
CI / CD
CI / CD
Pipelines
Jobs
Schedules
Charts
Wiki
Wiki
Snippets
Snippets
Members
Members
Collapse sidebar
Close sidebar
Activity
Graph
Charts
Create a new issue
Jobs
Commits
Issue Boards
Open sidebar
Elphel
tile_processor_gpu
Commits
f72160f5
Commit
f72160f5
authored
Apr 20, 2020
by
Andrey Filippov
Browse files
Options
Browse Files
Download
Email Patches
Plain Diff
cleanup
parent
70ff9492
Changes
2
Hide whitespace changes
Inline
Side-by-side
Showing
2 changed files
with
47 additions
and
54 deletions
+47
-54
TileProcessor.cuh
src/TileProcessor.cuh
+23
-29
TileProcessor.h
src/TileProcessor.h
+24
-25
No files found.
src/TileProcessor.cuh
View file @
f72160f5
...
@@ -1218,42 +1218,36 @@ extern "C" __global__ void correlate2D_inner(
...
@@ -1218,42 +1218,36 @@ extern "C" __global__ void correlate2D_inner(
#define USE_CDP
#define USE_CDP
#ifdef USE_CDP
#ifdef USE_CDP
extern "C" __global__ void generate_RBGA(
extern "C" __global__ void generate_RBGA(
// Parameters to generate texture tasks
// Parameters to generate texture tasks
struct tp_task * gpu_tasks,
struct tp_task * gpu_tasks,
int num_tiles, // number of tiles in task list
int num_tiles, // number of tiles in task list
// declare arrays in device code?
// declare arrays in device code?
int * gpu_texture_indices,// packed tile + bits (now only (1 << 7)
int * gpu_texture_indices,// packed tile + bits (now only (1 << 7)
int * num_texture_tiles, // number of texture tiles to process (8 separate elements for accumulation)
int * num_texture_tiles, // number of texture tiles to process (8 separate elements for accumulation)
int * woi, // x,y,width,height of the woi
int * woi, // x,y,width,height of the woi
int width, // <= TILESX, use for faster processing of LWIR images (should be actual + 1)
int width, // <= TILESX, use for faster processing of LWIR images (should be actual + 1)
int height, // <= TILESY, use for faster processing of LWIR images
int height, // <= TILESY, use for faster processing of LWIR images
// Parameters for the texture generation
// Parameters for the texture generation
float ** gpu_clt, // [NUM_CAMS] ->[TILESY][TILESX][NUM_COLORS][DTT_SIZE*DTT_SIZE]
float ** gpu_clt, // [NUM_CAMS] ->[TILESY][TILESX][NUM_COLORS][DTT_SIZE*DTT_SIZE]
// TODO: use geometry_correction rXY !
// TODO: use geometry_correction rXY !
struct gc * gpu_geometry_correction,
struct gc * gpu_geometry_correction,
int colors, // number of colors (3/1)
int colors, // number of colors (3/1)
int is_lwir, // do not perform shot correction
int is_lwir, // do not perform shot correction
float params[5],
float params[5], // mitigating CUDA_ERROR_INVALID_PTX
/*
/*
float min_shot, // 10.0
float min_shot, // 10.0
float scale_shot, // 3.0
float scale_shot, // 3.0
float diff_sigma, // pixel value/pixel change
float diff_sigma, // pixel value/pixel change
float diff_threshold, // pixel value/pixel change
float diff_threshold, // pixel value/pixel change
float min_agree, // minimal number of channels to agree on a point (real number to work with fuzzy averages)
float min_agree, // minimal number of channels to agree on a point (real number to work with fuzzy averages)
*/
*/
float weights[3], // scale for R,B,G
float weights[3], // scale for R,B,G
int dust_remove, // Do not reduce average weight when only one image differs much from the average
int dust_remove, // Do not reduce average weight when only one image differs much from the average
int keep_weights, // return channel weights after A in RGBA (was removed)
int keep_weights, // return channel weights after A in RGBA (was removed)
const size_t texture_rbga_stride, // in floats
const size_t texture_rbga_stride, // in floats
float * gpu_texture_tiles) // (number of colors +1 + ?)*16*16 rgba texture tiles
float * gpu_texture_tiles) // (number of colors +1 + ?)*16*16 rgba texture tiles
// float aaaa)
// float aaaa)
// float * gpu_diff_rgb_combo) // diff[NUM_CAMS], R[NUM_CAMS], B[NUM_CAMS],G[NUM_CAMS]
// float * gpu_diff_rgb_combo) // diff[NUM_CAMS], R[NUM_CAMS], B[NUM_CAMS],G[NUM_CAMS]
/*
const size_t texture_rbga_stride, // in floats
float * gpu_texture_tiles, // (number of colors +1 + ?)*16*16 rgba texture tiles
float aaaa)
*/
{
{
float min_shot = params[0]; // 10.0
float min_shot = params[0]; // 10.0
float scale_shot = params[1]; // 3.0
float scale_shot = params[1]; // 3.0
...
...
src/TileProcessor.h
View file @
f72160f5
...
@@ -122,34 +122,33 @@ extern "C" __global__ void imclt_rbg(
...
@@ -122,34 +122,33 @@ extern "C" __global__ void imclt_rbg(
int
woi_theight
,
int
woi_theight
,
const
size_t
dstride
);
// in floats (pixels)
const
size_t
dstride
);
// in floats (pixels)
/*
extern
"C"
__global__
void
generate_RBGA
(
extern
"C"
__global__
void
generate_RBGA
(
// Parameters to generate texture tasks
// Parameters to generate texture tasks
struct tp_task * gpu_tasks,
struct
tp_task
*
gpu_tasks
,
int num_tiles, // number of tiles in task list
int
num_tiles
,
// number of tiles in task list
// declare arrays in device code?
// declare arrays in device code?
int * gpu_texture_indices,// packed tile + bits (now only (1 << 7)
int
*
gpu_texture_indices
,
// packed tile + bits (now only (1 << 7)
int * num_texture_tiles, // number of texture tiles to process (8 separate elements for accumulation)
int
*
num_texture_tiles
,
// number of texture tiles to process (8 separate elements for accumulation)
int * woi, // x,y,width,height of the woi
int
*
woi
,
// x,y,width,height of the woi
int width, // <= TILESX, use for faster processing of LWIR images (should be actual + 1)
int
width
,
// <= TILESX, use for faster processing of LWIR images (should be actual + 1)
int height, // <= TILESY, use for faster processing of LWIR images
int
height
,
// <= TILESY, use for faster processing of LWIR images
// Parameters for the texture generation
// Parameters for the texture generation
float ** gpu_clt, // [NUM_CAMS] ->[TILESY][TILESX][NUM_COLORS][DTT_SIZE*DTT_SIZE]
float
**
gpu_clt
,
// [NUM_CAMS] ->[TILESY][TILESX][NUM_COLORS][DTT_SIZE*DTT_SIZE]
// TODO: use geometry_correction rXY !
// TODO: use geometry_correction rXY !
struct gc * gpu_geometry_correction,
struct
gc
*
gpu_geometry_correction
,
// float * gpu_geometry_correction,
int
colors
,
// number of colors (3/1)
// float * gpu_port_offsets, // relative ports x,y offsets - just to scale differences, may be approximate
int
is_lwir
,
// do not perform shot correction
int colors, // number of colors (3/1)
float
params
[
5
],
// mitigating CUDA_ERROR_INVALID_PTX
int is_lwir, // do not perform shot correction
/*
float min_shot, // 10.0
float min_shot, // 10.0
float scale_shot, // 3.0
float scale_shot, // 3.0
float diff_sigma, // pixel value/pixel change
float diff_sigma, // pixel value/pixel change
float diff_threshold, // pixel value/pixel change
float diff_threshold, // pixel value/pixel change
float min_agree, // minimal number of channels to agree on a point (real number to work with fuzzy averages)
float min_agree, // minimal number of channels to agree on a point (real number to work with fuzzy averages)
float weights[3], // scale for R,B,G
*/
int dust_remove, // Do not reduce average weight when only one image differs much from the average
float
weights
[
3
],
// scale for R,B,G
int keep_weights, // return channel weights after A in RGBA (was removed)
int
dust_remove
,
// Do not reduce average weight when only one image differs much from the average
const size_t texture_rbga_stride, // in floats
int
keep_weights
,
// return channel weights after A in RGBA (was removed)
float * gpu_texture_tiles, // (number of colors +1 + ?)*16*16 rgba texture tile
s
const
size_t
texture_rbga_stride
,
// in float
s
float * gpu_diff_rgb_combo); // diff[NUM_CAMS], R[NUM_CAMS], B[NUM_CAMS],G[NUM_CAMS]
float
*
gpu_texture_tiles
);
// (number of colors +1 + ?)*16*16 rgba texture tiles
*/
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