diff --git a/src/kernels/activation_elu.cu b/src/kernels/activation_elu.cu index b507b5a..a80e3ae 100644 --- a/src/kernels/activation_elu.cu +++ b/src/kernels/activation_elu.cu @@ -9,7 +9,7 @@ __global__ void activation_elu(value_type *input, value_type *output, int size) { - int i = threadIdx.x*(blockIdx.x +1); + int i = blockDim.x*blockIdx.x + threadIdx.x; if(i>>(srcData, dstData, size); + int blocks = (size+255)/256; + int threads = 256; + + activation_elu<<>>(srcData, dstData, size); checkCuda( cudaDeviceSynchronize() ); } \ No newline at end of file diff --git a/tests/simple_dense.py b/tests/simple_dense.py index 112b5f4..ab56f03 100644 --- a/tests/simple_dense.py +++ b/tests/simple_dense.py @@ -17,18 +17,19 @@ def dense_model(): batch_input_shape=(1, 100, 100, 4), # 100by100by2 output_shape=(100, 100, 4, 1))) # 100by100by2 model.add(Convolution3D(16, kernel_size=(8, 8, 2), subsample=(4, 4, 1), border_mode="valid", - bias_initializer="random_uniform", activation="relu")) - #model.add(ELU()) + bias_initializer="random_uniform")) + model.add(ELU()) model.add(AveragePooling3D(pool_size=(2, 2, 1))) model.add(Convolution3D(16, kernel_size=(4, 4, 2), subsample=(2, 2, 1), border_mode="valid", - bias_initializer="random_uniform", activation="relu")) + bias_initializer="random_uniform")) + model.add(ELU()) model.add(Convolution3D(24, kernel_size=(3, 3, 2), subsample=(1, 1, 1), border_mode="valid", - bias_initializer="random_uniform", activation="relu")) - #model.add(ELU()) + bias_initializer="random_uniform")) + model.add(ELU()) model.add(Flatten()) - model.add(Dense(256, activation="relu", bias_initializer="random_uniform")) + model.add(Dense(256, bias_initializer="random_uniform")) + model.add(ELU()) model.add(Dense(32, activation="relu", bias_initializer="random_uniform")) - #model.add(ELU()) model.add(Dense(2, bias_initializer="random_uniform")) sgd = keras.optimizers.Adam(lr=1e-4, decay=1e-8) diff --git a/tests/test.cpp b/tests/test.cpp index 09b4c33..d576155 100644 --- a/tests/test.cpp +++ b/tests/test.cpp @@ -22,15 +22,15 @@ int main() { tkDNN::dataDim_t dim(1, 1, 100, 100, 4); tkDNN::MulAdd m0 (&net, dim, 2, -1); tkDNN::Conv3d c0 (&net, m0.output_dim, 16, 8, 8, 2, 4, 4, 1, c0_bin, c0_bias_bin); - tkDNN::Activation a0 (&net, c0.output_dim, tkDNN::ACTIVATION_RELU); + tkDNN::Activation a0 (&net, c0.output_dim, tkDNN::ACTIVATION_ELU); tkDNN::Pooling p0 (&net, a0.output_dim, 2, 2, 2, 2, tkDNN::POOLING_AVERAGE); tkDNN::Conv3d c1 (&net, p0.output_dim, 16, 4, 4, 2, 2, 2, 1, c1_bin, c1_bias_bin); - tkDNN::Activation a1 (&net, c1.output_dim, tkDNN::ACTIVATION_RELU); + tkDNN::Activation a1 (&net, c1.output_dim, tkDNN::ACTIVATION_ELU); tkDNN::Conv3d c2 (&net, a1.output_dim, 24, 3, 3, 2, 1, 1, 1, c2_bin, c2_bias_bin); - tkDNN::Activation a2 (&net, c2.output_dim, tkDNN::ACTIVATION_RELU); + tkDNN::Activation a2 (&net, c2.output_dim, tkDNN::ACTIVATION_ELU); tkDNN::Flatten f2 (&net, a2.output_dim); tkDNN::Dense d3 (&net, f2.output_dim, 256, d3_bin, d3_bias_bin); - tkDNN::Activation a3 (&net, d3.output_dim, tkDNN::ACTIVATION_RELU); + tkDNN::Activation a3 (&net, d3.output_dim, tkDNN::ACTIVATION_ELU); tkDNN::Dense d4 (&net, a3.output_dim, 32, d4_bin, d4_bias_bin); tkDNN::Activation a4 (&net, d4.output_dim, tkDNN::ACTIVATION_RELU); tkDNN::Dense d5 (&net, a4.output_dim, 2, d5_bin, d5_bias_bin);