Skip to content

Commit

Permalink
Merge pull request #19 from ce1adon/activoff
Browse files Browse the repository at this point in the history 
Activoff
  • Loading branch information
@ce1adon
ce1adon authored Oct 21, 2017
2 parents 62f1e68 + f00e3bc commit 7c01fa1
Show file tree
Hide file tree
Showing 2 changed files with 77 additions and 4 deletions.
64 changes: 60 additions & 4 deletions src/kernels/MIOpenNeuron.cl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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)
{
Expand All @@ -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);
Expand All @@ -326,15 +354,43 @@ 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
#endif
{
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];
}
}
}
}
Expand Down
17 changes: 17 additions & 0 deletions src/ocl/activ_ocl.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -105,6 +105,23 @@ miopenStatus_t ActivationDescriptor::Forward(Handle& handle,
auto f_activ_beta = static_cast<float>(activ_beta);
auto f_activ_power = static_cast<float>(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,
Expand Down

0 comments on commit 7c01fa1

Please sign in to comment.