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
|
|
5e7c7581
Pavel Govyadinov
Debugging build f...
|
9
10
11
|
//#define
//#define DIM_Y 10890
//#define DIM_X 20
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
12
|
typedef unsigned char uchar;
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
13
14
15
|
//surface<void, 2> texOut; ///// maybe just do a normal array instead of a surface.
//we may not need a surface at all.
//texture<float, cudaTextureType2D, cudaReadModeElementType> texTemplate
|
32c433c7
Pavel Govyadinov
recovered the maj...
|
16
|
texture<uchar, cudaTextureType2D, cudaReadModeElementType> texIn;
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
17
18
19
|
float *result;
float* v_dif;
cudaArray* srcArray;
|
4f5b240a
Pavel Govyadinov
minor change to i...
|
20
|
bool testing = false;
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
|
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);
}
}
float get_sum(float *diff)
{
cublasStatus_t ret;
cublasHandle_t handle;
ret = cublasCreate(&handle);
|
a9f956be
Pavel Govyadinov
Fixed the cost fu...
|
39
|
ret = cublasSetVector(20*10, sizeof(*diff), diff, 1, v_dif, 1);
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
40
|
float out;
|
a9f956be
Pavel Govyadinov
Fixed the cost fu...
|
41
|
ret = cublasSasum(handle, 20*10, v_dif, 1, &out);
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
42
|
cublasDestroy(handle);
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
43
44
45
46
47
|
return out;
}
__device__ float Template(int x)
{
|
a9f956be
Pavel Govyadinov
Fixed the cost fu...
|
48
|
if(x < 20/6 || x > 20*5/6 || (x > 20*2/6 && x < 20*4/6)){
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
49
50
51
52
53
54
55
56
57
58
59
60
61
|
return 1.0;
}else{
return 0.0;
}
}
__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...
|
62
|
int idx = y*20+x;
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
63
|
|
a9f956be
Pavel Govyadinov
Fixed the cost fu...
|
64
65
|
//float valIn = tex2D(texIn, x, y);
float valIn = tex2D(texIn, x, y)/255.0;
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
66
67
68
69
70
|
//int idx = x*DIM_X+y;
//uchar4 color = tex2D(texIn, x, y);
//float3 tempcolor = make_float3(color.x, color.y, color.z);
//float valIn = tempcolor.x + tempcolor.y + tempcolor.z;
|
2a18be6d
Pavel Govyadinov
New comments and ...
|
71
|
//float valIn = tex2D(texIn, x, y);
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
72
|
float valTemp = Template(x);
|
2a18be6d
Pavel Govyadinov
New comments and ...
|
73
|
// result[idx] = (float) x/(blockDim.x*gridDim.x);//abs(x);
|
5e7c7581
Pavel Govyadinov
Debugging build f...
|
74
|
result[idx] = abs(valIn);
|
0fdb4ed4
Pavel Govyadinov
fixed the issue w...
|
75
76
77
|
// #if __CUDA_ARCH__>=200
// printf("Value is : %f\n and the result is : %f\n", valIn, result[idx]);
// #endif
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
78
79
80
81
82
|
}
|
5e7c7581
Pavel Govyadinov
Debugging build f...
|
83
|
void initArray(cudaGraphicsResource_t src, int DIM_Y)
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
|
{
//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...
|
100
|
cudaMalloc( (void**) &result, 20*DIM_Y*sizeof(float));
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
101
|
checkCUDAerrors("Memory Allocation Issue 1");
|
a9f956be
Pavel Govyadinov
Fixed the cost fu...
|
102
|
cudaMalloc((void **) &v_dif, 20*10*sizeof(float));
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
|
checkCUDAerrors("Memory Allocation Issue 2");
//HANDLE_ERROR(
// cudaBindTextureToArray(texIn, ptr, &channelDesc)
// );
}
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
124
125
|
}
extern "C"
|
5e7c7581
Pavel Govyadinov
Debugging build f...
|
126
|
int get_cost(cudaGraphicsResource_t src, int inter, int DIM_Y)
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
127
|
{
|
5e7c7581
Pavel Govyadinov
Debugging build f...
|
128
|
float output[DIM_Y];
|
a9f956be
Pavel Govyadinov
Fixed the cost fu...
|
129
130
|
float mini = 10000000000000000.0;
int idx;
|
42145f38
Pavel Govyadinov
Fixed the issues ...
|
131
|
stringstream name;
|
5e7c7581
Pavel Govyadinov
Debugging build f...
|
132
133
|
initArray(src, DIM_Y*10);
dim3 grid(20, DIM_Y*10);
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
134
|
dim3 block(1, 1);
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
135
|
get_diff <<< grid, block >>> (result);
|
5e7c7581
Pavel Govyadinov
Debugging build f...
|
136
137
|
stim::gpu2image<float>(result, "test.bmp", 20,DIM_Y*10,0,1);
for (int i = 0; i < DIM_Y; i++){
|
a9f956be
Pavel Govyadinov
Fixed the cost fu...
|
138
|
output[i] = get_sum(result+(20*10*i));
|
42145f38
Pavel Govyadinov
Fixed the issues ...
|
139
|
if(output[i] <= mini){
|
a9f956be
Pavel Govyadinov
Fixed the cost fu...
|
140
141
|
mini = output[i];
idx = i;
|
385d2447
Pavel Govyadinov
Checkpoint: Conve...
|
142
|
}
|
42145f38
Pavel Govyadinov
Fixed the issues ...
|
143
144
145
146
|
}
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...
|
147
|
cleanUP(src);
|
a9f956be
Pavel Govyadinov
Fixed the cost fu...
|
148
|
return idx;
|
7e099e80
Pavel Govyadinov
"lots of stuff, t...
|
149
|
}
|