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
|
#include <stdio.h>
|
edd4ab2d
Pavel Govyadinov
Cleaned up more a...
|
6
|
#include <stim/visualization/colormap.h>
|
42145f38
Pavel Govyadinov
Fixed the issues ...
|
7
|
#include <sstream>
|
edd4ab2d
Pavel Govyadinov
Cleaned up more a...
|
8
|
#include <stim/math/mathvec.h>
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
9
|
|
556c4e15
Pavel Govyadinov
Changed the handl...
|
10
11
|
///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...
|
12
|
typedef unsigned char uchar;
|
32c433c7
Pavel Govyadinov
recovered the maj...
|
13
|
texture<uchar, cudaTextureType2D, cudaReadModeElementType> texIn;
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
14
15
16
|
float *result;
float* v_dif;
cudaArray* srcArray;
|
4f5b240a
Pavel Govyadinov
minor change to i...
|
17
|
bool testing = false;
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
18
19
20
21
22
23
24
25
26
27
|
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...
|
28
29
30
|
///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...
|
31
32
33
34
35
36
37
|
float get_sum(float *diff)
{
cublasStatus_t ret;
cublasHandle_t handle;
ret = cublasCreate(&handle);
|
a9f956be
Pavel Govyadinov
Fixed the cost fu...
|
38
|
ret = cublasSetVector(20*10, sizeof(*diff), diff, 1, v_dif, 1);
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
39
|
float out;
|
a9f956be
Pavel Govyadinov
Fixed the cost fu...
|
40
|
ret = cublasSasum(handle, 20*10, v_dif, 1, &out);
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
41
|
cublasDestroy(handle);
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
42
43
44
|
return out;
}
|
556c4e15
Pavel Govyadinov
Changed the handl...
|
45
46
47
|
///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...
|
48
49
|
__device__ float Template(int x)
{
|
a9f956be
Pavel Govyadinov
Fixed the cost fu...
|
50
|
if(x < 20/6 || x > 20*5/6 || (x > 20*2/6 && x < 20*4/6)){
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
51
52
53
54
55
56
57
|
return 1.0;
}else{
return 0.0;
}
}
|
556c4e15
Pavel Govyadinov
Changed the handl...
|
58
59
60
|
///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...
|
61
62
63
64
65
66
|
__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...
|
67
|
int idx = y*20+x;
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
68
|
|
a9f956be
Pavel Govyadinov
Fixed the cost fu...
|
69
|
float valIn = tex2D(texIn, x, y)/255.0;
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
70
|
float valTemp = Template(x);
|
5eeaf941
Pavel Govyadinov
changer to the ba...
|
71
72
|
result[idx] = abs(valIn-valTemp);
//result[idx] = abs(valIn);
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
73
74
75
76
|
}
|
556c4e15
Pavel Govyadinov
Changed the handl...
|
77
78
79
80
|
///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...
|
81
|
void initArray(cudaGraphicsResource_t src, int DIM_Y)
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
|
{
//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...
|
98
|
cudaMalloc( (void**) &result, 20*DIM_Y*sizeof(float));
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
99
|
checkCUDAerrors("Memory Allocation Issue 1");
|
a9f956be
Pavel Govyadinov
Fixed the cost fu...
|
100
|
cudaMalloc((void **) &v_dif, 20*10*sizeof(float));
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
101
102
103
104
105
|
checkCUDAerrors("Memory Allocation Issue 2");
//HANDLE_ERROR(
// cudaBindTextureToArray(texIn, ptr, &channelDesc)
// );
}
|
556c4e15
Pavel Govyadinov
Changed the handl...
|
106
107
108
|
///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...
|
109
110
111
112
113
114
115
116
117
118
119
|
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...
|
120
121
122
|
HANDLE_ERROR(
cudaFree(v_dif)
);
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
123
|
}
|
556c4e15
Pavel Govyadinov
Changed the handl...
|
124
125
126
127
|
///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...
|
128
|
extern "C"
|
edd4ab2d
Pavel Govyadinov
Cleaned up more a...
|
129
|
stim::vec<int> get_cost(cudaGraphicsResource_t src, int DIM_Y)
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
130
|
{
|
5e7c7581
Pavel Govyadinov
Debugging build f...
|
131
|
float output[DIM_Y];
|
edd4ab2d
Pavel Govyadinov
Cleaned up more a...
|
132
|
stim::vec<int> ret(0, 0);
|
a9f956be
Pavel Govyadinov
Fixed the cost fu...
|
133
134
|
float mini = 10000000000000000.0;
int idx;
|
556c4e15
Pavel Govyadinov
Changed the handl...
|
135
|
stringstream name; //for debugging
|
edd4ab2d
Pavel Govyadinov
Cleaned up more a...
|
136
|
name << "Test.bmp";
|
5e7c7581
Pavel Govyadinov
Debugging build f...
|
137
138
|
initArray(src, DIM_Y*10);
dim3 grid(20, DIM_Y*10);
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
139
|
dim3 block(1, 1);
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
140
|
get_diff <<< grid, block >>> (result);
|
f31bf86d
Pavel Govyadinov
Added skeleton fu...
|
141
|
stim::gpu2image<float>(result, name.str(), 20,DIM_Y*10,0,1);
|
5e7c7581
Pavel Govyadinov
Debugging build f...
|
142
|
for (int i = 0; i < DIM_Y; i++){
|
a9f956be
Pavel Govyadinov
Fixed the cost fu...
|
143
|
output[i] = get_sum(result+(20*10*i));
|
42145f38
Pavel Govyadinov
Fixed the issues ...
|
144
|
if(output[i] <= mini){
|
a9f956be
Pavel Govyadinov
Fixed the cost fu...
|
145
146
|
mini = output[i];
idx = i;
|
385d2447
Pavel Govyadinov
Checkpoint: Conve...
|
147
|
}
|
42145f38
Pavel Govyadinov
Fixed the issues ...
|
148
|
}
|
edd4ab2d
Pavel Govyadinov
Cleaned up more a...
|
149
150
151
|
// name.clear();
// name << "sample_" << inter << "_" << idx << ".bmp";
|
42145f38
Pavel Govyadinov
Fixed the issues ...
|
152
|
output[idx] = get_sum(result+(20*10*idx));
|
edd4ab2d
Pavel Govyadinov
Cleaned up more a...
|
153
|
// stim::gpu2image<float>(v_dif, name.str(), 20,10,0,1);
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
154
|
cleanUP(src);
|
edd4ab2d
Pavel Govyadinov
Cleaned up more a...
|
155
156
|
ret[0] = idx; ret[1] = (int) output[idx];
return ret;
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
157
|
}
|