diff --git a/src/kernels/MIOpenNeuron.cl b/src/kernels/MIOpenNeuron.cl index 5aad79bee2..7fa1ffdf6d 100644 --- a/src/kernels/MIOpenNeuron.cl +++ b/src/kernels/MIOpenNeuron.cl @@ -303,7 +303,21 @@ MIOpenNeuronFwd(const __global _FLOAT* bot, int i = 0; for(; i < MLO_N_PIXS_OFF; ++i) { - data[i] = bot[xOffset + x * MLO_READ_UNIT + i]; + if(MLO_N_IN_STRIDE > MLO_C_IN * MLO_H_IN * MLO_W_IN || MLO_C_IN_STRIDE > MLO_H_IN * MLO_W_IN || MLO_H_IN_STRIDE > MLO_W_IN || MLO_W_IN_STRIDE > 1) + { + int loc, n_loc, c_loc, h_loc, w_loc; + loc = x * MLO_READ_UNIT + i; + n_loc = loc / MLO_N_IN; + c_loc = (loc % MLO_N_IN) / MLO_C_IN; + h_loc = ((loc % MLO_N_IN) % MLO_C_IN) / MLO_H_IN; + w_loc = ((loc % MLO_N_IN) % MLO_C_IN) % MLO_H_IN; + + data[i] = bot[xOffset + n_loc * MLO_N_IN_STRIDE + c_loc * MLO_C_IN_STRIDE + h_loc * MLO_H_IN_STRIDE + w_loc * MLO_W_IN_STRIDE]; + } + else + { + data[i] = bot[xOffset + x * MLO_READ_UNIT + i]; + } } for(; i < MLO_READ_UNIT; ++i) { @@ -315,7 +329,21 @@ MIOpenNeuronFwd(const __global _FLOAT* bot, { for(int i = 0; i < MLO_READ_UNIT; ++i) { - data[i] = bot[xOffset + x * MLO_READ_UNIT + i]; + if(MLO_N_IN_STRIDE > MLO_C_IN * MLO_H_IN * MLO_W_IN || MLO_C_IN_STRIDE > MLO_H_IN * MLO_W_IN || MLO_H_IN_STRIDE > MLO_W_IN || MLO_W_IN_STRIDE > 1) + { + int loc, n_loc, c_loc, h_loc, w_loc; + loc = x * MLO_READ_UNIT + i; + n_loc = loc / MLO_N_IN; + c_loc = (loc % MLO_N_IN) / MLO_C_IN; + h_loc = ((loc % MLO_N_IN) % MLO_C_IN) / MLO_H_IN; + w_loc = ((loc % MLO_N_IN) % MLO_C_IN) % MLO_H_IN; + + data[i] = bot[xOffset + n_loc * MLO_N_IN_STRIDE + c_loc * MLO_C_IN_STRIDE + h_loc * MLO_H_IN_STRIDE + w_loc * MLO_W_IN_STRIDE]; + } + else + { + data[i] = bot[xOffset + x * MLO_READ_UNIT + i]; + } } } ActivationFunction(MLO_READ_UNIT, response, (const _FLOAT*)data, power, scale, shift); @@ -326,7 +354,21 @@ MIOpenNeuronFwd(const __global _FLOAT* bot, int i = 0; for(; i < MLO_N_PIXS_OFF; ++i) { - top[yOffset + x * MLO_READ_UNIT + i] = response[i]; + if(MLO_N_OUT_STRIDE > MLO_C_OUT * MLO_H_OUT * MLO_W_OUT || MLO_C_OUT_STRIDE > MLO_H_OUT * MLO_W_OUT || MLO_H_OUT_STRIDE > MLO_W_OUT || MLO_W_OUT_STRIDE > 1) + { + int loc, n_loc, c_loc, h_loc, w_loc; + loc = x * MLO_READ_UNIT + i; + n_loc = loc / MLO_N_OUT; + c_loc = (loc % MLO_N_OUT) / MLO_C_OUT; + h_loc = ((loc % MLO_N_OUT) % MLO_C_OUT) / MLO_H_OUT; + w_loc = ((loc % MLO_N_OUT) % MLO_C_OUT) % MLO_H_OUT; + + top[yOffset + n_loc * MLO_N_OUT_STRIDE + c_loc * MLO_C_OUT_STRIDE + h_loc * MLO_H_OUT_STRIDE + w_loc * MLO_W_OUT_STRIDE] = response[i]; + } + else + { + top[yOffset + x * MLO_READ_UNIT + i] = response[i]; + } } } else @@ -334,7 +376,21 @@ MIOpenNeuronFwd(const __global _FLOAT* bot, { for(int i = 0; i < MLO_READ_UNIT; ++i) { - top[yOffset + x * MLO_READ_UNIT + i] = response[i]; + if(MLO_N_OUT_STRIDE > MLO_C_OUT * MLO_H_OUT * MLO_W_OUT || MLO_C_OUT_STRIDE > MLO_H_OUT * MLO_W_OUT || MLO_H_OUT_STRIDE > MLO_W_OUT || MLO_W_OUT_STRIDE > 1) + { + int loc, n_loc, c_loc, h_loc, w_loc; + loc = x * MLO_READ_UNIT + i; + n_loc = loc / MLO_N_OUT; + c_loc = (loc % MLO_N_OUT) / MLO_C_OUT; + h_loc = ((loc % MLO_N_OUT) % MLO_C_OUT) / MLO_H_OUT; + w_loc = ((loc % MLO_N_OUT) % MLO_C_OUT) % MLO_H_OUT; + + top[yOffset + n_loc * MLO_N_OUT_STRIDE + c_loc * MLO_C_OUT_STRIDE + h_loc * MLO_H_OUT_STRIDE + w_loc * MLO_W_OUT_STRIDE] = response[i]; + } + else + { + top[yOffset + x * MLO_READ_UNIT + i] = response[i]; + } } } } diff --git a/src/ocl/activ_ocl.cpp b/src/ocl/activ_ocl.cpp index f964875104..4944c523e9 100644 --- a/src/ocl/activ_ocl.cpp +++ b/src/ocl/activ_ocl.cpp @@ -105,6 +105,23 @@ miopenStatus_t ActivationDescriptor::Forward(Handle& handle, auto f_activ_beta = static_cast(activ_beta); auto f_activ_power = static_cast(activ_power); + compiler_options += " -DMLO_N_IN=" + std::to_string(nIn) + + " -DMLO_C_IN=" + std::to_string(cIn) + + " -DMLO_H_IN=" + std::to_string(hIn) + + " -DMLO_W_IN=" + std::to_string(wIn) + + " -DMLO_N_IN_STRIDE=" + std::to_string(nInStride) + + " -DMLO_C_IN_STRIDE=" + std::to_string(cInStride) + + " -DMLO_H_IN_STRIDE=" + std::to_string(hInStride) + + " -DMLO_W_IN_STRIDE=" + std::to_string(wInStride) + + " -DMLO_N_OUT=" + std::to_string(nOut) + + " -DMLO_C_OUT=" + std::to_string(cOut) + + " -DMLO_H_OUT=" + std::to_string(hOut) + + " -DMLO_W_OUT=" + std::to_string(wOut) + + " -DMLO_N_OUT_STRIDE=" + std::to_string(nOutStride) + + " -DMLO_C_OUT_STRIDE=" + std::to_string(cOutStride) + + " -DMLO_H_OUT_STRIDE=" + std::to_string(hOutStride) + + " -DMLO_W_OUT_STRIDE=" + std::to_string(wOutStride); + handle.GetKernel("miopenActivationForward", network_config, program_name,