Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Softmax log backward : Increase precision of fp16's accumulator to fp32 #3427

Open
wants to merge 6 commits into
base: develop
Choose a base branch
from
+27 −8
Open

Softmax log backward : Increase precision of fp16's accumulator to fp32 #3427

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
f086165
- initialize y for softmax backward
bghimireamd Dec 6, 2024
df21f02
fix typo
bghimireamd Dec 6, 2024
8680a49
introduce _FLOAT_ACCUM
bghimireamd Dec 8, 2024
6103106
move definition
bghimireamd Dec 9, 2024
9208e47
fix local variable
bghimireamd Dec 10, 2024
bcf5ab5
Merge branch 'develop' of github.com:ROCm/MIOpen into bg/softmax_log_…
bghimireamd Dec 11, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
12 changes: 12 additions & 0 deletions driver/softmax_driver.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -115,6 +115,7 @@ class SoftmaxDriver : public Driver
float beta;
miopenSoftmaxAlgorithm_t algo;
miopenSoftmaxMode_t mode;
bool isForward = false;
};

template <typename Tgpu, typename Tref>
Expand Down Expand Up @@ -181,6 +182,8 @@ std::vector<int> SoftmaxDriver<Tgpu, Tref>::GetInputTensorLengthsFromCmdLine()
int in_h = inflags.GetValueInt("in_h");
int in_w = inflags.GetValueInt("in_w");

isForward = inflags.GetValueInt("forw") == 1;

return std::vector<int>({in_n, in_c, in_h, in_w});
}

Expand Down Expand Up @@ -220,6 +223,15 @@ int SoftmaxDriver<Tgpu, Tref>::AllocateBuffersAndCopy()
dout[i] = Data_scale * prng::gen_A_to_B(static_cast<Tgpu>(-0.5), static_cast<Tgpu>(0.5));
}

// if bwd then initialize the y
if(!isForward)
{
for(int i = 0; i < out_sz; i++)
{
out[i] = prng::gen_A_to_B(static_cast<Tgpu>(-5.0), static_cast<Tgpu>(5.0));
}
}

status_t status;
status = in_dev->ToGPU(q, in.data());
status |= out_dev->ToGPU(q, out.data());
Expand Down
23 changes: 15 additions & 8 deletions src/kernels/MIOpenSoftmax.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -65,6 +65,12 @@
#define USE_SOFTMAX_LOG 0
#endif

#if USE_SOFTMAX_LOG && MIOPEN_USE_FP16
#define _FLOAT_ACCUM float
#else
#define _FLOAT_ACCUM _FLOAT
#endif

#ifndef USE_SOFTMAX_ACCURATE
#define USE_SOFTMAX_ACCURATE 0
#endif
Expand Down Expand Up @@ -689,7 +695,8 @@ __kernel void SoftmaxBackward(global _FLOAT* y,
{

#if NUM_BATCH == 1 // CSR-Vector like appraoch
local _FLOAT l_helper[256];

local _FLOAT_ACCUM l_helper[256];

int gid = get_group_id(0);
int lid = get_local_id(0);
Expand All @@ -705,7 +712,7 @@ __kernel void SoftmaxBackward(global _FLOAT* y,
int s1 = s % input_w;
#endif

_FLOAT channel_dot = (_FLOAT)0; // thread_local helper var
_FLOAT_ACCUM channel_dot = (_FLOAT_ACCUM)0; // thread_local helper var

// Compute dot product per channel
// Iterate over all the channels one thread is supposed to loop over
Expand Down Expand Up @@ -835,7 +842,7 @@ __kernel void SoftmaxBackward(global _FLOAT* y,

#else

local _FLOAT l_helper[256];
local _FLOAT_ACCUM l_helper[256];

int gid = get_group_id(0);
int lid = get_local_id(0);
Expand All @@ -845,13 +852,13 @@ __kernel void SoftmaxBackward(global _FLOAT* y,
int batch = lid / BATCH_SIZE; // which spatial_dim or pixel

// Batch specific n and s
int batch_n = (NUM_BATCH * gid + batch) / spatial_dim; // nth image
int batch_s = (NUM_BATCH * gid + batch) % spatial_dim; // which spatial_dim/pixel
int batch_n = (NUM_BATCH * gid + batch) / spatial_dim; // nth image
int batch_s = (NUM_BATCH * gid + batch) % spatial_dim; // which spatial_dim/pixel
#if(!IS_DINPUT_PACKED || !IS_DOUTPUT_PACKED || !IS_OUTPUT_PACKED) && USE_SOFTMAX_MODE_CHANNEL
int batch_s0 = batch_s / input_w;
int batch_s1 = batch_s % input_w;
int batch_s0 = batch_s / input_w;
int batch_s1 = batch_s % input_w;
#endif
_FLOAT channel_dot = (_FLOAT)(0); // thread_local helper var
_FLOAT_ACCUM channel_dot = (_FLOAT)0; // thread_local helper var

// stores all the values touched by one thread so that we do not have load
// again as the CSR-Vector approach
Expand Down
Loading