mirror/darknet
1
0
Fork 0
mirror of https://github.com/pjreddie/darknet.git synced 2023-08-10 21:13:14 +03:00
Code Issues Projects Releases Wiki Activity

Merge branch 'master' of pjreddie.com:jnet

Browse Source 
Conflicts:
	src/tests.c
This commit is contained in:
Joseph Redmon 2014-04-17 09:57:30 -07:00
commit b4b729a15e
31 changed files with 1877 additions and 725 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

13
Makefile
View File

@ -1,18 +1,21 @@
CC=gcc
COMMON=-Wall `pkg-config --cflags opencv`
UNAME = $(shell uname)
OPTS=-O3
ifeq ($(UNAME), Darwin)
COMMON += -isystem /usr/local/Cellar/opencv/2.4.6.1/include/opencv -isystem /usr/local/Cellar/opencv/2.4.6.1/include
COMMON+= -isystem /usr/local/Cellar/opencv/2.4.6.1/include/opencv -isystem /usr/local/Cellar/opencv/2.4.6.1/include
LDFLAGS= -framework OpenCL
else
COMMON += -march=native
OPTS+= -march=native -flto
LDFLAGS= -lOpenCL
endif
CFLAGS= $(COMMON) -Ofast -flto
CFLAGS= $(COMMON) $(OPTS)
#CFLAGS= $(COMMON) -O0 -g
LDFLAGS=`pkg-config --libs opencv` -lm
LDFLAGS+=`pkg-config --libs opencv` -lm
VPATH=./src/
EXEC=cnn
OBJ=network.o image.o tests.o connected_layer.o maxpool_layer.o activations.o list.o option_list.o parser.o utils.o data.o matrix.o softmax_layer.o mini_blas.o convolutional_layer.o
OBJ=network.o image.o tests.o connected_layer.o maxpool_layer.o activations.o list.o option_list.o parser.o utils.o data.o matrix.o softmax_layer.o mini_blas.o convolutional_layer.o opencl.o gpu_gemm.o cpu_gemm.o normalization_layer.o
all: $(EXEC)

BIN
dog.jpg
View File

Binary file not shown.
Side by Side

Before

Width:  |  Height:  |  Size: 22 KiB

19
src/connected_layer.c
View File

@ -7,16 +7,17 @@
#include <stdlib.h>
#include <string.h>
connected_layer *make_connected_layer(int inputs, int outputs, ACTIVATION activation)
connected_layer *make_connected_layer(int batch, int inputs, int outputs, ACTIVATION activation)
{
fprintf(stderr, "Connected Layer: %d inputs, %d outputs\n", inputs, outputs);
int i;
connected_layer *layer = calloc(1, sizeof(connected_layer));
layer->inputs = inputs;
layer->outputs = outputs;
layer->batch=batch;
layer->output = calloc(outputs, sizeof(float*));
layer->delta = calloc(outputs, sizeof(float*));
layer->output = calloc(batch*outputs, sizeof(float*));
layer->delta = calloc(batch*outputs, sizeof(float*));
layer->weight_updates = calloc(inputs*outputs, sizeof(float));
layer->weight_adapt = calloc(inputs*outputs, sizeof(float));
@ -78,14 +79,14 @@ void forward_connected_layer(connected_layer layer, float *input)
{
int i;
memcpy(layer.output, layer.biases, layer.outputs*sizeof(float));
int m = 1;
int m = layer.batch;
int k = layer.inputs;
int n = layer.outputs;
float *a = input;
float *b = layer.weights;
float *c = layer.output;
gemm(0,0,m,n,k,1,a,k,b,n,1,c,n);
for(i = 0; i < layer.outputs; ++i){
for(i = 0; i < layer.outputs*layer.batch; ++i){
layer.output[i] = activate(layer.output[i], layer.activation);
}
//for(i = 0; i < layer.outputs; ++i) if(i%(layer.outputs/10+1)==0) printf("%f, ", layer.output[i]); printf("\n");
@ -94,12 +95,12 @@ void forward_connected_layer(connected_layer layer, float *input)
void learn_connected_layer(connected_layer layer, float *input)
{
int i;
for(i = 0; i < layer.outputs; ++i){
for(i = 0; i < layer.outputs*layer.batch; ++i){
layer.delta[i] *= gradient(layer.output[i], layer.activation);
layer.bias_updates[i] += layer.delta[i];
layer.bias_updates[i%layer.batch] += layer.delta[i]/layer.batch;
}
int m = layer.inputs;
int k = 1;
int k = layer.batch;
int n = layer.outputs;
float *a = input;
float *b = layer.delta;
@ -113,7 +114,7 @@ void backward_connected_layer(connected_layer layer, float *input, float *delta)
int m = layer.inputs;
int k = layer.outputs;
int n = 1;
int n = layer.batch;
float *a = layer.weights;
float *b = layer.delta;

3
src/connected_layer.h
View File

@ -4,6 +4,7 @@
#include "activations.h"
typedef struct{
int batch;
int inputs;
int outputs;
float *weights;
@ -25,7 +26,7 @@ typedef struct{
} connected_layer;
connected_layer *make_connected_layer(int inputs, int outputs, ACTIVATION activation);
connected_layer *make_connected_layer(int batch, int inputs, int outputs, ACTIVATION activation);
void forward_connected_layer(connected_layer layer, float *input);
void backward_connected_layer(connected_layer layer, float *input, float *delta);

159
src/convolutional_layer.c
View File

@ -31,7 +31,7 @@ image get_convolutional_delta(convolutional_layer layer)
return float_to_image(h,w,c,layer.delta);
}
convolutional_layer *make_convolutional_layer(int h, int w, int c, int n, int size, int stride, ACTIVATION activation)
convolutional_layer *make_convolutional_layer(int batch, int h, int w, int c, int n, int size, int stride, ACTIVATION activation)
{
int i;
size = 2*(size/2)+1; //HA! And you thought you'd use an even sized filter...
@ -40,6 +40,7 @@ convolutional_layer *make_convolutional_layer(int h, int w, int c, int n, int si
layer->w = w;
layer->c = c;
layer->n = n;
layer->batch = batch;
layer->stride = stride;
layer->size = size;
@ -56,12 +57,12 @@ convolutional_layer *make_convolutional_layer(int h, int w, int c, int n, int si
//layer->biases[i] = rand_normal()*scale + scale;
layer->biases[i] = 0;
}
int out_h = (h-size)/stride + 1;
int out_w = (w-size)/stride + 1;
int out_h = convolutional_out_height(*layer);
int out_w = convolutional_out_width(*layer);
layer->col_image = calloc(out_h*out_w*size*size*c, sizeof(float));
layer->output = calloc(out_h * out_w * n, sizeof(float));
layer->delta = calloc(out_h * out_w * n, sizeof(float));
layer->col_image = calloc(layer->batch*out_h*out_w*size*size*c, sizeof(float));
layer->output = calloc(layer->batch*out_h * out_w * n, sizeof(float));
layer->delta = calloc(layer->batch*out_h * out_w * n, sizeof(float));
layer->activation = activation;
fprintf(stderr, "Convolutional Layer: %d x %d x %d image, %d filters -> %d x %d x %d image\n", h,w,c,n, out_h, out_w, n);
@ -70,21 +71,39 @@ convolutional_layer *make_convolutional_layer(int h, int w, int c, int n, int si
return layer;
}
void resize_convolutional_layer(convolutional_layer *layer, int h, int w, int c)
{
layer->h = h;
layer->w = w;
layer->c = c;
int out_h = convolutional_out_height(*layer);
int out_w = convolutional_out_width(*layer);
layer->col_image = realloc(layer->col_image,
layer->batch*out_h*out_w*layer->size*layer->size*layer->c*sizeof(float));
layer->output = realloc(layer->output,
layer->batch*out_h * out_w * layer->n*sizeof(float));
layer->delta = realloc(layer->delta,
layer->batch*out_h * out_w * layer->n*sizeof(float));
}
void forward_convolutional_layer(const convolutional_layer layer, float *in)
{
int i;
int m = layer.n;
int k = layer.size*layer.size*layer.c;
int n = ((layer.h-layer.size)/layer.stride + 1)*
((layer.w-layer.size)/layer.stride + 1);
int n = convolutional_out_height(layer)*
convolutional_out_width(layer)*
layer.batch;
memset(layer.output, 0, m*n*sizeof(float));
float *a = layer.filters;
float *b = layer.col_image;
float *c = layer.output;
im2col_cpu(in, layer.c, layer.h, layer.w, layer.size, layer.stride, b);
for(i = 0; i < layer.batch; ++i){
im2col_cpu(in+i*(n/layer.batch), layer.c, layer.h, layer.w, layer.size, layer.stride, b+i*(n/layer.batch));
}
gemm(0,0,m,n,k,1,a,k,b,n,1,c,n);
for(i = 0; i < m*n; ++i){
@ -97,9 +116,10 @@ void forward_convolutional_layer(const convolutional_layer layer, float *in)
void gradient_delta_convolutional_layer(convolutional_layer layer)
{
int i;
int size = convolutional_out_height(layer)
*convolutional_out_width(layer)
*layer.n;
int size = convolutional_out_height(layer)*
convolutional_out_width(layer)*
layer.n*
layer.batch;
for(i = 0; i < size; ++i){
layer.delta[i] *= gradient(layer.output[i], layer.activation);
}
@ -107,15 +127,17 @@ void gradient_delta_convolutional_layer(convolutional_layer layer)
void learn_bias_convolutional_layer(convolutional_layer layer)
{
int i,j;
int i,j,b;
int size = convolutional_out_height(layer)
*convolutional_out_width(layer);
for(i = 0; i < layer.n; ++i){
float sum = 0;
for(j = 0; j < size; ++j){
sum += layer.delta[j+i*size];
for(b = 0; b < layer.batch; ++b){
for(i = 0; i < layer.n; ++i){
float sum = 0;
for(j = 0; j < size; ++j){
sum += layer.delta[j+size*(i+b*layer.n)];
}
layer.bias_updates[i] += sum/size;
}
layer.bias_updates[i] += sum/size;
}
}
@ -125,8 +147,9 @@ void learn_convolutional_layer(convolutional_layer layer)
learn_bias_convolutional_layer(layer);
int m = layer.n;
int n = layer.size*layer.size*layer.c;
int k = ((layer.h-layer.size)/layer.stride + 1)*
((layer.w-layer.size)/layer.stride + 1);
int k = convolutional_out_height(layer)*
convolutional_out_width(layer)*
layer.batch;
float *a = layer.delta;
float *b = layer.col_image;
@ -137,10 +160,12 @@ void learn_convolutional_layer(convolutional_layer layer)
void backward_convolutional_layer(convolutional_layer layer, float *delta)
{
int i;
int m = layer.size*layer.size*layer.c;
int k = layer.n;
int n = ((layer.h-layer.size)/layer.stride + 1)*
((layer.w-layer.size)/layer.stride + 1);
int n = convolutional_out_height(layer)*
convolutional_out_width(layer)*
layer.batch;
float *a = layer.filters;
float *b = layer.delta;
@ -150,8 +175,10 @@ void backward_convolutional_layer(convolutional_layer layer, float *delta)
memset(c, 0, m*n*sizeof(float));
gemm(1,0,m,n,k,1,a,m,b,n,1,c,n);
memset(delta, 0, layer.h*layer.w*layer.c*sizeof(float));
col2im_cpu(c, layer.c, layer.h, layer.w, layer.size, layer.stride, delta);
memset(delta, 0, layer.batch*layer.h*layer.w*layer.c*sizeof(float));
for(i = 0; i < layer.batch; ++i){
col2im_cpu(c+i*n/layer.batch, layer.c, layer.h, layer.w, layer.size, layer.stride, delta+i*n/layer.batch);
}
}
void update_convolutional_layer(convolutional_layer layer, float step, float momentum, float decay)
@ -225,7 +252,7 @@ void update_convolutional_layer(convolutional_layer layer, float step, float mom
void test_convolutional_layer()
{
convolutional_layer l = *make_convolutional_layer(4,4,1,1,3,1,LINEAR);
convolutional_layer l = *make_convolutional_layer(1,4,4,1,1,3,1,LINEAR);
float input[] = {1,2,3,4,
5,6,7,8,
9,10,11,12,
@ -258,52 +285,48 @@ image get_convolutional_filter(convolutional_layer layer, int i)
return float_to_image(h,w,c,layer.filters+i*h*w*c);
}
void visualize_convolutional_layer(convolutional_layer layer, char *window)
image *weighted_sum_filters(convolutional_layer layer, image *prev_filters)
{
int color = 1;
int border = 1;
int h,w,c;
int size = layer.size;
h = size;
w = (size + border) * layer.n - border;
c = layer.c;
if(c != 3 || !color){
h = (h+border)*c - border;
c = 1;
image *filters = calloc(layer.n, sizeof(image));
int i,j,k,c;
if(!prev_filters){
for(i = 0; i < layer.n; ++i){
filters[i] = copy_image(get_convolutional_filter(layer, i));
}
}
image filters = make_image(h,w,c);
int i,j;
for(i = 0; i < layer.n; ++i){
int w_offset = i*(size+border);
image k = get_convolutional_filter(layer, i);
//printf("%f ** ", layer.biases[i]);
//print_image(k);
image copy = copy_image(k);
normalize_image(copy);
for(j = 0; j < k.c; ++j){
//set_pixel(copy,0,0,j,layer.biases[i]);
}
if(c == 3 && color){
embed_image(copy, filters, 0, w_offset);
}
else{
for(j = 0; j < k.c; ++j){
int h_offset = j*(size+border);
image layer = get_image_layer(k, j);
embed_image(layer, filters, h_offset, w_offset);
free_image(layer);
else{
image base = prev_filters[0];
for(i = 0; i < layer.n; ++i){
image filter = get_convolutional_filter(layer, i);
filters[i] = make_image(base.h, base.w, base.c);
for(j = 0; j < layer.size; ++j){
for(k = 0; k < layer.size; ++k){
for(c = 0; c < layer.c; ++c){
float weight = get_pixel(filter, j, k, c);
image prev_filter = copy_image(prev_filters[c]);
scale_image(prev_filter, weight);
add_into_image(prev_filter, filters[i], 0,0);
free_image(prev_filter);
}
}
}
}
free_image(copy);
}
image delta = get_convolutional_delta(layer);
image dc = collapse_image_layers(delta, 1);
char buff[256];
sprintf(buff, "%s: Delta", window);
show_image(dc, buff);
free_image(dc);
show_image(filters, window);
free_image(filters);
return filters;
}
image *visualize_convolutional_layer(convolutional_layer layer, char *window, image *prev_filters)
{
image *single_filters = weighted_sum_filters(layer, 0);
show_images(single_filters, layer.n, window);
image delta = get_convolutional_image(layer);
image dc = collapse_image_layers(delta, 1);
char buff[256];
sprintf(buff, "%s: Output", window);
show_image(dc, buff);
save_image(dc, buff);
free_image(dc);
return single_filters;
}

6
src/convolutional_layer.h
View File

@ -5,6 +5,7 @@
#include "activations.h"
typedef struct {
int batch;
int h,w,c;
int n;
int size;
@ -24,11 +25,12 @@ typedef struct {
ACTIVATION activation;
} convolutional_layer;
convolutional_layer *make_convolutional_layer(int h, int w, int c, int n, int size, int stride, ACTIVATION activation);
convolutional_layer *make_convolutional_layer(int batch, int h, int w, int c, int n, int size, int stride, ACTIVATION activation);
void resize_convolutional_layer(convolutional_layer *layer, int h, int w, int c);
void forward_convolutional_layer(const convolutional_layer layer, float *in);
void learn_convolutional_layer(convolutional_layer layer);
void update_convolutional_layer(convolutional_layer layer, float step, float momentum, float decay);
void visualize_convolutional_layer(convolutional_layer layer, char *window);
image *visualize_convolutional_layer(convolutional_layer layer, char *window, image *prev_filters);
void backward_convolutional_layer(convolutional_layer layer, float *delta);

86
src/cpu_gemm.c Normal file
View File

@ -0,0 +1,86 @@
#include "mini_blas.h"
void cpu_gemm_nn(int TA, int TB, int M, int N, int K, float ALPHA,
float *A, int lda,
float *B, int ldb,
float BETA,
float *C, int ldc)
{
int i,j,k;
for(i = 0; i < M; ++i){
for(k = 0; k < K; ++k){
register float A_PART = ALPHA*A[i*lda+k];
for(j = 0; j < N; ++j){
C[i*ldc+j] += A_PART*B[k*ldb+j];
}
}
}
}
void cpu_gemm_nt(int TA, int TB, int M, int N, int K, float ALPHA,
float *A, int lda,
float *B, int ldb,
float BETA,
float *C, int ldc)
{
int i,j,k;
for(i = 0; i < M; ++i){
for(j = 0; j < N; ++j){
register float sum = 0;
for(k = 0; k < K; ++k){
sum += ALPHA*A[i*lda+k]*B[k+j*ldb];
}
C[i*ldc+j] += sum;
}
}
}
void cpu_gemm_tn(int TA, int TB, int M, int N, int K, float ALPHA,
float *A, int lda,
float *B, int ldb,
float BETA,
float *C, int ldc)
{
int i,j,k;
for(i = 0; i < M; ++i){
for(k = 0; k < K; ++k){
register float A_PART = ALPHA*A[k*lda+i];
for(j = 0; j < N; ++j){
C[i*ldc+j] += A_PART*B[k*ldb+j];
}
}
}
}
void cpu_gemm_tt(int TA, int TB, int M, int N, int K, float ALPHA,
float *A, int lda,
float *B, int ldb,
float BETA,
float *C, int ldc)
{
int i,j,k;
for(i = 0; i < M; ++i){
for(j = 0; j < N; ++j){
for(k = 0; k < K; ++k){
C[i*ldc+j] += ALPHA*A[i+k*lda]*B[k+j*ldb];
}
}
}
}
void cpu_gemm(int TA, int TB, int M, int N, int K, float ALPHA,
float *A, int lda,
float *B, int ldb,
float BETA,
float *C, int ldc)
{
// Assume beta = 1 LULZ
if(!TA && !TB)
cpu_gemm_nn( TA, TB, M, N, K, ALPHA,A,lda, B, ldb,BETA,C,ldc);
else if(TA && !TB)
cpu_gemm_tn( TA, TB, M, N, K, ALPHA,A,lda, B, ldb,BETA,C,ldc);
else if(!TA && TB)
cpu_gemm_nt( TA, TB, M, N, K, ALPHA,A,lda, B, ldb,BETA,C,ldc);
else
cpu_gemm_tt( TA, TB, M, N, K, ALPHA,A,lda, B, ldb,BETA,C,ldc);
}

24
src/data.c
View File

@ -119,6 +119,30 @@ data load_categorical_data_csv(char *filename, int target, int k)
return d;
}
data load_cifar10_data(char *filename)
{
data d;
d.shallow = 0;
unsigned long i,j;
matrix X = make_matrix(10000, 3072);
matrix y = make_matrix(10000, 10);
d.X = X;
d.y = y;
FILE *fp = fopen(filename, "rb");
for(i = 0; i < 10000; ++i){
unsigned char bytes[3073];
fread(bytes, 1, 3073, fp);
int class = bytes[0];
y.vals[i][class] = 1;
for(j = 0; j < X.cols; ++j){
X.vals[i][j] = (double)bytes[j+1];
}
}
fclose(fp);
return d;
}
void randomize_data(data d)
{
int i;

1
src/data.h
View File

@ -17,6 +17,7 @@ data load_data_image_pathfile_part(char *filename, int part, int total,
char **labels, int k, int h, int w);
data load_data_image_pathfile_random(char *filename, int n, char **labels,
int k, int h, int w);
data load_cifar10_data(char *filename);
list *get_paths(char *filename);
data load_categorical_data_csv(char *filename, int target, int k);
void normalize_data_rows(data d);

72
src/gemm.cl Normal file
View File

@ -0,0 +1,72 @@
__kernel void gemm(int TA, int TB, int M, int N, int K, float ALPHA,
__global float *A, int lda,
__global float *B, int ldb,
float BETA,
__global float *C, int ldc)
{
__local float Asub[BLOCK][BLOCK];
__local float Bsub[BLOCK][BLOCK];
float val = 0;
int row_block = get_group_id(0);
int col_block = get_group_id(1);
int sub_row = get_local_id(0);
int sub_col = get_local_id(1);
int row = row_block*BLOCK + sub_row;
int col = col_block*BLOCK + sub_col;
int i,j;
for(i = 0; i < K; i += BLOCK){
int arow = row_block*BLOCK + sub_row;
int acol = i + sub_col;
int brow = i + sub_row;
int bcol = col_block*BLOCK + sub_col;
Asub[sub_row][sub_col] = TA ? A[arow + acol*lda] : A[arow*lda + acol];
Bsub[sub_row][sub_col] = TB ? B[brow + bcol*ldb] : B[brow*ldb + bcol];
barrier(CLK_LOCAL_MEM_FENCE);
for(j = 0; j < BLOCK && i+j<K; ++j){
val += Asub[sub_row][j]*Bsub[j][sub_col];
}
barrier(CLK_LOCAL_MEM_FENCE);
}
if(row < M && col < N){
C[row*ldc+col] = val;
}
}
/*
__kernel void gemm_slow(int TA, int TB, int M, int N, int K, float ALPHA,
__global float *A, int lda,
__global float *B, int ldb,
float BETA,
__global float *C, int ldc)
{
float val = 0;
int row = get_global_id(0);
int col = get_global_id(1);
int i;
for(i = 0; i < K; ++i){
float Aval;
if(TA) Aval = A[i*lda+row];
else Aval = A[row*lda+i];
float Bval;
if(TB) Bval = B[col*ldb+i];
else Bval = B[col+i*ldb];
val += Aval*Bval;
}
C[row*ldc+col] = val;
}
*/

153
src/gpu_gemm.c Normal file
View File

@ -0,0 +1,153 @@
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <time.h>
#include <math.h>
#include "opencl.h"
#include "mini_blas.h"
#define STR_HELPER(x) #x
#define STR(x) STR_HELPER(x)
#define BLOCK 8
cl_kernel get_gemm_kernel()
{
static int init = 0;
static cl_kernel gemm_kernel;
if(!init){
gemm_kernel = get_kernel("src/gemm.cl", "gemm", "-D BLOCK=" STR(BLOCK) );
init = 1;
}
return gemm_kernel;
}
void gpu_gemm(int TA, int TB, int M, int N, int K, float ALPHA,
float *A, int lda,
float *B, int ldb,
float BETA,
float *C, int ldc)
{
cl_setup();
cl_kernel gemm_kernel = get_gemm_kernel();
cl_context context = cl.context;
cl_command_queue queue = cl.queue;
size_t size = sizeof(float)*(TA ? lda*K:lda*M);
cl_mem A_gpu = clCreateBuffer(context,
CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,
size, A, &cl.error);
check_error(cl);
size = sizeof(float)*(TB ? ldb*N:ldb*K);
cl_mem B_gpu = clCreateBuffer(context,
CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,
size, B, &cl.error);
check_error(cl);
size = sizeof(float)*(ldc*M);
cl_mem C_gpu = clCreateBuffer(context,
CL_MEM_WRITE_ONLY|CL_MEM_COPY_HOST_PTR,
size, C, &cl.error);
check_error(cl);
cl_uint i = 0;
cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(TA), (void*) &TA);
cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(TB), (void*) &TB);
cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(M), (void*) &M);
cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(N), (void*) &N);
cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(K), (void*) &K);
cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(ALPHA), (void*) &ALPHA);
cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(A_gpu), (void*) &A_gpu);
cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(lda), (void*) &lda);
cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(B_gpu), (void*) &B_gpu);
cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(ldb), (void*) &ldb);
cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(BETA), (void*) &BETA);
cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(C_gpu), (void*) &C_gpu);
cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(ldc), (void*) &ldc);
check_error(cl);
const size_t global_size[] = {ceil((float)M/BLOCK)*BLOCK, ceil((float)N/BLOCK)*BLOCK};
const size_t local_size[] = {BLOCK, BLOCK};
//printf("%zd %zd %zd %zd\n", global_size[0], global_size[1], local_size[0], local_size[1]);
clEnqueueNDRangeKernel(queue, gemm_kernel, 2, 0, global_size, local_size, 0, 0, 0);
check_error(cl);
clEnqueueReadBuffer(queue, C_gpu, CL_TRUE, 0, size, C, 0, 0, 0);
check_error(cl);
clReleaseMemObject(A_gpu);
clReleaseMemObject(B_gpu);
clReleaseMemObject(C_gpu);
}
/*
cl_kernel get_gemm_kernel_slow()
{
static int init = 0;
static cl_kernel gemm_kernel;
if(!init){
gemm_kernel = get_kernel("src/gemm.cl", "gemm_slow");
init = 1;
}
return gemm_kernel;
}
void gpu_gemm_slow(int TA, int TB, int M, int N, int K, float ALPHA,
float *A, int lda,
float *B, int ldb,
float BETA,
float *C, int ldc)
{
cl_setup();
cl_kernel gemm_kernel = get_gemm_kernel_slow();
cl_context context = cl.context;
cl_command_queue queue = cl.queue;
size_t size = sizeof(float)*(TA ? lda*K:lda*M);
cl_mem A_gpu = clCreateBuffer(context,
CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,
size, A, &cl.error);
check_error(cl);
size = sizeof(float)*(TB ? ldb*N:ldb*K);
cl_mem B_gpu = clCreateBuffer(context,
CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,
size, B, &cl.error);
check_error(cl);
size = sizeof(float)*(ldc*M);
cl_mem C_gpu = clCreateBuffer(context,
CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,
size, C, &cl.error);
check_error(cl);
cl_uint i = 0;
cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(TA), (void*) &TA);
cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(TB), (void*) &TB);
cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(M), (void*) &M);
cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(N), (void*) &N);
cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(K), (void*) &K);
cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(ALPHA), (void*) &ALPHA);
cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(A_gpu), (void*) &A_gpu);
cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(lda), (void*) &lda);
cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(B_gpu), (void*) &B_gpu);
cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(ldb), (void*) &ldb);
cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(BETA), (void*) &BETA);
cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(C_gpu), (void*) &C_gpu);
cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(ldc), (void*) &ldc);
check_error(cl);
const size_t global_size[] = {M, N};
clEnqueueNDRangeKernel(queue, gemm_kernel, 2, 0, global_size, 0, 0, 0, 0);
clEnqueueReadBuffer(queue, C_gpu, CL_TRUE, 0, size, C, 0, 0, 0);
clReleaseMemObject(A_gpu);
clReleaseMemObject(B_gpu);
clReleaseMemObject(C_gpu);
}
*/

170
src/image.c
View File

@ -113,6 +113,7 @@ image copy_image(image p)
return copy;
}
void show_image(image p, char *name)
{
int i,j,k;
@ -136,7 +137,7 @@ void show_image(image p, char *name)
}
}
free_image(copy);
if(disp->height < 500 || disp->width < 500){
if(disp->height < 500 || disp->width < 500 || disp->height > 1000){
int w = 1500;
int h = w*p.h/p.w;
if(h > 1000){
@ -152,6 +153,30 @@ void show_image(image p, char *name)
cvReleaseImage(&disp);
}
void save_image(image p, char *name)
{
int i,j,k;
image copy = copy_image(p);
normalize_image(copy);
char buff[256];
//sprintf(buff, "%s (%d)", name, windows);
sprintf(buff, "%s.png", name);
IplImage *disp = cvCreateImage(cvSize(p.w,p.h), IPL_DEPTH_8U, p.c);
int step = disp->widthStep;
for(i = 0; i < p.h; ++i){
for(j = 0; j < p.w; ++j){
for(k= 0; k < p.c; ++k){
disp->imageData[i*step + j*p.c + k] = (unsigned char)(get_pixel(copy,i,j,k)*255);
}
}
}
free_image(copy);
cvSaveImage(buff, disp,0);
cvReleaseImage(&disp);
}
void show_image_layers(image p, char *name)
{
int i;
@ -227,7 +252,19 @@ image make_random_image(int h, int w, int c)
return out;
}
void add_scalar_image(image m, float s)
void add_into_image(image src, image dest, int h, int w)
{
int i,j,k;
for(k = 0; k < src.c; ++k){
for(i = 0; i < src.h; ++i){
for(j = 0; j < src.w; ++j){
add_pixel(dest, h+i, w+j, k, get_pixel(src, i, j, k));
}
}
}
}
void translate_image(image m, float s)
{
int i;
for(i = 0; i < m.h*m.w*m.c; ++i) m.data[i] += s;
@ -404,6 +441,20 @@ image get_image_layer(image m, int l)
}
return out;
}
image get_sub_image(image m, int h, int w, int dh, int dw)
{
image out = make_image(dh, dw, m.c);
int i,j,k;
for(k = 0; k < out.c; ++k){
for(i = 0; i < dh; ++i){
for(j = 0; j < dw; ++j){
float val = get_pixel(m, h+i, w+j, k);
set_pixel(out, i, j, k, val);
}
}
}
return out;
}
float get_pixel(image m, int x, int y, int c)
{
@ -594,6 +645,121 @@ void print_image(image m)
for(i =0 ; i < m.h*m.w*m.c; ++i) printf("%lf, ", m.data[i]);
printf("\n");
}
image collapse_images_vert(image *ims, int n)
{
int color = 1;
int border = 1;
int h,w,c;
w = ims[0].w;
h = (ims[0].h + border) * n - border;
c = ims[0].c;
if(c != 3 || !color){
w = (w+border)*c - border;
c = 1;
}
image filters = make_image(h,w,c);
int i,j;
for(i = 0; i < n; ++i){
int h_offset = i*(ims[0].h+border);
image copy = copy_image(ims[i]);
//normalize_image(copy);
if(c == 3 && color){
embed_image(copy, filters, h_offset, 0);
}
else{
for(j = 0; j < copy.c; ++j){
int w_offset = j*(ims[0].w+border);
image layer = get_image_layer(copy, j);
embed_image(layer, filters, h_offset, w_offset);
free_image(layer);
}
}
free_image(copy);
}
return filters;
}
image collapse_images_horz(image *ims, int n)
{
int color = 1;
int border = 1;
int h,w,c;
int size = ims[0].h;
h = size;
w = (ims[0].w + border) * n - border;
c = ims[0].c;
if(c != 3 || !color){
h = (h+border)*c - border;
c = 1;
}
image filters = make_image(h,w,c);
int i,j;
for(i = 0; i < n; ++i){
int w_offset = i*(size+border);
image copy = copy_image(ims[i]);
//normalize_image(copy);
if(c == 3 && color){
embed_image(copy, filters, 0, w_offset);
}
else{
for(j = 0; j < copy.c; ++j){
int h_offset = j*(size+border);
image layer = get_image_layer(copy, j);
embed_image(layer, filters, h_offset, w_offset);
free_image(layer);
}
}
free_image(copy);
}
return filters;
}
void show_images(image *ims, int n, char *window)
{
image m = collapse_images_vert(ims, n);
save_image(m, window);
show_image(m, window);
free_image(m);
}
image grid_images(image **ims, int h, int w)
{
int i;
image *rows = calloc(h, sizeof(image));
for(i = 0; i < h; ++i){
rows[i] = collapse_images_horz(ims[i], w);
}
image out = collapse_images_vert(rows, h);
for(i = 0; i < h; ++i){
free_image(rows[i]);
}
free(rows);
return out;
}
void test_grid()
{
int i,j;
int num = 3;
int topk = 3;
image **vizs = calloc(num, sizeof(image*));
for(i = 0; i < num; ++i){
vizs[i] = calloc(topk, sizeof(image));
for(j = 0; j < topk; ++j) vizs[i][j] = make_image(3,3,3);
}
image grid = grid_images(vizs, num, topk);
save_image(grid, "Test Grid");
free_image(grid);
}
void show_images_grid(image **ims, int h, int w, char *window)
{
image out = grid_images(ims, h, w);
show_image(out, window);
free_image(out);
}
void free_image(image m)
{

13
src/image.h
View File

@ -1,6 +1,7 @@
#ifndef IMAGE_H
#define IMAGE_H
#include "opencv2/highgui/highgui_c.h"
#include "opencv2/imgproc/imgproc_c.h"
typedef struct {
@ -12,7 +13,7 @@ typedef struct {
image image_distance(image a, image b);
void scale_image(image m, float s);
void add_scalar_image(image m, float s);
void translate_image(image m, float s);
void normalize_image(image p);
void z_normalize_image(image p);
void threshold_image(image p, float t);
@ -21,11 +22,20 @@ void rotate_image(image m);
void subtract_image(image a, image b);
float avg_image_layer(image m, int l);
void embed_image(image source, image dest, int h, int w);
void add_into_image(image src, image dest, int h, int w);
image collapse_image_layers(image source, int border);
image collapse_images_horz(image *ims, int n);
image collapse_images_vert(image *ims, int n);
image get_sub_image(image m, int h, int w, int dh, int dw);
void show_image(image p, char *name);
void save_image(image p, char *name);
void show_images(image *ims, int n, char *window);
void show_image_layers(image p, char *name);
void show_image_collapsed(image p, char *name);
void show_images_grid(image **ims, int h, int w, char *window);
void test_grid();
image grid_images(image **ims, int h, int w);
void print_image(image m);
image make_image(int h, int w, int c);
@ -39,6 +49,7 @@ image ipl_to_image(IplImage* src);
float get_pixel(image m, int x, int y, int c);
float get_pixel_extend(image m, int x, int y, int c);
void add_pixel(image m, int x, int y, int c, float val);
void set_pixel(image m, int x, int y, int c, float val);
image get_image_layer(image m, int l);

13
src/maxpool_layer.c
View File

@ -17,10 +17,12 @@ image get_maxpool_delta(maxpool_layer layer)
return float_to_image(h,w,c,layer.delta);
}
maxpool_layer *make_maxpool_layer(int h, int w, int c, int stride)
maxpool_layer *make_maxpool_layer(int batch, int h, int w, int c, int stride)
{
c = c*batch;
fprintf(stderr, "Maxpool Layer: %d x %d x %d image, %d stride\n", h,w,c,stride);
maxpool_layer *layer = calloc(1, sizeof(maxpool_layer));
layer->batch = batch;
layer->h = h;
layer->w = w;
layer->c = c;
@ -30,6 +32,15 @@ maxpool_layer *make_maxpool_layer(int h, int w, int c, int stride)
return layer;
}
void resize_maxpool_layer(maxpool_layer *layer, int h, int w, int c)
{
layer->h = h;
layer->w = w;
layer->c = c;
layer->output = realloc(layer->output, ((h-1)/layer->stride+1) * ((w-1)/layer->stride+1) * c * sizeof(float));
layer->delta = realloc(layer->delta, ((h-1)/layer->stride+1) * ((w-1)/layer->stride+1) * c * sizeof(float));
}
void forward_maxpool_layer(const maxpool_layer layer, float *in)
{
image input = float_to_image(layer.h, layer.w, layer.c, in);

4
src/maxpool_layer.h
View File

@ -4,6 +4,7 @@
#include "image.h"
typedef struct {
int batch;
int h,w,c;
int stride;
float *delta;
@ -11,7 +12,8 @@ typedef struct {
} maxpool_layer;
image get_maxpool_image(maxpool_layer layer);
maxpool_layer *make_maxpool_layer(int h, int w, int c, int stride);
maxpool_layer *make_maxpool_layer(int batch, int h, int w, int c, int stride);
void resize_maxpool_layer(maxpool_layer *layer, int h, int w, int c);
void forward_maxpool_layer(const maxpool_layer layer, float *in);
void backward_maxpool_layer(const maxpool_layer layer, float *in, float *delta);

148
src/mini_blas.c
View File

@ -3,6 +3,8 @@
#include <stdio.h>
#include <math.h>
#include <time.h>
#include <string.h>
#include "mini_blas.h"
void pm(int M, int N, float *A)
{
@ -17,42 +19,12 @@ void pm(int M, int N, float *A)
}
void gemm(int TA, int TB, int M, int N, int K, float ALPHA,
float *A, int lda,
float *B, int ldb,
float BETA,
float *C, int ldc)
float *A, int lda,
float *B, int ldb,
float BETA,
float *C, int ldc)
{
// Assume beta = 1 LULZ
int i,j,k;
if(TB && !TA){
for(i = 0; i < M; ++i){
for(j = 0; j < N; ++j){
register float sum = 0;
for(k = 0; k < K; ++k){
sum += ALPHA*A[i*lda+k]*B[k+j*ldb];
}
C[i*ldc+j] += sum;
}
}
}else if(TA && !TB){
for(i = 0; i < M; ++i){
for(k = 0; k < K; ++k){
register float A_PART = ALPHA*A[k*lda+i];
for(j = 0; j < N; ++j){
C[i*ldc+j] += A_PART*B[k*ldb+j];
}
}
}
}else{
for(i = 0; i < M; ++i){
for(k = 0; k < K; ++k){
register float A_PART = ALPHA*A[i*lda+k];
for(j = 0; j < N; ++j){
C[i*ldc+j] += A_PART*B[k*ldb+j];
}
}
}
}
gpu_gemm( TA, TB, M, N, K, ALPHA,A,lda, B, ldb,BETA,C,ldc);
}
void im2row(float *image, int h, int w, int c, int size, int stride, float *matrix)
@ -150,16 +122,26 @@ float *random_matrix(int rows, int cols)
void time_random_matrix(int TA, int TB, int m, int k, int n)
{
float *a = random_matrix(m,k);
float *b = random_matrix(k,n);
float *a;
if(!TA) a = random_matrix(m,k);
else a = random_matrix(k,m);
int lda = (!TA)?k:m;
float *b;
if(!TB) b = random_matrix(k,n);
else b = random_matrix(n,k);
int ldb = (!TB)?n:k;
float *c = random_matrix(m,n);
int i;
clock_t start = clock(), end;
for(i = 0; i<1000; ++i){
gemm(TA,TB,m,n,k,1,a,k,b,n,1,c,n);
cpu_gemm(TA,TB,m,n,k,1,a,lda,b,ldb,1,c,n);
}
end = clock();
printf("Matrix Multiplication %dx%d * %dx%d, TA=%d, TB=%d: %lf ms\n",m,k,k,n, TA, TB, (float)(end-start)/CLOCKS_PER_SEC);
free(a);
free(b);
free(c);
}
void test_blas()
@ -167,9 +149,97 @@ void test_blas()
time_random_matrix(0,0,100,100,100);
time_random_matrix(1,0,100,100,100);
time_random_matrix(0,1,100,100,100);
time_random_matrix(1,1,100,100,100);
time_random_matrix(0,1,1000,100,100);
time_random_matrix(0,0,1000,100,100);
time_random_matrix(1,0,1000,100,100);
time_random_matrix(0,1,1000,100,100);
time_random_matrix(1,1,1000,100,100);
}
void time_gpu_random_matrix(int TA, int TB, int m, int k, int n)
{
float *a;
if(!TA) a = random_matrix(m,k);
else a = random_matrix(k,m);
int lda = (!TA)?k:m;
float *b;
if(!TB) b = random_matrix(k,n);
else b = random_matrix(n,k);
int ldb = (!TB)?n:k;
float *c = random_matrix(m,n);
int i;
clock_t start = clock(), end;
for(i = 0; i<1000; ++i){
gpu_gemm(TA,TB,m,n,k,1,a,lda,b,ldb,1,c,n);
}
end = clock();
printf("Matrix Multiplication %dx%d * %dx%d, TA=%d, TB=%d: %lf ms\n",m,k,k,n, TA, TB, (float)(end-start)/CLOCKS_PER_SEC);
free(a);
free(b);
free(c);
}
void test_gpu_accuracy(int TA, int TB, int m, int k, int n)
{
srand(0);
float *a;
if(!TA) a = random_matrix(m,k);
else a = random_matrix(k,m);
int lda = (!TA)?k:m;
float *b;
if(!TB) b = random_matrix(k,n);
else b = random_matrix(n,k);
int ldb = (!TB)?n:k;
float *c = random_matrix(m,n);
float *c_gpu = random_matrix(m,n);
memset(c, 0, m*n*sizeof(float));
memset(c_gpu, 0, m*n*sizeof(float));
int i;
//pm(m,k,b);
gpu_gemm(TA,TB,m,n,k,1,a,lda,b,ldb,1,c_gpu,n);
//pm(m, n, c_gpu);
cpu_gemm(TA,TB,m,n,k,1,a,lda,b,ldb,1,c,n);
//pm(m, n, c);
double sse = 0;
for(i = 0; i < m*n; ++i) {
//printf("%f %f\n", c[i], c_gpu[i]);
sse += pow(c[i]-c_gpu[i], 2);
}
printf("Matrix Multiplication %dx%d * %dx%d, TA=%d, TB=%d: %g MSE\n",m,k,k,n, TA, TB, sse/(m*n));
free(a);
free(b);
free(c);
}
void test_gpu_blas()
{
test_gpu_accuracy(0,0,17,10,10);
test_gpu_accuracy(1,0,17,10,10);
test_gpu_accuracy(0,1,17,10,10);
test_gpu_accuracy(1,1,17,10,10);
test_gpu_accuracy(0,0,1000,10,100);
test_gpu_accuracy(1,0,1000,10,100);
test_gpu_accuracy(0,1,1000,10,100);
test_gpu_accuracy(1,1,1000,10,100);
time_gpu_random_matrix(0,0,1000,1000,100);
time_random_matrix(0,0,1000,1000,100);
time_gpu_random_matrix(0,1,1000,1000,100);
time_random_matrix(0,1,1000,1000,100);
time_gpu_random_matrix(1,0,1000,1000,100);
time_random_matrix(1,0,1000,1000,100);
time_gpu_random_matrix(1,1,1000,1000,100);
time_random_matrix(1,1,1000,1000,100);
}

13
src/mini_blas.h
View File

@ -4,6 +4,7 @@ void gemm(int TA, int TB, int M, int N, int K, float ALPHA,
float *B, int ldb,
float BETA,
float *C, int ldc);
float *random_matrix(int rows, int cols);
void im2row(float *image, int h, int w, int c, int size, int stride, float *matrix);
void im2col(float *image, int h, int w, int c, int size, int stride, float *matrix);
void im2col_cpu(float* data_im, const int channels,
@ -13,3 +14,15 @@ void col2im_cpu(float* data_col, const int channels,
const int height, const int width, const int ksize, const int stride,
float* data_im);
void test_blas();
void gpu_gemm(int TA, int TB, int M, int N, int K, float ALPHA,
float *A, int lda,
float *B, int ldb,
float BETA,
float *C, int ldc);
void cpu_gemm(int TA, int TB, int M, int N, int K, float ALPHA,
float *A, int lda,
float *B, int ldb,
float BETA,
float *C, int ldc);
void test_gpu_blas();

153
src/network.c
View File

@ -8,12 +8,14 @@
#include "convolutional_layer.h"
//#include "old_conv.h"
#include "maxpool_layer.h"
#include "normalization_layer.h"
#include "softmax_layer.h"
network make_network(int n)
network make_network(int n, int batch)
{
network net;
net.n = n;
net.batch = batch;
net.layers = calloc(net.n, sizeof(void *));
net.types = calloc(net.n, sizeof(LAYER_TYPE));
net.outputs = 0;
@ -25,10 +27,11 @@ void print_convolutional_cfg(FILE *fp, convolutional_layer *l, int first)
{
int i;
fprintf(fp, "[convolutional]\n");
if(first) fprintf(fp, "height=%d\n"
if(first) fprintf(fp, "batch=%d\n"
"height=%d\n"
"width=%d\n"
"channels=%d\n",
l->h, l->w, l->c);
l->batch,l->h, l->w, l->c);
fprintf(fp, "filters=%d\n"
"size=%d\n"
"stride=%d\n"
@ -38,17 +41,28 @@ void print_convolutional_cfg(FILE *fp, convolutional_layer *l, int first)
fprintf(fp, "data=");
for(i = 0; i < l->n; ++i) fprintf(fp, "%g,", l->biases[i]);
for(i = 0; i < l->n*l->c*l->size*l->size; ++i) fprintf(fp, "%g,", l->filters[i]);
/*
int j,k;
for(i = 0; i < l->n; ++i) fprintf(fp, "%g,", l->biases[i]);
for(i = 0; i < l->n; ++i){
for(j = l->c-1; j >= 0; --j){
for(k = 0; k < l->size*l->size; ++k){
fprintf(fp, "%g,", l->filters[i*(l->c*l->size*l->size)+j*l->size*l->size+k]);
}
}
}
*/
fprintf(fp, "\n\n");
}
void print_connected_cfg(FILE *fp, connected_layer *l, int first)
{
int i;
fprintf(fp, "[connected]\n");
if(first) fprintf(fp, "input=%d\n", l->inputs);
if(first) fprintf(fp, "batch=%d\ninput=%d\n", l->batch, l->inputs);
fprintf(fp, "output=%d\n"
"activation=%s\n",
l->outputs,
get_activation_string(l->activation));
"activation=%s\n",
l->outputs,
get_activation_string(l->activation));
fprintf(fp, "data=");
for(i = 0; i < l->outputs; ++i) fprintf(fp, "%g,", l->biases[i]);
for(i = 0; i < l->inputs*l->outputs; ++i) fprintf(fp, "%g,", l->weights[i]);
@ -58,17 +72,32 @@ void print_connected_cfg(FILE *fp, connected_layer *l, int first)
void print_maxpool_cfg(FILE *fp, maxpool_layer *l, int first)
{
fprintf(fp, "[maxpool]\n");
if(first) fprintf(fp, "height=%d\n"
"width=%d\n"
"channels=%d\n",
l->h, l->w, l->c);
if(first) fprintf(fp, "batch=%d\n"
"height=%d\n"
"width=%d\n"
"channels=%d\n",
l->batch,l->h, l->w, l->c);
fprintf(fp, "stride=%d\n\n", l->stride);
}
void print_normalization_cfg(FILE *fp, normalization_layer *l, int first)
{
fprintf(fp, "[localresponsenormalization]\n");
if(first) fprintf(fp, "batch=%d\n"
"height=%d\n"
"width=%d\n"
"channels=%d\n",
l->batch,l->h, l->w, l->c);
fprintf(fp, "size=%d\n"
"alpha=%g\n"
"beta=%g\n"
"kappa=%g\n\n", l->size, l->alpha, l->beta, l->kappa);
}
void print_softmax_cfg(FILE *fp, softmax_layer *l, int first)
{
fprintf(fp, "[softmax]\n");
if(first) fprintf(fp, "input=%d\n", l->inputs);
if(first) fprintf(fp, "batch=%d\ninput=%d\n", l->batch, l->inputs);
fprintf(fp, "\n");
}
@ -85,6 +114,8 @@ void save_network(network net, char *filename)
print_connected_cfg(fp, (connected_layer *)net.layers[i], i==0);
else if(net.types[i] == MAXPOOL)
print_maxpool_cfg(fp, (maxpool_layer *)net.layers[i], i==0);
else if(net.types[i] == NORMALIZATION)
print_normalization_cfg(fp, (normalization_layer *)net.layers[i], i==0);
else if(net.types[i] == SOFTMAX)
print_softmax_cfg(fp, (softmax_layer *)net.layers[i], i==0);
}
@ -115,6 +146,11 @@ void forward_network(network net, float *input)
forward_maxpool_layer(layer, input);
input = layer.output;
}
else if(net.types[i] == NORMALIZATION){
normalization_layer layer = *(normalization_layer *)net.layers[i];
forward_normalization_layer(layer, input);
input = layer.output;
}
}
}
@ -132,6 +168,9 @@ void update_network(network net, float step, float momentum, float decay)
else if(net.types[i] == SOFTMAX){
//maxpool_layer layer = *(maxpool_layer *)net.layers[i];
}
else if(net.types[i] == NORMALIZATION){
//maxpool_layer layer = *(maxpool_layer *)net.layers[i];
}
else if(net.types[i] == CONNECTED){
connected_layer layer = *(connected_layer *)net.layers[i];
update_connected_layer(layer, step, momentum, decay);
@ -153,6 +192,9 @@ float *get_network_output_layer(network net, int i)
} else if(net.types[i] == CONNECTED){
connected_layer layer = *(connected_layer *)net.layers[i];
return layer.output;
} else if(net.types[i] == NORMALIZATION){
normalization_layer layer = *(normalization_layer *)net.layers[i];
return layer.output;
}
return 0;
}
@ -191,11 +233,11 @@ float calculate_error_network(network net, float *truth)
float *out = get_network_output(net);
int i, k = get_network_output_size(net);
for(i = 0; i < k; ++i){
printf("%f, ", out[i]);
//printf("%f, ", out[i]);
delta[i] = truth[i] - out[i];
sum += delta[i]*delta[i];
}
printf("\n");
//printf("\n");
return sum;
}
@ -230,6 +272,10 @@ float backward_network(network net, float *input, float *truth)
maxpool_layer layer = *(maxpool_layer *)net.layers[i];
if(i != 0) backward_maxpool_layer(layer, prev_input, prev_delta);
}
else if(net.types[i] == NORMALIZATION){
normalization_layer layer = *(normalization_layer *)net.layers[i];
if(i != 0) backward_normalization_layer(layer, prev_input, prev_delta);
}
else if(net.types[i] == SOFTMAX){
softmax_layer layer = *(softmax_layer *)net.layers[i];
if(i != 0) backward_softmax_layer(layer, prev_input, prev_delta);
@ -258,19 +304,26 @@ float train_network_sgd(network net, data d, int n, float step, float momentum,f
int i;
float error = 0;
int correct = 0;
int pos = 0;
for(i = 0; i < n; ++i){
int index = rand()%d.X.rows;
error += train_network_datum(net, d.X.vals[index], d.y.vals[index], step, momentum, decay);
float err = train_network_datum(net, d.X.vals[index], d.y.vals[index], step, momentum, decay);
float *y = d.y.vals[index];
int class = get_predicted_class_network(net);
correct += (y[class]?1:0);
if(y[1]){
error += err;
++pos;
}
//printf("%d %f %f\n", i,net.output[0], d.y.vals[index][0]);
//if((i+1)%10 == 0){
// printf("%d: %f\n", (i+1), (float)correct/(i+1));
//}
}
printf("Accuracy: %f\n",(float) correct/n);
return error/n;
//printf("Accuracy: %f\n",(float) correct/n);
return error/pos;
}
float train_network_batch(network net, data d, int n, float step, float momentum,float decay)
{
@ -304,7 +357,7 @@ void train_network(network net, data d, float step, float momentum, float decay)
}
visualize_network(net);
cvWaitKey(100);
printf("Accuracy: %f\n", (float)correct/d.X.rows);
fprintf(stderr, "Accuracy: %f\n", (float)correct/d.X.rows);
}
int get_network_output_size_layer(network net, int i)
@ -330,29 +383,63 @@ int get_network_output_size_layer(network net, int i)
return 0;
}
int reset_network_size(network net, int h, int w, int c)
/*
int resize_network(network net, int h, int w, int c)
{
int i;
for (i = 0; i < net.n; ++i){
if(net.types[i] == CONVOLUTIONAL){
convolutional_layer *layer = (convolutional_layer *)net.layers[i];
layer->h = h;
layer->w = w;
layer->c = c;
image output = get_convolutional_image(*layer);
h = output.h;
w = output.w;
c = output.c;
}
else if(net.types[i] == MAXPOOL){
maxpool_layer *layer = (maxpool_layer *)net.layers[i];
layer->h = h;
layer->w = w;
layer->c = c;
image output = get_maxpool_image(*layer);
h = output.h;
w = output.w;
c = output.c;
}
}
return 0;
}
*/
int resize_network(network net, int h, int w, int c)
{
int i;
for (i = 0; i < net.n; ++i){
if(net.types[i] == CONVOLUTIONAL){
convolutional_layer *layer = (convolutional_layer *)net.layers[i];
layer->h = h;
layer->w = w;
layer->c = c;
resize_convolutional_layer(layer, h, w, c);
image output = get_convolutional_image(*layer);
h = output.h;
w = output.w;
c = output.c;
}
else if(net.types[i] == MAXPOOL){
}else if(net.types[i] == MAXPOOL){
maxpool_layer *layer = (maxpool_layer *)net.layers[i];
layer->h = h;
layer->w = w;
layer->c = c;
resize_maxpool_layer(layer, h, w, c);
image output = get_maxpool_image(*layer);
h = output.h;
w = output.w;
c = output.c;
}else if(net.types[i] == NORMALIZATION){
normalization_layer *layer = (normalization_layer *)net.layers[i];
resize_normalization_layer(layer, h, w, c);
image output = get_normalization_image(*layer);
h = output.h;
w = output.w;
c = output.c;
}else{
error("Cannot resize this type of layer");
}
}
return 0;
@ -374,6 +461,10 @@ image get_network_image_layer(network net, int i)
maxpool_layer layer = *(maxpool_layer *)net.layers[i];
return get_maxpool_image(layer);
}
else if(net.types[i] == NORMALIZATION){
normalization_layer layer = *(normalization_layer *)net.layers[i];
return get_normalization_image(layer);
}
return make_empty_image(0,0,0);
}
@ -389,13 +480,18 @@ image get_network_image(network net)
void visualize_network(network net)
{
image *prev = 0;
int i;
char buff[256];
for(i = 0; i < net.n; ++i){
sprintf(buff, "Layer %d", i);
if(net.types[i] == CONVOLUTIONAL){
convolutional_layer layer = *(convolutional_layer *)net.layers[i];
visualize_convolutional_layer(layer, buff);
prev = visualize_convolutional_layer(layer, buff, prev);
}
if(net.types[i] == NORMALIZATION){
normalization_layer layer = *(normalization_layer *)net.layers[i];
visualize_normalization_layer(layer, buff);
}
}
}
@ -467,3 +563,4 @@ float network_accuracy(network net, data d)
return acc;
}

8
src/network.h
View File

@ -9,18 +9,20 @@ typedef enum {
CONVOLUTIONAL,
CONNECTED,
MAXPOOL,
SOFTMAX
SOFTMAX,
NORMALIZATION
} LAYER_TYPE;
typedef struct {
int n;
int batch;
void **layers;
LAYER_TYPE *types;
int outputs;
float *output;
} network;
network make_network(int n);
network make_network(int n, int batch);
void forward_network(network net, float *input);
float backward_network(network net, float *input, float *truth);
void update_network(network net, float step, float momentum, float decay);
@ -41,7 +43,7 @@ int get_predicted_class_network(network net);
void print_network(network net);
void visualize_network(network net);
void save_network(network net, char *filename);
int reset_network_size(network net, int h, int w, int c);
int resize_network(network net, int h, int w, int c);
#endif

96
src/normalization_layer.c Normal file
View File

@ -0,0 +1,96 @@
#include "normalization_layer.h"
#include <stdio.h>
image get_normalization_image(normalization_layer layer)
{
int h = layer.h;
int w = layer.w;
int c = layer.c;
return float_to_image(h,w,c,layer.output);
}
image get_normalization_delta(normalization_layer layer)
{
int h = layer.h;
int w = layer.w;
int c = layer.c;
return float_to_image(h,w,c,layer.delta);
}
normalization_layer *make_normalization_layer(int batch, int h, int w, int c, int size, float alpha, float beta, float kappa)
{
fprintf(stderr, "Local Response Normalization Layer: %d x %d x %d image, %d size\n", h,w,c,size);
normalization_layer *layer = calloc(1, sizeof(normalization_layer));
layer->batch = batch;
layer->h = h;
layer->w = w;
layer->c = c;
layer->kappa = kappa;
layer->size = size;
layer->alpha = alpha;
layer->beta = beta;
layer->output = calloc(h * w * c * batch, sizeof(float));
layer->delta = calloc(h * w * c * batch, sizeof(float));
layer->sums = calloc(h*w, sizeof(float));
return layer;
}
void resize_normalization_layer(normalization_layer *layer, int h, int w, int c)
{
layer->h = h;
layer->w = w;
layer->c = c;
layer->output = realloc(layer->output, h * w * c * layer->batch * sizeof(float));
layer->delta = realloc(layer->delta, h * w * c * layer->batch * sizeof(float));
layer->sums = realloc(layer->sums, h*w * sizeof(float));
}
void add_square_array(float *src, float *dest, int n)
{
int i;
for(i = 0; i < n; ++i){
dest[i] += src[i]*src[i];
}
}
void sub_square_array(float *src, float *dest, int n)
{
int i;
for(i = 0; i < n; ++i){
dest[i] -= src[i]*src[i];
}
}
void forward_normalization_layer(const normalization_layer layer, float *in)
{
int i,j,k;
memset(layer.sums, 0, layer.h*layer.w*sizeof(float));
int imsize = layer.h*layer.w;
for(j = 0; j < layer.size/2; ++j){
if(j < layer.c) add_square_array(in+j*imsize, layer.sums, imsize);
}
for(k = 0; k < layer.c; ++k){
int next = k+layer.size/2;
int prev = k-layer.size/2-1;
if(next < layer.c) add_square_array(in+next*imsize, layer.sums, imsize);
if(prev > 0) sub_square_array(in+prev*imsize, layer.sums, imsize);
for(i = 0; i < imsize; ++i){
layer.output[k*imsize + i] = in[k*imsize+i] / pow(layer.kappa + layer.alpha * layer.sums[i], layer.beta);
}
}
}
void backward_normalization_layer(const normalization_layer layer, float *in, float *delta)
{
//TODO!
}
void visualize_normalization_layer(normalization_layer layer, char *window)
{
image delta = get_normalization_image(layer);
image dc = collapse_image_layers(delta, 1);
char buff[256];
sprintf(buff, "%s: Output", window);
show_image(dc, buff);
save_image(dc, buff);
free_image(dc);
}

26
src/normalization_layer.h Normal file
View File

@ -0,0 +1,26 @@
#ifndef NORMALIZATION_LAYER_H
#define NORMALIZATION_LAYER_H
#include "image.h"
typedef struct {
int batch;
int h,w,c;
int size;
float alpha;
float beta;
float kappa;
float *delta;
float *output;
float *sums;
} normalization_layer;
image get_normalization_image(normalization_layer layer);
normalization_layer *make_normalization_layer(int batch, int h, int w, int c, int size, float alpha, float beta, float kappa);
void resize_normalization_layer(normalization_layer *layer, int h, int w, int c);
void forward_normalization_layer(const normalization_layer layer, float *in);
void backward_normalization_layer(const normalization_layer layer, float *in, float *delta);
void visualize_normalization_layer(normalization_layer layer, char *window);
#endif

77
src/opencl.c Normal file
View File

@ -0,0 +1,77 @@
#include "opencl.h"
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
cl_info cl = {0};
void check_error(cl_info info)
{
if (info.error != CL_SUCCESS) {
printf("\n Error number %d", info.error);
}
}
cl_info cl_init()
{
cl_info info;
info.initialized = 0;
cl_uint platforms, devices;
// Fetch the Platform and Device IDs; we only want one.
info.error=clGetPlatformIDs(1, &info.platform, &platforms);
check_error(info);
info.error=clGetDeviceIDs(info.platform, CL_DEVICE_TYPE_ALL, 1, &info.device, &devices);
check_error(info);
cl_context_properties properties[]={
CL_CONTEXT_PLATFORM, (cl_context_properties)info.platform,
0};
// Note that nVidia's OpenCL requires the platform property
info.context=clCreateContext(properties, 1, &info.device, 0, 0, &info.error);
check_error(info);
info.queue = clCreateCommandQueue(info.context, info.device, 0, &info.error);
check_error(info);
info.initialized = 1;
return info;
}
cl_program cl_fprog(char *filename, char *options, cl_info info)
{
size_t srcsize;
char src[8192];
memset(src, 0, 8192);
FILE *fil=fopen(filename,"r");
srcsize=fread(src, sizeof src, 1, fil);
fclose(fil);
const char *srcptr[]={src};
// Submit the source code of the example kernel to OpenCL
cl_program prog=clCreateProgramWithSource(info.context,1, srcptr, &srcsize, &info.error);
check_error(info);
char build_c[4096];
// and compile it (after this we could extract the compiled version)
info.error=clBuildProgram(prog, 0, 0, options, 0, 0);
if ( info.error != CL_SUCCESS ) {
fprintf(stderr, "Error Building Program: %d\n", info.error);
clGetProgramBuildInfo( prog, info.device, CL_PROGRAM_BUILD_LOG, 4096, build_c, 0);
fprintf(stderr, "Build Log for %s program:\n%s\n", filename, build_c);
}
return prog;
}
void cl_setup()
{
if(!cl.initialized){
cl = cl_init();
}
}
cl_kernel get_kernel(char *filename, char *kernelname, char *options)
{
cl_setup();
cl_program prog = cl_fprog(filename, options, cl);
cl_kernel kernel=clCreateKernel(prog, kernelname, &cl.error);
check_error(cl);
return kernel;
}

21
src/opencl.h Normal file
View File

@ -0,0 +1,21 @@
#ifdef __APPLE__
#include <OpenCL/opencl.h>
#else
#include <CL/cl.h>
#endif
typedef struct {
int initialized;
cl_int error;
cl_platform_id platform;
cl_device_id device;
cl_context context;
cl_command_queue queue;
}cl_info;
extern cl_info cl;
void cl_setup();
void check_error(cl_info info);
cl_kernel get_kernel(char *filename, char *kernelname, char *options);

54
src/parser.c
View File

@ -7,6 +7,7 @@
#include "convolutional_layer.h"
#include "connected_layer.h"
#include "maxpool_layer.h"
#include "normalization_layer.h"
#include "softmax_layer.h"
#include "list.h"
#include "option_list.h"
@ -21,6 +22,7 @@ int is_convolutional(section *s);
int is_connected(section *s);
int is_maxpool(section *s);
int is_softmax(section *s);
int is_normalization(section *s);
list *read_cfg(char *filename);
void free_section(section *s)
@ -52,6 +54,7 @@ convolutional_layer *parse_convolutional(list *options, network net, int count)
h = option_find_int(options, "height",1);
w = option_find_int(options, "width",1);
c = option_find_int(options, "channels",1);
net.batch = option_find_int(options, "batch",1);
}else{
image m = get_network_image_layer(net, count-1);
h = m.h;
@ -59,7 +62,7 @@ convolutional_layer *parse_convolutional(list *options, network net, int count)
c = m.c;
if(h == 0) error("Layer before convolutional layer must output image.");
}
convolutional_layer *layer = make_convolutional_layer(h,w,c,n,size,stride, activation);
convolutional_layer *layer = make_convolutional_layer(net.batch,h,w,c,n,size,stride, activation);
char *data = option_find_str(options, "data", 0);
if(data){
char *curr = data;
@ -90,10 +93,11 @@ connected_layer *parse_connected(list *options, network net, int count)
ACTIVATION activation = get_activation(activation_s);
if(count == 0){
input = option_find_int(options, "input",1);
net.batch = option_find_int(options, "batch",1);
}else{
input = get_network_output_size_layer(net, count-1);
}
connected_layer *layer = make_connected_layer(input, output, activation);
connected_layer *layer = make_connected_layer(net.batch, input, output, activation);
char *data = option_find_str(options, "data", 0);
if(data){
char *curr = data;
@ -120,10 +124,11 @@ softmax_layer *parse_softmax(list *options, network net, int count)
int input;
if(count == 0){
input = option_find_int(options, "input",1);
net.batch = option_find_int(options, "batch",1);
}else{
input = get_network_output_size_layer(net, count-1);
}
softmax_layer *layer = make_softmax_layer(input);
softmax_layer *layer = make_softmax_layer(net.batch, input);
option_unused(options);
return layer;
}
@ -136,6 +141,7 @@ maxpool_layer *parse_maxpool(list *options, network net, int count)
h = option_find_int(options, "height",1);
w = option_find_int(options, "width",1);
c = option_find_int(options, "channels",1);
net.batch = option_find_int(options, "batch",1);
}else{
image m = get_network_image_layer(net, count-1);
h = m.h;
@ -143,7 +149,31 @@ maxpool_layer *parse_maxpool(list *options, network net, int count)
c = m.c;
if(h == 0) error("Layer before convolutional layer must output image.");
}
maxpool_layer *layer = make_maxpool_layer(h,w,c,stride);
maxpool_layer *layer = make_maxpool_layer(net.batch,h,w,c,stride);
option_unused(options);
return layer;
}
normalization_layer *parse_normalization(list *options, network net, int count)
{
int h,w,c;
int size = option_find_int(options, "size",1);
float alpha = option_find_float(options, "alpha", 0.);
float beta = option_find_float(options, "beta", 1.);
float kappa = option_find_float(options, "kappa", 1.);
if(count == 0){
h = option_find_int(options, "height",1);
w = option_find_int(options, "width",1);
c = option_find_int(options, "channels",1);
net.batch = option_find_int(options, "batch",1);
}else{
image m = get_network_image_layer(net, count-1);
h = m.h;
w = m.w;
c = m.c;
if(h == 0) error("Layer before convolutional layer must output image.");
}
normalization_layer *layer = make_normalization_layer(net.batch,h,w,c,size, alpha, beta, kappa);
option_unused(options);
return layer;
}
@ -151,7 +181,7 @@ maxpool_layer *parse_maxpool(list *options, network net, int count)
network parse_network_cfg(char *filename)
{
list *sections = read_cfg(filename);
network net = make_network(sections->size);
network net = make_network(sections->size, 0);
node *n = sections->front;
int count = 0;
@ -162,18 +192,27 @@ network parse_network_cfg(char *filename)
convolutional_layer *layer = parse_convolutional(options, net, count);
net.types[count] = CONVOLUTIONAL;
net.layers[count] = layer;
net.batch = layer->batch;
}else if(is_connected(s)){
connected_layer *layer = parse_connected(options, net, count);
net.types[count] = CONNECTED;
net.layers[count] = layer;
net.batch = layer->batch;
}else if(is_softmax(s)){
softmax_layer *layer = parse_softmax(options, net, count);
net.types[count] = SOFTMAX;
net.layers[count] = layer;
net.batch = layer->batch;
}else if(is_maxpool(s)){
maxpool_layer *layer = parse_maxpool(options, net, count);
net.types[count] = MAXPOOL;
net.layers[count] = layer;
net.batch = layer->batch;
}else if(is_normalization(s)){
normalization_layer *layer = parse_normalization(options, net, count);
net.types[count] = NORMALIZATION;
net.layers[count] = layer;
net.batch = layer->batch;
}else{
fprintf(stderr, "Type not recognized: %s\n", s->type);
}
@ -208,6 +247,11 @@ int is_softmax(section *s)
return (strcmp(s->type, "[soft]")==0
|| strcmp(s->type, "[softmax]")==0);
}
int is_normalization(section *s)
{
return (strcmp(s->type, "[lrnorm]")==0
|| strcmp(s->type, "[localresponsenormalization]")==0);
}
int read_option(char *s, list *options)
{

41
src/softmax_layer.c
View File

@ -3,13 +3,14 @@
#include <stdlib.h>
#include <stdio.h>
softmax_layer *make_softmax_layer(int inputs)
softmax_layer *make_softmax_layer(int batch, int inputs)
{
fprintf(stderr, "Softmax Layer: %d inputs\n", inputs);
softmax_layer *layer = calloc(1, sizeof(softmax_layer));
layer->batch = batch;
layer->inputs = inputs;
layer->output = calloc(inputs, sizeof(float));
layer->delta = calloc(inputs, sizeof(float));
layer->output = calloc(inputs*batch, sizeof(float));
layer->delta = calloc(inputs*batch, sizeof(float));
return layer;
}
@ -28,28 +29,30 @@ void forward_softmax_layer(const softmax_layer layer, float *input)
*/
void forward_softmax_layer(const softmax_layer layer, float *input)
{
int i;
float sum = 0;
float largest = 0;
for(i = 0; i < layer.inputs; ++i){
if(input[i] > largest) largest = input[i];
}
for(i = 0; i < layer.inputs; ++i){
sum += exp(input[i]-largest);
//printf("%f, ", input[i]);
}
//printf("\n");
if(sum) sum = largest+log(sum);
else sum = largest-100;
for(i = 0; i < layer.inputs; ++i){
layer.output[i] = exp(input[i]-sum);
int i,b;
for(b = 0; b < layer.batch; ++b){
float sum = 0;
float largest = 0;
for(i = 0; i < layer.inputs; ++i){
if(input[i+b*layer.inputs] > largest) largest = input[i+b*layer.inputs];
}
for(i = 0; i < layer.inputs; ++i){
sum += exp(input[i+b*layer.inputs]-largest);
//printf("%f, ", input[i]);
}
//printf("\n");
if(sum) sum = largest+log(sum);
else sum = largest-100;
for(i = 0; i < layer.inputs; ++i){
layer.output[i+b*layer.inputs] = exp(input[i+b*layer.inputs]-sum);
}
}
}
void backward_softmax_layer(const softmax_layer layer, float *input, float *delta)
{
int i;
for(i = 0; i < layer.inputs; ++i){
for(i = 0; i < layer.inputs*layer.batch; ++i){
delta[i] = layer.delta[i];
}
}

3
src/softmax_layer.h
View File

@ -3,11 +3,12 @@
typedef struct {
int inputs;
int batch;
float *delta;
float *output;
} softmax_layer;
softmax_layer *make_softmax_layer(int inputs);
softmax_layer *make_softmax_layer(int batch, int inputs);
void forward_softmax_layer(const softmax_layer layer, float *input);
void backward_softmax_layer(const softmax_layer layer, float *input, float *delta);

1226
src/tests.c
View File

File diff suppressed because it is too large Load Diff

BIN
test.jpg
View File

Binary file not shown.
Side by Side

Before

Width:  |  Height:  |  Size: 9.5 KiB

BIN
test_color.png
View File

Binary file not shown.
Side by Side

Before

Width:  |  Height:  |  Size: 1.1 KiB

BIN
test_dog.jpg
View File

Binary file not shown.
Side by Side

Before

Width:  |  Height:  |  Size: 30 KiB

BIN
test_hinton.jpg
View File

Binary file not shown.
Side by Side

Before

Width:  |  Height:  |  Size: 16 KiB