7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
1
2
3
4
|
#include <assert.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <cublas_v2.h>
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
5
6
|
#include <stdio.h>
#include "../visualization/colormap.h"
|
42145f38
Pavel Govyadinov
Fixed the issues ...
|
7
|
#include <sstream>
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
8
|
|
556c4e15
Pavel Govyadinov
Changed the handl...
|
9
10
|
///Cost function that works with the gl-spider class to find index of the item with min-cost.
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
11
|
typedef unsigned char uchar;
|
32c433c7
Pavel Govyadinov
recovered the maj...
|
12
|
texture<uchar, cudaTextureType2D, cudaReadModeElementType> texIn;
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
13
14
15
|
float *result;
float* v_dif;
cudaArray* srcArray;
|
4f5b240a
Pavel Govyadinov
minor change to i...
|
16
|
bool testing = false;
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
17
18
19
20
21
22
23
24
25
26
|
inline void checkCUDAerrors(const char *msg)
{
cudaError_t err = cudaGetLastError();
if (cudaSuccess != err){
fprintf(stderr, "Cuda error: %s: %s.\n", msg, cudaGetErrorString(err) );
exit(1);
}
}
|
556c4e15
Pavel Govyadinov
Changed the handl...
|
27
28
29
|
///Finds the sum of all the pixes in a gives template element.
///Returns the abosolute value.
///@param *diff, a pointer to the memory block that holds the pixel-differences.
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
30
31
32
33
34
35
36
|
float get_sum(float *diff)
{
cublasStatus_t ret;
cublasHandle_t handle;
ret = cublasCreate(&handle);
|
a9f956be
Pavel Govyadinov
Fixed the cost fu...
|
37
|
ret = cublasSetVector(20*10, sizeof(*diff), diff, 1, v_dif, 1);
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
38
|
float out;
|
a9f956be
Pavel Govyadinov
Fixed the cost fu...
|
39
|
ret = cublasSasum(handle, 20*10, v_dif, 1, &out);
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
40
|
cublasDestroy(handle);
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
41
42
43
|
return out;
}
|
556c4e15
Pavel Govyadinov
Changed the handl...
|
44
45
46
|
///A virtual representation of a uniform template.
///Returns the value of the template pixel.
///@param x, location of a pixel.
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
47
48
|
__device__ float Template(int x)
{
|
a9f956be
Pavel Govyadinov
Fixed the cost fu...
|
49
|
if(x < 20/6 || x > 20*5/6 || (x > 20*2/6 && x < 20*4/6)){
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
50
51
52
53
54
55
56
|
return 1.0;
}else{
return 0.0;
}
}
|
556c4e15
Pavel Govyadinov
Changed the handl...
|
57
58
59
|
///Find the difference of the given set of samples and the template
///using cuda acceleration.
///@param *result, a pointer to the memory that stores the result.
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
60
61
62
63
64
65
|
__global__
void get_diff (float *result)
{
//cuPrintf("Hello");
int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
|
5e7c7581
Pavel Govyadinov
Debugging build f...
|
66
|
int idx = y*20+x;
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
67
|
|
a9f956be
Pavel Govyadinov
Fixed the cost fu...
|
68
|
float valIn = tex2D(texIn, x, y)/255.0;
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
69
|
float valTemp = Template(x);
|
5eeaf941
Pavel Govyadinov
changer to the ba...
|
70
71
|
result[idx] = abs(valIn-valTemp);
//result[idx] = abs(valIn);
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
72
73
74
75
|
}
|
556c4e15
Pavel Govyadinov
Changed the handl...
|
76
77
78
79
|
///Initialization function, allocates the memory and passes the necessary
///handles from OpenGL and Cuda.
///@param src, cudaGraphicsResource that handles the shared OpenGL/Cuda Texture
///@param DIM_Y, integer controlling how much memory to allocate.
|
5e7c7581
Pavel Govyadinov
Debugging build f...
|
80
|
void initArray(cudaGraphicsResource_t src, int DIM_Y)
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
|
{
//cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<uchar> ();
//cudaMallocArray(&result, &channelDesc, DIM_X, DIM_Y, 0);
//HANDLE_ERROR(
// cudaGraphicsGLRegisterImage(&src,
// fboID,
// GL_TEXTURE_2D,
HANDLE_ERROR(
cudaGraphicsMapResources(1, &src)
);
HANDLE_ERROR(
cudaGraphicsSubResourceGetMappedArray(&srcArray, src,0,0)
);
HANDLE_ERROR(
cudaBindTextureToArray(texIn, srcArray)
);
|
5e7c7581
Pavel Govyadinov
Debugging build f...
|
97
|
cudaMalloc( (void**) &result, 20*DIM_Y*sizeof(float));
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
98
|
checkCUDAerrors("Memory Allocation Issue 1");
|
a9f956be
Pavel Govyadinov
Fixed the cost fu...
|
99
|
cudaMalloc((void **) &v_dif, 20*10*sizeof(float));
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
100
101
102
103
104
|
checkCUDAerrors("Memory Allocation Issue 2");
//HANDLE_ERROR(
// cudaBindTextureToArray(texIn, ptr, &channelDesc)
// );
}
|
556c4e15
Pavel Govyadinov
Changed the handl...
|
105
106
107
|
///Deinit function that frees the memery used and releases the texture resource
///back to OpenGL.
///@param src, cudaGraphicsResource that handles the shared OpenGL/Cuda Texture
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
108
109
110
111
112
113
114
115
116
117
118
|
void cleanUP(cudaGraphicsResource_t src)
{
HANDLE_ERROR(
cudaUnbindTexture(texIn)
);
HANDLE_ERROR(
cudaFree(result)
);
HANDLE_ERROR(
cudaGraphicsUnmapResources(1,&src)
);
|
a9f956be
Pavel Govyadinov
Fixed the cost fu...
|
119
120
121
|
HANDLE_ERROR(
cudaFree(v_dif)
);
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
122
|
}
|
556c4e15
Pavel Govyadinov
Changed the handl...
|
123
124
125
126
|
///External access-point to the cuda function
///@param src, cudaGraphicsResource that handles the shared OpenGL/Cuda Texture
///@param DIM_Y, the number of samples in the template.
///@inter temporary paramenter that tracks the number of times cost.h was called.
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
127
|
extern "C"
|
5e7c7581
Pavel Govyadinov
Debugging build f...
|
128
|
int get_cost(cudaGraphicsResource_t src, int inter, int DIM_Y)
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
129
|
{
|
5e7c7581
Pavel Govyadinov
Debugging build f...
|
130
|
float output[DIM_Y];
|
a9f956be
Pavel Govyadinov
Fixed the cost fu...
|
131
132
|
float mini = 10000000000000000.0;
int idx;
|
556c4e15
Pavel Govyadinov
Changed the handl...
|
133
|
stringstream name; //for debugging
|
5e7c7581
Pavel Govyadinov
Debugging build f...
|
134
135
|
initArray(src, DIM_Y*10);
dim3 grid(20, DIM_Y*10);
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
136
|
dim3 block(1, 1);
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
137
|
get_diff <<< grid, block >>> (result);
|
f31bf86d
Pavel Govyadinov
Added skeleton fu...
|
138
139
|
name << "temp_diff_" << inter << ".bmp";
stim::gpu2image<float>(result, name.str(), 20,DIM_Y*10,0,1);
|
5e7c7581
Pavel Govyadinov
Debugging build f...
|
140
|
for (int i = 0; i < DIM_Y; i++){
|
a9f956be
Pavel Govyadinov
Fixed the cost fu...
|
141
|
output[i] = get_sum(result+(20*10*i));
|
42145f38
Pavel Govyadinov
Fixed the issues ...
|
142
|
if(output[i] <= mini){
|
a9f956be
Pavel Govyadinov
Fixed the cost fu...
|
143
144
|
mini = output[i];
idx = i;
|
385d2447
Pavel Govyadinov
Checkpoint: Conve...
|
145
|
}
|
42145f38
Pavel Govyadinov
Fixed the issues ...
|
146
|
}
|
f31bf86d
Pavel Govyadinov
Added skeleton fu...
|
147
|
name.clear();
|
42145f38
Pavel Govyadinov
Fixed the issues ...
|
148
149
150
|
name << "sample_" << inter << "_" << idx << ".bmp";
output[idx] = get_sum(result+(20*10*idx));
stim::gpu2image<float>(v_dif, name.str(), 20,10,0,1);
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
151
|
cleanUP(src);
|
a9f956be
Pavel Govyadinov
Fixed the cost fu...
|
152
|
return idx;
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
153
|
}
|