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.

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

Implement NLLLoss (contiguous, no reduction, forward only) #3

Closed
wants to merge 25 commits into from

Conversation

hieule88
Copy link
Collaborator

@hieule88 hieule88 commented Apr 15, 2024

  • Added no reduction NLLLoss forward operation and kernel.
  • Added driver test in the new format and gtest for forwarding NLLLoss .
  • New API is guarded by MIOPEN_BETA_API macro.
  • Compared to ROCm pytorch:
float16
op_name dtype size (N C D1 D2) ignore_index model direction ROCm pytorch MIOpen HIP Improvement
NLLLoss float16 [1 2 2 2] -100 random fwd 21166 5866 3.61
NLLLoss float16 [2 10 128 128] 3 random fwd 39069 7004 5.58
NLLLoss float16 [5 13 17 11] 5 random fwd 20318 7679 2.65
NLLLoss float16 [8 12 256 256] -1 random fwd 108630 19749 5.50
NLLLoss float16 [8 16 512 512] 10 random fwd 217693 76579 2.84
NLLLoss float16 [16 21 512 512] 255 stdc fwd 344451 192497 1.79
float32
op_name dtype size (N C D1 D2) ignore_index model direction ROCm pytorch MIOpen HIP Improvement
NLLLoss float32 [1 2 2 2] -100 random fwd 23742 5510 4.31
NLLLoss float32 [2 10 128 128] 3 random fwd 39677 6862 5.78
NLLLoss float32 [5 13 17 11] 5 random fwd 19662 8017 2.45
NLLLoss float32 [8 12 256 256] -1 random fwd 120486 27144 4.44
NLLLoss float32 [8 16 512 512] 10 random fwd 247707 113375 2.18
NLLLoss float32 [16 21 512 512] 255 stdc fwd 371121 275084 1.35
bfloat16
op_name dtype size (N C D1 D2) ignore_index model direction ROCm pytorch MIOpen HIP Improvement
NLLLoss bfloat16 [1 2 2 2] -100 random fwd 24574 5600 4.39
NLLLoss bfloat16 [2 10 128 128] 3 random fwd 39980 6275 6.37
NLLLoss bfloat16 [5 13 17 11] 5 random fwd 20494 7822 2.62
NLLLoss bfloat16 [8 12 256 256] -1 random fwd 110822 19962 5.55
NLLLoss bfloat16 [8 16 512 512] 10 random fwd 230461 74642 3.09
NLLLoss bfloat16 [16 21 512 512] 255 stdc fwd 371521 191911 1.94
  • Average over all cases:
type average
float16 3.66
float32 3.42
bfloat16 3.99

test/gtest/nllloss.cpp Outdated Show resolved Hide resolved
test/gtest/nllloss.hpp Outdated Show resolved Hide resolved
src/CMakeLists.txt Outdated Show resolved Hide resolved
driver/nllloss_driver.hpp Show resolved Hide resolved
src/kernels/MIOpenNLLLoss.cpp Outdated Show resolved Hide resolved
src/solver/nllloss/forward_nllloss.cpp Show resolved Hide resolved
test/gtest/nllloss.hpp Outdated Show resolved Hide resolved
driver/CMakeLists.txt Outdated Show resolved Hide resolved
driver/nllloss_driver.hpp Outdated Show resolved Hide resolved
include/miopen/miopen.h Outdated Show resolved Hide resolved
src/solver/nllloss/forward_nllloss.cpp Outdated Show resolved Hide resolved
@kyeonghwanryu
Copy link

The performance looks good, enough to go to https://github.com/ROCm/MIOpen .
Read the codes of the other cases in NLLLoss like non-contiguous and backward. I'll give you a new ticket soon.
You can change your Jira ticket's status to done. Good work.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

4 participants