92989af0
 ming
更新解码器
|
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
|
#include "cuda_kernels.h"
typedef unsigned char uchar;
typedef unsigned int uint32;
typedef int int32;
namespace cuda_common
{
__global__ void kernel_bilinear(float *src_img, float *dst_img,
int src_width, int src_height, int dst_width, int dst_height)
{
const int x = blockIdx.x * blockDim.x + threadIdx.x;
const int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x < dst_width && y < dst_height)
{
float fx = (x + 0.5)*src_width / (float)dst_width - 0.5;
float fy = (y + 0.5)*src_height / (float)dst_height - 0.5;
int ax = floor(fx);
int ay = floor(fy);
if (ax < 0)
{
ax = 0;
}
else if (ax > src_width - 2)
{
ax = src_width - 2;
}
if (ay < 0){
ay = 0;
}
else if (ay > src_height - 2)
{
ay = src_height - 2;
}
int A = ax + ay*src_width;
int B = ax + ay*src_width + 1;
int C = ax + ay*src_width + src_width;
int D = ax + ay*src_width + src_width + 1;
float w1, w2, w3, w4;
w1 = fx - ax;
w2 = 1 - w1;
w3 = fy - ay;
w4 = 1 - w3;
float blue = src_img[A] * w2*w4 + src_img[B] * w1*w4 + src_img[C] * w2*w3 + src_img[D] * w1*w3;
float green = src_img[src_width * src_height + A] * w2*w4 + src_img[src_width * src_height + B] * w1*w4
+ src_img[src_width * src_height + C] * w2*w3 + src_img[src_width * src_height + D] * w1*w3;
float red = src_img[src_width * src_height * 2 + A] * w2*w4 + src_img[src_width * src_height * 2 + B] * w1*w4
+ src_img[src_width * src_height * 2 + C] * w2*w3 + src_img[src_width * src_height * 2 + D] * w1*w3;
dst_img[y * dst_width + x] = blue;
dst_img[dst_width * dst_height + y * dst_width + x] = green;
dst_img[dst_width * dst_height * 2 + y * dst_width + x] = red;
}
}
cudaError_t ResizeImage(float* d_srcRGB, int src_width, int src_height, float* d_dstRGB, int dst_width, int dst_height)
{
dim3 block(32, 16, 1);
dim3 grid((dst_width + (block.x - 1)) / block.x, (dst_height + (block.y - 1)) / block.y, 1);
kernel_bilinear << < grid, block >> >(d_srcRGB, d_dstRGB, src_width, src_height, dst_width, dst_height);
cudaError_t cudaStatus = cudaGetLastError();
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "kernel_bilinear launch failed: %s\n", cudaGetErrorString(cudaStatus));
return cudaStatus;
}
cudaStatus = cudaDeviceSynchronize();
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching kernel_bilinear!\n", cudaStatus);
return cudaStatus;
}
return cudaStatus;
}
}
|
