cl_kernel get_kernel(char *kern_name, cl_context context, cl_device_id device) { cl_program program; cl_kernel kernel; cl_int err = 0; char *fin_program_src; const char *program_source = "void sum_reduce_and_store(__local float *sdata,\n" "__global float *store_arr,\n" "float value,\n" "int store_off)\n" "{\n" //Note that this draws from NVIDIA's reduction example: //- Doesn't use % operator. //- Uses contiguous threads. //- Uses sequential addressing -- no divergence or bank conflicts. //- Is completely unrolled. // local size must be a power of 2 and (>= 64 or == 1) "unsigned int lsz = get_local_size(0);\n" "unsigned int lid = get_local_id(0);\n" "sdata[lid] = value;\n" "barrier(CLK_LOCAL_MEM_FENCE);\n" // do reduction in shared mem "if (lsz != 1) {\n" "if (lsz >= 512) { if (lid < 256) { sdata[lid] += sdata[lid + 256]; } barrier(CLK_LOCAL_MEM_FENCE); }\n" "if (lsz >= 256) { if (lid < 128) { sdata[lid] += sdata[lid + 128]; } barrier(CLK_LOCAL_MEM_FENCE); }\n" "if (lsz >= 128) { if (lid < 64) { sdata[lid] += sdata[lid + 64]; } barrier(CLK_LOCAL_MEM_FENCE); }\n" //Avoid extra if statements by only using local size >= 64 "if (lid < 32) { sdata[lid] += sdata[lid + 32]; } barrier(CLK_LOCAL_MEM_FENCE);\n" "if (lid < 16) { sdata[lid] += sdata[lid + 16]; } barrier(CLK_LOCAL_MEM_FENCE);\n" "if (lid < 8) { sdata[lid] += sdata[lid + 8]; } barrier(CLK_LOCAL_MEM_FENCE);\n" "if (lid < 4) { sdata[lid] += sdata[lid + 4]; } barrier(CLK_LOCAL_MEM_FENCE);\n" "if (lid < 2) { sdata[lid] += sdata[lid + 2]; } barrier(CLK_LOCAL_MEM_FENCE);\n" "if (lid < 1) { sdata[lid] += sdata[lid + 1]; } barrier(CLK_LOCAL_MEM_FENCE);\n" "}\n" // write result for this block to global mem "if (lid == 0) store_arr[store_off] = sdata[0];\n" "barrier(CLK_LOCAL_MEM_FENCE);\n" "}\n" "void global_sum_and_reduce(__local float *reduce_s,\n" "__global float *reduce_arr,\n" "int beg_off,\n" "int arr_len)\n" "{\n" "unsigned int lsz = get_local_size(0);\n" "unsigned int i;" "float value = 0.0f;\n" "if (get_group_id(0) != 0)\n" "return;\n" //Reduce the entire array using one work group "for(i = get_local_id(0); i < arr_len; i += lsz)\n" "if (i < arr_len)\n" "value += reduce_arr[beg_off+i];\n" "sum_reduce_and_store(reduce_s, reduce_arr, value, beg_off);\n" "}\n" "__kernel void consolidate_train(__constant unsigned int *sizes,\n" "__global unsigned int *num_layers,\n" "__global unsigned int *num_neurons, \n" "__global unsigned int *num_inputs,\n" "__global unsigned int *num_outputs,\n" "__global float *MSE_values,\n" "__global float *num_bit_fail,\n" "__global float *train_errors,\n" "__global float *weights_deltas,\n" "__local float *reduce_s)\n" "{\n" "unsigned int input_sz = get_global_size(0);\n" "unsigned int gnum;\n" "unsigned int l;\n" //Calculate the number of groups used in the training run "if (sizes[5] %% sizes[7])\n" "gnum = 1 + (sizes[5] / sizes[7]);\n" "else\n" "gnum = sizes[5] / sizes[7];\n" //Calculate for all layers "for(l = 0; l < num_layers[get_global_id(1)]; ++l) {\n" "unsigned int part_layer_off = get_global_id(1)*sizes[1]+l;\n" "unsigned int num_neurons_l = num_neurons[part_layer_off];\n" "unsigned int n_layer_off = sizes[2]*part_layer_off;\n" "unsigned int o_layer_off = sizes[4]*part_layer_off;\n" "unsigned int n;\n" //Calcalate for all neurons "for(n = 0; n < num_neurons[part_layer_off]; ++n) {\n" "unsigned int num_outputs_l = num_outputs[n_layer_off+n];\n" "unsigned int num_inputs_l = num_inputs[n_layer_off+n];\n" "unsigned int o;\n" //Calculate for all outputs "for(o = 0; o < num_outputs_l; ++o) {\n" "unsigned int i;\n" //Sum delta data "for(i = 0; i < num_inputs_l; ++i) {\n" "global_sum_and_reduce(reduce_s, weights_deltas,\n" "((o_layer_off+o)*sizes[3]+i)*gnum, gnum);\n" // "if(get_global_id(0) == 0)\n" // "printf(\"d (l, n, o, i) (delta): (%%5d %%5d %%5d %%5d) (%%5d %%10f)\\n\", l, n, o, i, ((o_layer_off+o)*sizes[3]+i)*gsz, weights_deltas[((o_layer_off+o)*sizes[3]+i)*gsz]);\n" "}\n" "global_sum_and_reduce(reduce_s, train_errors,\n" "(o_layer_off+o)*sizes[5], sizes[5]);\n" // "printf(\"e (l, n, o) (errs): (%%5d %%5d %%5d) (%%5d %%10f)\\n\", l, n, o, (o_layer_off+o)*input_sz, train_errors[(o_layer_off+o)*input_sz]);\n" "}\n" "o_layer_off += num_outputs_l;\n" "}\n" "}\n" "global_sum_and_reduce(reduce_s, MSE_values, get_global_id(1)*sizes[5], sizes[5]);\n" "global_sum_and_reduce(reduce_s, num_bit_fail, get_global_id(1)*sizes[5], sizes[5]);\n" // "printf(\"m (msev): (%%10f)\\n\", MSE_values[get_global_id(1)*input_sz]);\n" // "printf(\"n (fail): (%%10f)\\n\", num_bit_fail[get_global_id(1)*input_sz]);\n" "}\n" "float activation_derived(float steepness, int act_func,\n" "__global float *outputs,\n" "__global float *sums, int o_i)\n" "{\n" "switch (act_func)\n" "{\n" "case %d:\n" "case %d:\n" "case %d:\n" "return " QUOTEME(fann_linear_derive(steepness, outputs[o_i])) ";\n" "case %d:\n" "case %d:\n" "return " QUOTEME(fann_sigmoid_derive(steepness, fann_clip(outputs[o_i], 0.01f, 0.99f))) ";\n" "case %d:\n" "case %d:\n" "return " QUOTEME(fann_sigmoid_symmetric_derive(steepness, fann_clip(outputs[o_i], -0.98f, 0.98f))) ";\n" "case %d:\n" "return " QUOTEME(fann_gaussian_derive(steepness, outputs[o_i], sums[o_i])) ";\n" "case %d:\n" "return " QUOTEME(fann_gaussian_symmetric_derive(steepness, outputs[o_i], sums[o_i])) ";\n" "case %d:\n" "return " QUOTEME(fann_elliot_derive(steepness, fann_clip(outputs[o_i], 0.01f, 0.99f), sums[o_i])) ";\n" "case %d:\n" "return " QUOTEME(fann_elliot_symmetric_derive(steepness, fann_clip(outputs[o_i], -0.98f, 0.98f), sums[o_i])) ";\n" "case %d:\n" "return " QUOTEME(fann_sin_symmetric_derive(steepness, sums[o_i])) ";\n" "case %d:\n" "return " QUOTEME(fann_cos_symmetric_derive(steepness, sums[o_i])) ";\n" "case %d:\n" "return " QUOTEME(fann_sin_derive(steepness, sums[o_i])) ";\n" "case %d:\n" "return " QUOTEME(fann_cos_derive(steepness, sums[o_i])) ";\n" "case %d: //This should be an error\n" "case %d: //This should be an error\n" "case %d: //FIXME\n" "return -99.0;\n" "}\n" "}\n" "void backpropagate_MSE(__constant unsigned int *sizes,\n" "__global unsigned int *num_layers,\n" "__global unsigned int *num_neurons, \n" "__global unsigned int *num_inputs,\n" "__global unsigned int *num_outputs,\n" "__global float *steepness,\n" "__global int *activation,\n" "__global float *weights,\n" "__global float *inputs,\n" "__global float *sums,\n" "__global float *outputs,\n" "__global float *train_errors,\n" "__global float *weights_deltas,\n" //Shared areas for caching "__local float *steep_s,\n" "__local int *act_s,\n" "__local float *weights_s,\n" "__local float *reduce_s )\n" "{\n" "unsigned int input_id = get_global_id(0);\n" "unsigned int lid = get_local_id(0);\n" "unsigned int lsz = get_local_size(0);\n" "unsigned int gnum;\n" "unsigned int gid = get_group_id(0);\n" "int l;\n" //Calculate the number of groups used in the training run "if (sizes[5] %% sizes[7])\n" "gnum = 1 + (sizes[5] / sizes[7]);\n" "else\n" "gnum = sizes[5] / sizes[7];\n" //Calculate for all layers "for(l = num_layers[get_global_id(1)]-1; l >= 0; --l) {\n" "unsigned int part_layer_off = get_global_id(1)*sizes[1]+l;\n" "unsigned int num_neurons_l = num_neurons[part_layer_off];\n" "unsigned int n_layer_off = sizes[2]*part_layer_off;\n" "unsigned int o_layer_off = sizes[4]*part_layer_off;\n" "unsigned int output_off = o_layer_off-sizes[4];\n" "unsigned int n;\n" //Copy steepness & activation to shared mem "barrier(CLK_LOCAL_MEM_FENCE);\n" "for(n = 0; n < num_neurons_l; n += lsz) {\n" "unsigned int neuron_num = n+lid;\n" "if (neuron_num < num_neurons[part_layer_off]){\n" "steep_s[neuron_num] = steepness[n_layer_off+neuron_num];\n" "act_s[neuron_num] = activation[n_layer_off+neuron_num];\n" "}\n" "}\n" "barrier(CLK_LOCAL_MEM_FENCE);\n" //Clear all the previous layer's train_errors "for(n = 0; n < num_neurons_l && l != 0; ++n) {\n" "unsigned int num_outputs_l = num_outputs[n_layer_off+n];\n" "unsigned int o;\n" //Zero all outputs "for(o = 0; o < num_outputs_l; ++o) {\n" //Don't overrun data "if (sizes[5] > input_id)\n" "train_errors[output_off*sizes[5]+input_id] = 0.0f;\n" "++output_off;\n" "}\n" "}\n" //Reset "output_off = o_layer_off;\n" //Calcalate for all neurons "for(n = 0; n < num_neurons[part_layer_off]; ++n) {\n" "unsigned int num_outputs_l = num_outputs[n_layer_off+n];\n" "unsigned int num_inputs_l = num_inputs[n_layer_off+n];\n" "unsigned int o;\n" //Calculate for all outputs "for(o = 0; o < num_outputs_l; ++o) {\n" "unsigned int i;\n" "unsigned int o_i = output_off*sizes[5]+input_id;\n" "float error;\n" //Multiply errors with the activation function derivative. "if (sizes[5] > input_id)\n" "train_errors[o_i] = error =\n" "train_errors[o_i]*activation_derived(steep_s[n], act_s[n], outputs, sums, o_i);\n" //Weight & sum data from the inputs & bias "for(i = 0; i < num_inputs_l; ++i) {\n" "unsigned int weights_i = 0;\n" "unsigned int prev_output_i = 0;\n" "float delta = 0.0f;\n" //Weights aren't used for first layer "if (l != 0) {\n" "weights_i = (sizes[3]*o+i) %% lsz;\n" //Load shared memory as appropriate "if (weights_i == 0) {\n" "barrier(CLK_LOCAL_MEM_FENCE);\n" "if (sizes[3]*o+i+lid < sizes[3]*num_outputs_l)\n" "weights_s[lid] = weights[output_off*sizes[3]+i+lid];\n" "barrier(CLK_LOCAL_MEM_FENCE);\n" "}\n" "}\n" //Don't overrun data "if (sizes[5] > input_id) {\n" //Figure out what input was used "if(i == num_inputs_l-1){\n" "prev_output_i = (o_layer_off-sizes[4]+i)*sizes[5]+input_id;\n" "delta = error;\n" "} else if(l == 0) {\n" "delta = inputs[i*sizes[5]+input_id] * error;\n" "} else {\n" "prev_output_i = (o_layer_off-sizes[4]+i)*sizes[5]+input_id;\n" "delta = outputs[prev_output_i] * error;\n" "}\n" "}\n" // "printf(\"(id, l, n, o, i) (delta): (%%5d %%5d %%5d %%5d %%5d) (%%10f)\\n\", input_id, l, n, o, i, error*input);\n" // Calculate the weight deltas // Due to memory requirements we're reducing the deltas here "sum_reduce_and_store(reduce_s, weights_deltas, delta,\n" "(output_off*sizes[3]+i)*gnum+gid);\n" // "weights_deltas[(output_off*sizes[3]+i)*gnum+gid] = gnum;\n" // "printf(\"(id, l, n, o, i) (fin delta): (%%5d %%5d %%5d %%5d %%5d) (%%10f)\\n\", input_id, l, n, o, i, weights_deltas[(output_off*sizes[3]+i)*gnum+gid]);\n" //Calculate the error for previous layer "if(l != 0 && sizes[5] > input_id)\n" "train_errors[prev_output_i] += error * weights_s[weights_i];\n" "}\n" "++output_off;\n" "}\n" "}\n" "}\n" "}\n" "void compute_MSE(__constant unsigned int *sizes,\n" "__global float *f_params,\n" "__global unsigned int *num_layers,\n" "__global unsigned int *num_neurons, \n" "__global unsigned int *num_outputs,\n" "__global int *activation,\n" "__global float *outputs,\n" "__global float *train_errors,\n" "__global float *actual_outputs,\n" "__global float *MSE_values,\n" "__global float *num_bit_fail,\n" "__local int *act_s)\n" "{\n" "unsigned int ann_id = get_global_id(1);\n" "unsigned int input_id = get_global_id(0);\n" "unsigned int out_neuron_index = ann_id*sizes[1]+num_layers[ann_id]-1;\n" "unsigned int out_off = out_neuron_index*sizes[4]*sizes[5]+input_id;\n" "unsigned int num_neurons_l = num_neurons[out_neuron_index];\n" "unsigned int layer_off = sizes[2]*out_neuron_index;\n" "unsigned int n;\n" "unsigned int layer_o = 0;\n" "unsigned int bit_fail = 0;\n" "float MSE_value = 0.0f;\n" //Copy steepness & activation to shared mem "for(n = 0; n < num_neurons_l; n += get_local_size(0)) {\n" "unsigned int neuron_off = n+get_local_id(0);\n" "if (neuron_off < num_neurons_l)\n" "act_s[neuron_off] = activation[layer_off+neuron_off];\n" "}\n" "barrier(CLK_LOCAL_MEM_FENCE);\n" //Don't use the extra threads "if (input_id >= sizes[5])\n" "return;\n" //Calcalate for all neurons "for(n = 0; n < num_neurons_l; ++n) {\n" "unsigned int num_outputs_l = num_outputs[layer_off+n];\n" "unsigned int act_out_off = (ann_id*sizes[4]+layer_o)*sizes[5]+input_id;\n" "unsigned int o;\n" //Calculate for all outputs "for(o = 0; o < num_outputs_l; ++o) {\n" "unsigned int out_index = out_off+(layer_o+o)*sizes[5];\n" "float neuron_diff = actual_outputs[act_out_off+o*sizes[5]] - outputs[out_index];\n" //Update MSE macro follows "switch (act_s[n]) {\n" "case %d:\n" "case %d:\n" "case %d:\n" "case %d:\n" "case %d:\n" "case %d:\n" "case %d:\n" "case %d:\n" "neuron_diff *= 0.5f;\n" "}\n" "MSE_value += neuron_diff * neuron_diff;\n" "if(fabs(neuron_diff) >= f_params[0])\n" "++bit_fail;\n" // "printf(\"neuron_diff MSE_value: %%10f %%10f\\n\", neuron_diff, MSE_value);\n" //Update error "if (sizes[9]) {\n" "if(neuron_diff < -.9999999f)\n" "neuron_diff = -17.0f;\n" "else if(neuron_diff > .9999999f)\n" "neuron_diff = 17.0f;\n" "else\n" "neuron_diff = log((1.0f + neuron_diff) / (1.0f - neuron_diff));\n" "}\n" // "printf(\"train_error out_index: %%10f %%5d\\n\", neuron_diff, out_index);\n" "train_errors[out_index] = neuron_diff;\n" //Don't update ann->training_params->num_MSE because it can be calculated later // "printf(\"(%%5d %%5d %%5d) train_errors actual_output neuron_value: %%10f %%10f %%10f\\n\", input_id, n, o, train_errors[out_index], actual_outputs[act_out_off+o*sizes[5]], outputs[out_index]);\n" "}\n" "layer_o += num_outputs_l;\n" "}\n" "unsigned int net_index = ann_id*sizes[5]+input_id;\n" "num_bit_fail[net_index] = bit_fail;\n" "MSE_values[net_index] = MSE_value;\n" "}\n" "float calc_act(float sum, float steepness, int act)\n" "{\n" "float max_sum;\n" "sum *= steepness;\n" "max_sum = 150.0f/steepness;\n" "if(sum > max_sum)\n" "sum = max_sum;\n" "else if(sum < -max_sum)\n" "sum = -max_sum;\n" "switch(act)\n" "{\n" "case %d:\n" "return sum;\n" "case %d:\n" "return ((sum < 0.0f) ? 0.0f : (sum > 1.0f) ? 1.0f : sum);\n" "case %d:\n" "return ((sum < -1.0f) ? -1.0f : (sum > 1.0f) ? 1.0f : sum);\n" "case %d:\n" "return " QUOTEME(fann_sigmoid_real(sum)) ";\n" "case %d:\n" "return " QUOTEME(fann_sigmoid_symmetric_real(sum)) ";\n" "case %d:\n" "return " QUOTEME(fann_stepwise(-2.64665293693542480469e+00f, -1.47221934795379638672e+00f, -5.49306154251098632812e-01f, 5.49306154251098632812e-01f, 1.47221934795379638672e+00f, 2.64665293693542480469e+00f, -9.90000009536743164062e-01f, -8.99999976158142089844e-01f, -5.00000000000000000000e-01f, 5.00000000000000000000e-01f, 8.99999976158142089844e-01f, 9.90000009536743164062e-01f, -1.0f, 1.0f, sum)) ";\n" "case %d:\n" "return " QUOTEME(fann_stepwise(-2.64665246009826660156e+00f, -1.47221946716308593750e+00f, -5.49306154251098632812e-01f, 5.49306154251098632812e-01f, 1.47221934795379638672e+00f, 2.64665293693542480469e+00f, 4.99999988824129104614e-03f, 5.00000007450580596924e-02f, 2.50000000000000000000e-01f, 7.50000000000000000000e-01f, 9.49999988079071044922e-01f, 9.95000004768371582031e-01f, 0.0f, 1.0f, sum)) ";\n" "case %d:\n" "return ((sum < 0.0f) ? 0.0f : 1.0f);\n" "case %d:\n" "return ((sum < 0.0f) ? -1.0f : 1.0f);\n" "case %d:\n" "return " QUOTEME(fann_gaussian_real(sum)) ";\n" "case %d:\n" "return " QUOTEME(fann_gaussian_symmetric_real(sum)) ";\n" "case %d:\n" "return " QUOTEME(fann_elliot_real(sum)) ";\n" "case %d:\n" "return " QUOTEME(fann_elliot_symmetric_real(sum)) ";\n" "case %d:\n" "return " QUOTEME(fann_sin_symmetric_real(sum)) ";\n" "case %d:\n" "return " QUOTEME(fann_cos_symmetric_real(sum)) ";\n" "case %d:\n" "return " QUOTEME(fann_sin_real(sum)) ";\n" "case %d:\n" "return " QUOTEME(fann_cos_real(sum)) ";\n" "case %d:\n" "return 0;\n" "}\n" "}\n" "__kernel void run(__constant unsigned int *sizes,\n" "__global float *f_params,\n" "__global unsigned int *num_layers,\n" "__global unsigned int *num_neurons,\n" "__global unsigned int *num_inputs,\n" "__global unsigned int *num_outputs,\n" "__global float *steepness,\n" "__global int *activation,\n" "__global float *weights,\n" "__global float *inputs,\n" "__global float *sums,\n" "__global float *outputs,\n" "__local float *steep_s,\n" "__local int *act_s,\n" "__local float *weights_s )\n" "{\n" "unsigned int input_id = get_global_id(0);\n" "unsigned int ann_id = get_global_id(1);\n" "unsigned int lid = get_local_id(0);\n" "unsigned int lsz = get_local_size(0);\n" "unsigned int l;\n" //Calculate for all layers "for(l = 0; l < num_layers[ann_id]; ++l) {\n" "unsigned int part_layer_off = ann_id*sizes[1]+l;\n" "unsigned int n;\n" "unsigned int num_neurons_l = num_neurons[part_layer_off];\n" "unsigned int n_layer_off = sizes[2]*part_layer_off;\n" "unsigned int o_layer_off = sizes[4]*part_layer_off;\n" "unsigned int output_off = o_layer_off;\n" //Copy steepness & activation to shared mem "barrier(CLK_LOCAL_MEM_FENCE);\n" "for(n = 0; n < num_neurons[part_layer_off]; n += lsz) {\n" "unsigned int neuron_num = n+lid;\n" "if (neuron_num < num_neurons[part_layer_off]){\n" "steep_s[neuron_num] = steepness[n_layer_off+neuron_num];\n" "act_s[neuron_num] = activation[n_layer_off+neuron_num];\n" "}\n" "}\n" "barrier(CLK_LOCAL_MEM_FENCE);\n" //Calcalate for all neurons "for(n = 0; n < num_neurons[part_layer_off]; ++n) {\n" "unsigned int num_outputs_l = num_outputs[n_layer_off+n];\n" "unsigned int num_inputs_l = num_inputs[n_layer_off+n];\n" "unsigned int o;\n" //Calculate for all outputs "for(o = 0; o < num_outputs_l; ++o) {\n" "unsigned int i;\n" "float sum = 0.0f;\n" //Weight & sum data from the inputs & bias "for(i = 0; i < num_inputs_l; ++i) {\n" "float in_val;\n" "unsigned int weights_i = (sizes[3]*o+i) %% lsz;\n" //Load shared memory as appropriate "if (weights_i == 0) {\n" "barrier(CLK_LOCAL_MEM_FENCE);\n" "if (sizes[3]*o+i+lid < sizes[3]*num_outputs_l)\n" "weights_s[lid] = weights[output_off*sizes[3]+i+lid];\n" "barrier(CLK_LOCAL_MEM_FENCE);\n" "}\n" //Don't overrun data "if (sizes[5] <= input_id)\n" "continue;\n" "if (i == num_inputs_l-1) {\n" //Bias "in_val = 1.0f;\n" // "printf(\"%%5d %%2d %%2d %%2d %%2d: %%15fCL %%15fCL\\n\", input_id, l, n, o ,i, weights_s[weights_i], in_val);\n" "} else if (l == 0) {\n" //Handle input from user "in_val = inputs[i*sizes[5]+input_id];\n" // "printf(\"%%5d %%2d %%2d %%2d %%2d: %%15fCL %%15fCL\\n\", input_id, l, n, o ,i, weights_s[weights_i], in_val);\n" "} else {\n" //Handle input from neurons "in_val = outputs[(o_layer_off-sizes[4]+i)*sizes[5]+input_id];\n" // "printf(\"%%5d %%2d %%2d %%2d %%2d: %%15fCL %%15fCL N:%%5d\\n\", input_id, l, n, o ,i, weights_s[weights_i], in_val, (o_layer_off-sizes[4]+i)*sizes[5]+input_id);\n" "}\n" //Weight & sum "sum += weights_s[weights_i]*in_val;\n" "}\n" //Don't overrun data "if (sizes[5] > input_id){\n" //Save into output data array "sums[output_off*sizes[5]+input_id] = sum;\n" "outputs[output_off*sizes[5]+input_id] = calc_act(sum, steep_s[n], act_s[n]);\n" // "printf(\"%%5d %%2d %%2d %%2d %%2d: %%15f %%15f N:%%5d\\n\", input_id, l, n, o ,i, sums[output_off*sizes[5]+input_id], outputs[output_off*sizes[5]+input_id], output_off*sizes[5]+input_id);\n" "}\n" "++output_off;\n" "}\n" "}\n" "}\n" "}" /* + remember to set ann->training_params->num_MSE on return + remember to init train errors like in fann_compute_MSE() + remember that the group size must be a power of two >= 64 or == 1 for reduce to work remember to set neuron->num_backprop_done on return */ "__kernel void train_batch(\n"//ANN structure "__constant unsigned int *sizes,\n" "__global float *f_params,\n" "__global unsigned int *num_layers,\n" "__global unsigned int *num_neurons,\n" "__global unsigned int *num_inputs,\n" "__global unsigned int *num_outputs,\n" //Network values "__global float *steepness,\n" "__global int *activation,\n" "__global float *weights,\n" //Per-run data "__global float *inputs,\n" "__global float *sums,\n" "__global float *outputs,\n" //Per-run training memory "__global float *train_errors,\n" "__global float *actual_outputs,\n" "__global float *MSE_values,\n" "__global float *num_bit_fail,\n" "__global float *weights_deltas,\n" //Shared areas "__local float *steep_s,\n" "__local int *act_s,\n" "__local float *weights_s,\n" "__local float *reduce_s)\n" "{\n" //Do the normal procedure of an epoch "run(sizes, f_params, num_layers, num_neurons, num_inputs, num_outputs,\n" "steepness, activation, weights, inputs, sums, outputs,\n" "steep_s, act_s, weights_s);\n" "compute_MSE(sizes, f_params, num_layers, num_neurons, num_outputs,\n" "activation, outputs, train_errors, actual_outputs,\n" "MSE_values, num_bit_fail, act_s);\n" "backpropagate_MSE(sizes, num_layers, num_neurons, num_inputs, num_outputs,\n" "steepness, activation, weights, inputs, sums,\n" "outputs, train_errors, weights_deltas,\n" "steep_s, act_s, weights_s, reduce_s);\n" "}\n"; //Insert enum values here because I can't seem to do it at compile time fin_program_src = calloc(128000, sizeof(char)); sprintf(fin_program_src, program_source, FANN_LINEAR, FANN_LINEAR_PIECE, FANN_LINEAR_PIECE_SYMMETRIC, FANN_SIGMOID, FANN_SIGMOID_STEPWISE, FANN_SIGMOID_SYMMETRIC, FANN_SIGMOID_SYMMETRIC_STEPWISE, FANN_GAUSSIAN, FANN_GAUSSIAN_SYMMETRIC, FANN_ELLIOT, FANN_ELLIOT_SYMMETRIC, FANN_SIN_SYMMETRIC, FANN_COS_SYMMETRIC, FANN_SIN, FANN_COS, FANN_THRESHOLD_SYMMETRIC, FANN_THRESHOLD, FANN_GAUSSIAN_STEPWISE, FANN_LINEAR_PIECE_SYMMETRIC, FANN_THRESHOLD_SYMMETRIC, FANN_SIGMOID_SYMMETRIC, FANN_SIGMOID_SYMMETRIC_STEPWISE, FANN_ELLIOT_SYMMETRIC, FANN_GAUSSIAN_SYMMETRIC, FANN_SIN_SYMMETRIC, FANN_COS_SYMMETRIC, FANN_LINEAR, FANN_LINEAR_PIECE, FANN_LINEAR_PIECE_SYMMETRIC, FANN_SIGMOID, FANN_SIGMOID_SYMMETRIC, FANN_SIGMOID_SYMMETRIC_STEPWISE, FANN_SIGMOID_STEPWISE, FANN_THRESHOLD, FANN_THRESHOLD_SYMMETRIC, FANN_GAUSSIAN, FANN_GAUSSIAN_SYMMETRIC, FANN_ELLIOT, FANN_ELLIOT_SYMMETRIC, FANN_SIN_SYMMETRIC, FANN_COS_SYMMETRIC, FANN_SIN, FANN_COS, FANN_GAUSSIAN_STEPWISE); program = clCreateProgramWithSource(context, 1, (const char**)&fin_program_src, NULL, &err); assert(err == CL_SUCCESS); err = clBuildProgram(program, 1, &device, "-cl-fast-relaxed-math -Werror", NULL, NULL); //Detailed debugging info if (err != CL_SUCCESS) { size_t len; char *buffer = (char *)calloc(128000, sizeof(char)); printf("Error: Failed to build program executable!\n"); printf("clBuildProgram return:\n"); if(err == CL_INVALID_PROGRAM) printf("CL_INVALID_PROGRAM\n"); else if(err == CL_INVALID_VALUE) printf("CL_INVALID_VALUE\n"); else if(err == CL_INVALID_BINARY) printf("CL_INVALID_BINARY\n"); else if(err == CL_INVALID_BUILD_OPTIONS) printf("CL_INVALID_BUILD_OPTIONS\n"); else if(err == CL_INVALID_OPERATION) printf("CL_INVALID_OPERATION\n"); else if(err == CL_COMPILER_NOT_AVAILABLE) printf("CL_COMPILER_NOT_AVAILABLE\n"); else if(err == CL_BUILD_PROGRAM_FAILURE) printf("CL_BUILD_PROGRAM_FAILURE\n"); else if(err == CL_OUT_OF_HOST_MEMORY) printf("CL_OUT_OF_HOST_MEMORY\n"); err = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_STATUS, 128000*sizeof(char), buffer, &len); assert(err == CL_SUCCESS); printf("Build Status:\n"); if(buffer[0] == CL_BUILD_NONE) printf("CL_BUILD_NONE\n"); else if(buffer[0] == CL_BUILD_ERROR) printf("CL_BUILD_ERROR\n"); else if(buffer[0] == CL_BUILD_SUCCESS) printf("CL_BUILD_SUCCESS\n"); else if(buffer[0] == CL_BUILD_IN_PROGRESS) printf("CL_BUILD_IN_PROGRESS\n"); err = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 128000*sizeof(char), buffer, &len); printf("Get Build Info:\n"); switch (err) { case CL_INVALID_DEVICE: printf("CL_INVALID_DEVICE\n"); break; case CL_INVALID_VALUE: printf("CL_INVALID_VALUE\n"); break; case CL_INVALID_PROGRAM: printf("CL_INVALID_PROGRAM\n"); break; } printf("Build Info:\n%s\nProgram Source:\n%s\n", buffer, fin_program_src); free(buffer); exit(1); } kernel = clCreateKernel(program, kern_name, &err); clReleaseProgram(program); free(fin_program_src); return kernel; }
static __inline void MAKE_NAME(base_neuron_run)(struct fann * ann, struct fann_neuron * neuron) { unsigned int i, o, num_connections, num_inputs, num_outputs; fann_type *neuron_sums, *outputs, *inputs, *weights; unsigned int activation_function; fann_type steepness; int multiplier = ann->fixed_params->multiplier; unsigned int decimal_point = ann->fixed_params->decimal_point; /* values used for the stepwise linear sigmoid function */ fann_type r1 = 0, r2 = 0, r3 = 0, r4 = 0, r5 = 0, r6 = 0; fann_type v1 = 0, v2 = 0, v3 = 0, v4 = 0, v5 = 0, v6 = 0; fann_type last_steepness = 0; unsigned int last_activation_function = 0; /* Algorith for fully connected networks */ activation_function = neuron->activation_function; steepness = neuron->activation_steepness; num_inputs = neuron->num_inputs; inputs = neuron->inputs; num_outputs = neuron->num_outputs; outputs = neuron->outputs; num_connections = neuron->num_inputs; weights = neuron->weights; neuron_sums=neuron->sums; for (o=0; o<num_outputs ; o++) { neuron_sums[o]=0; /* unrolled loop start */ i = num_connections & 3; /* same as modulo 4 */ switch (i) { case 3: neuron_sums[o] += fann_mult(weights[2], inputs[2]); case 2: neuron_sums[o] += fann_mult(weights[1], inputs[1]); case 1: neuron_sums[o] += fann_mult(weights[0], inputs[0]); case 0: break; } for(; i != num_connections; i += 4) { neuron_sums[o] += fann_mult(weights[i] , inputs[i] ) + fann_mult(weights[i + 1], inputs[i + 1]) + fann_mult(weights[i + 2], inputs[i + 2]) + fann_mult(weights[i + 3], inputs[i + 3]); } weights += num_connections; /* unrolled loop end */ neuron->sums[o] = fann_mult(steepness, neuron_sums[o]); if(activation_function != last_activation_function || steepness != last_steepness) { switch (activation_function) { case FANN_SIGMOID: case FANN_SIGMOID_STEPWISE: r1 = ann->fixed_params->sigmoid_results[0]; r2 = ann->fixed_params->sigmoid_results[1]; r3 = ann->fixed_params->sigmoid_results[2]; r4 = ann->fixed_params->sigmoid_results[3]; r5 = ann->fixed_params->sigmoid_results[4]; r6 = ann->fixed_params->sigmoid_results[5]; v1 = ann->fixed_params->sigmoid_values[0] / steepness; v2 = ann->fixed_params->sigmoid_values[1] / steepness; v3 = ann->fixed_params->sigmoid_values[2] / steepness; v4 = ann->fixed_params->sigmoid_values[3] / steepness; v5 = ann->fixed_params->sigmoid_values[4] / steepness; v6 = ann->fixed_params->sigmoid_values[5] / steepness; break; case FANN_SIGMOID_SYMMETRIC: case FANN_SIGMOID_SYMMETRIC_STEPWISE: r1 = ann->fixed_params->sigmoid_symmetric_results[0]; r2 = ann->fixed_params->sigmoid_symmetric_results[1]; r3 = ann->fixed_params->sigmoid_symmetric_results[2]; r4 = ann->fixed_params->sigmoid_symmetric_results[3]; r5 = ann->fixed_params->sigmoid_symmetric_results[4]; r6 = ann->fixed_params->sigmoid_symmetric_results[5]; v1 = ann->fixed_params->sigmoid_symmetric_values[0] / steepness; v2 = ann->fixed_params->sigmoid_symmetric_values[1] / steepness; v3 = ann->fixed_params->sigmoid_symmetric_values[2] / steepness; v4 = ann->fixed_params->sigmoid_symmetric_values[3] / steepness; v5 = ann->fixed_params->sigmoid_symmetric_values[4] / steepness; v6 = ann->fixed_params->sigmoid_symmetric_values[5] / steepness; break; case FANN_THRESHOLD: break; } } switch (activation_function) { case FANN_SIGMOID: case FANN_SIGMOID_STEPWISE: neuron->outputs[o] = (fann_type) fann_stepwise(v1, v2, v3, v4, v5, v6, r1, r2, r3, r4, r5, r6, 0, multiplier, neuron_sums[o]); break; case FANN_SIGMOID_SYMMETRIC: case FANN_SIGMOID_SYMMETRIC_STEPWISE: neuron->outputs[o] = (fann_type) fann_stepwise(v1, v2, v3, v4, v5, v6, r1, r2, r3, r4, r5, r6, -multiplier, multiplier, neuron_sums[o]); break; case FANN_THRESHOLD: neuron->outputs[o] = (fann_type) ((neuron_sums[o] < 0) ? 0 : multiplier); break; case FANN_THRESHOLD_SYMMETRIC: neuron->outputs[o] = (fann_type) ((neuron_sums[o] < 0) ? -multiplier : multiplier); break; case FANN_LINEAR: neuron->outputs[o] = neuron_sums[o]; break; case FANN_LINEAR_PIECE: neuron->outputs[o] = (fann_type)((neuron_sums[o] < 0) ? 0 : (neuron_sums[o] > multiplier) ? multiplier : neuron_sums[o]); break; case FANN_LINEAR_PIECE_SYMMETRIC: neuron->outputs[o] = (fann_type)((neuron_sums[o] < -multiplier) ? -multiplier : (neuron_sums[o] > multiplier) ? multiplier : neuron_sums[o]); break; case FANN_ELLIOT: case FANN_ELLIOT_SYMMETRIC: case FANN_GAUSSIAN: case FANN_GAUSSIAN_SYMMETRIC: case FANN_GAUSSIAN_STEPWISE: case FANN_SIN_SYMMETRIC: case FANN_COS_SYMMETRIC: fann_error((struct fann_error *) ann, FANN_E_CANT_USE_ACTIVATION); break; } last_steepness = steepness; last_activation_function = activation_function; } }
FANN_EXTERNAL fann_type *FANN_API fann_run(struct fann * ann, fann_type * input) { struct fann_neuron *neuron_it, *last_neuron, *neurons, **neuron_pointers; unsigned int i, num_connections, num_input, num_output; fann_type neuron_sum, *output; fann_type *weights; struct fann_layer *layer_it, *last_layer; unsigned int activation_function; fann_type steepness; /* store some variabels local for fast access */ struct fann_neuron *first_neuron = ann->first_layer->first_neuron; #if 0 __m128 xmm_weight, xmm_neurons, xmm_sum; #endif #ifdef FIXEDFANN int multiplier = ann->multiplier; unsigned int decimal_point = ann->decimal_point; /* values used for the stepwise linear sigmoid function */ fann_type r1 = 0, r2 = 0, r3 = 0, r4 = 0, r5 = 0, r6 = 0; fann_type v1 = 0, v2 = 0, v3 = 0, v4 = 0, v5 = 0, v6 = 0; fann_type last_steepness = 0; unsigned int last_activation_function = 0; #else fann_type max_sum; #endif /* first set the input */ num_input = ann->num_input; for(i = 0; i != num_input; i++) { #ifdef FIXEDFANN if(fann_abs(input[i]) > multiplier) { printf ("Warning input number %d is out of range -%d - %d with value %d, integer overflow may occur.\n", i, multiplier, multiplier, input[i]); } #endif first_neuron[i].value = input[i]; } /* Set the bias neuron in the input layer */ #ifdef FIXEDFANN (ann->first_layer->last_neuron - 1)->value = multiplier; #else (ann->first_layer->last_neuron - 1)->value = 1; #endif last_layer = ann->last_layer; for(layer_it = ann->first_layer + 1; layer_it != last_layer; layer_it++) { last_neuron = layer_it->last_neuron; for(neuron_it = layer_it->first_neuron; neuron_it != last_neuron; neuron_it++) { if(neuron_it->first_con == neuron_it->last_con) { /* bias neurons */ #ifdef FIXEDFANN neuron_it->value = multiplier; #else neuron_it->value = 1; #endif continue; } activation_function = neuron_it->activation_function; steepness = neuron_it->activation_steepness; neuron_sum = 0; num_connections = neuron_it->last_con - neuron_it->first_con; weights = ann->weights + neuron_it->first_con; if(ann->connection_rate >= 1) { if(ann->network_type == FANN_NETTYPE_SHORTCUT) { neurons = ann->first_layer->first_neuron; } else { neurons = (layer_it - 1)->first_neuron; } /* unrolled loop start */ i = num_connections & 3; /* same as modulo 4 */ switch (i) { case 3: neuron_sum += fann_mult(weights[2], neurons[2].value); case 2: neuron_sum += fann_mult(weights[1], neurons[1].value); case 1: neuron_sum += fann_mult(weights[0], neurons[0].value); case 0: break; } #if 0 xmm_sum = _mm_setzero_ps(); for(; i != num_connections; i += 4) { xmm_weight = _mm_loadu_ps(weights + i); xmm_neurons = _mm_set_ps(neurons[i + 3].value,neurons[i + 2].value, neurons[i + 1].value,neurons[i].value); xmm_weight = _mm_mul_ps(xmm_weight,xmm_neurons); xmm_sum =_mm_add_ps(xmm_sum,xmm_weight); } neuron_sum += xmm_sum.m128_f32[3] + xmm_sum.m128_f32[2] + xmm_sum.m128_f32[1] + xmm_sum.m128_f32[0]; #else for(; i != num_connections; i += 4) { neuron_sum += fann_mult(weights[i], neurons[i].value) + fann_mult(weights[i + 1], neurons[i + 1].value) + fann_mult(weights[i + 2], neurons[i + 2].value) + fann_mult(weights[i + 3], neurons[i + 3].value); } #endif /* unrolled loop end */ /* * for(i = 0;i != num_connections; i++){ * printf("%f += %f*%f, ", neuron_sum, weights[i], neurons[i].value); * neuron_sum += fann_mult(weights[i], neurons[i].value); * } */ } else { neuron_pointers = ann->connections + neuron_it->first_con; i = num_connections & 3; /* same as modulo 4 */ switch (i) { case 3: neuron_sum += fann_mult(weights[2], neuron_pointers[2]->value); case 2: neuron_sum += fann_mult(weights[1], neuron_pointers[1]->value); case 1: neuron_sum += fann_mult(weights[0], neuron_pointers[0]->value); case 0: break; } #if 0 xmm_sum = _mm_setzero_ps(); for(; i != num_connections; i += 4) { xmm_weight = _mm_loadu_ps(weights + i); xmm_neurons = _mm_set_ps(neuron_pointers[i + 3]->value,neuron_pointers[i + 2]->value, neuron_pointers[i + 1]->value,neuron_pointers[i + 0]->value); xmm_weight = _mm_mul_ps(xmm_weight,xmm_neurons); xmm_sum =_mm_add_ps(xmm_sum,xmm_weight); } neuron_sum += xmm_sum.m128_f32[3] + xmm_sum.m128_f32[2] + xmm_sum.m128_f32[1] + xmm_sum.m128_f32[0]; #else for(; i != num_connections; i += 4) { neuron_sum += fann_mult(weights[i], neuron_pointers[i]->value) + fann_mult(weights[i + 1], neuron_pointers[i + 1]->value) + fann_mult(weights[i + 2], neuron_pointers[i + 2]->value) + fann_mult(weights[i + 3], neuron_pointers[i + 3]->value); } #endif } #ifdef FIXEDFANN neuron_it->sum = fann_mult(steepness, neuron_sum); if(activation_function != last_activation_function || steepness != last_steepness) { switch (activation_function) { case FANN_SIGMOID: case FANN_SIGMOID_STEPWISE: r1 = ann->sigmoid_results[0]; r2 = ann->sigmoid_results[1]; r3 = ann->sigmoid_results[2]; r4 = ann->sigmoid_results[3]; r5 = ann->sigmoid_results[4]; r6 = ann->sigmoid_results[5]; v1 = ann->sigmoid_values[0] / steepness; v2 = ann->sigmoid_values[1] / steepness; v3 = ann->sigmoid_values[2] / steepness; v4 = ann->sigmoid_values[3] / steepness; v5 = ann->sigmoid_values[4] / steepness; v6 = ann->sigmoid_values[5] / steepness; break; case FANN_SIGMOID_SYMMETRIC: case FANN_SIGMOID_SYMMETRIC_STEPWISE: r1 = ann->sigmoid_symmetric_results[0]; r2 = ann->sigmoid_symmetric_results[1]; r3 = ann->sigmoid_symmetric_results[2]; r4 = ann->sigmoid_symmetric_results[3]; r5 = ann->sigmoid_symmetric_results[4]; r6 = ann->sigmoid_symmetric_results[5]; v1 = ann->sigmoid_symmetric_values[0] / steepness; v2 = ann->sigmoid_symmetric_values[1] / steepness; v3 = ann->sigmoid_symmetric_values[2] / steepness; v4 = ann->sigmoid_symmetric_values[3] / steepness; v5 = ann->sigmoid_symmetric_values[4] / steepness; v6 = ann->sigmoid_symmetric_values[5] / steepness; break; case FANN_THRESHOLD: break; } } switch (activation_function) { case FANN_SIGMOID: case FANN_SIGMOID_STEPWISE: neuron_it->value = (fann_type) fann_stepwise(v1, v2, v3, v4, v5, v6, r1, r2, r3, r4, r5, r6, 0, multiplier, neuron_sum); break; case FANN_SIGMOID_SYMMETRIC: case FANN_SIGMOID_SYMMETRIC_STEPWISE: neuron_it->value = (fann_type) fann_stepwise(v1, v2, v3, v4, v5, v6, r1, r2, r3, r4, r5, r6, -multiplier, multiplier, neuron_sum); break; case FANN_THRESHOLD: neuron_it->value = (fann_type) ((neuron_sum < 0) ? 0 : multiplier); break; case FANN_THRESHOLD_SYMMETRIC: neuron_it->value = (fann_type) ((neuron_sum < 0) ? -multiplier : multiplier); break; case FANN_LINEAR: neuron_it->value = neuron_sum; break; case FANN_LINEAR_PIECE: neuron_it->value = (fann_type)((neuron_sum < 0) ? 0 : (neuron_sum > multiplier) ? multiplier : neuron_sum); break; case FANN_LINEAR_PIECE_SYMMETRIC: neuron_it->value = (fann_type)((neuron_sum < -multiplier) ? -multiplier : (neuron_sum > multiplier) ? multiplier : neuron_sum); break; case FANN_ELLIOT: case FANN_ELLIOT_SYMMETRIC: case FANN_GAUSSIAN: case FANN_GAUSSIAN_SYMMETRIC: case FANN_GAUSSIAN_STEPWISE: case FANN_SIN_SYMMETRIC: case FANN_COS_SYMMETRIC: fann_error((struct fann_error *) ann, FANN_E_CANT_USE_ACTIVATION); break; } last_steepness = steepness; last_activation_function = activation_function; #else neuron_sum = fann_mult(steepness, neuron_sum); max_sum = 150/steepness; if(neuron_sum > max_sum) neuron_sum = max_sum; else if(neuron_sum < -max_sum) neuron_sum = -max_sum; neuron_it->sum = neuron_sum; fann_activation_switch(activation_function, neuron_sum, neuron_it->value); #endif } } /* set the output */ output = ann->output; num_output = ann->num_output; neurons = (ann->last_layer - 1)->first_neuron; for(i = 0; i != num_output; i++) { output[i] = neurons[i].value; } return ann->output; }