-
Notifications
You must be signed in to change notification settings - Fork 75
[AUTOGENERATED] rocm7.1_internal_testing_IFU_2025-09-09 #2625
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
Changes from 1 commit
34aa782
e532c9d
b7dad7d
601ae8e
6b8b3ac
3a20a20
cc5bdd1
1ebd70d
248355f
81aeefa
019fed3
43b7c86
ba7f546
f36f285
0c0e056
869cbcc
8bb213b
8a736fa
61fb632
9eadb37
3302859
d1a15ab
6f7608d
9480cdc
6b1900c
1f51056
3dde5d7
c371032
9e5247f
afa6e56
c3d54de
dbec087
95ee0bf
ef3be67
a3d72b0
48bedd7
d5b3841
b9ba612
9bdcee0
89d41d3
0d71a9d
b04e922
1ec2c15
0d84ff3
09be189
3a20781
9499c87
c7e4107
d2d4c8e
5c67426
2928086
73eb451
be5b03d
b67c410
3bbc2e3
494878a
5da573c
5c473e9
bffc7dd
a714437
2dd529d
06da7c0
f3cebec
b2c7b9a
c2a3024
9837461
261a84a
d711f27
b18bb67
2ef665a
9602590
4902c76
af590cb
a301dc3
031d79c
d63ad53
e02e9ed
9a8d454
c321111
88d94d1
70f865a
adae7f6
3771380
a3c7f77
6087ef4
de893e9
01edcd4
2fa0520
92a4302
771f369
c10195e
a00cdc1
70d36e0
79fcd52
01ab325
e0a62b2
8d50355
9c03d6b
4d4abec
486b20b
081cab0
1463714
4f72d93
0f45aaf
7f4ff79
291cd11
145a3a7
c3ceca2
20629b1
b2b4add
c0983e6
a3e5466
da4db4b
20b47ac
aac1a50
bc50597
c98ddac
28f4ab0
0ff8eab
9aedb3c
5985e28
b6d0a9e
ae0edc1
047603d
541aa23
1a588ac
5927a70
48e3be3
2b8a839
5211f1f
e3068cd
b919560
2a45837
fea2077
eac3d6f
104f268
93fb23d
ada43ed
7a83cf4
9ad5e8e
4348db0
093ab5f
df59c21
e246a85
8235c4f
ff2de5d
ec2e368
eb9073a
5babb4d
103f725
c9ac8c2
29e09a6
1e0656f
fb0afa8
31d5c67
5b90e85
32911ff
e101411
3f59933
25c170b
53297f6
a92773e
f044fa2
8e076d8
49c446c
5793dd7
ebd29a1
72e6717
de5dc1f
bc4176c
314d47a
fbcabb4
d80297a
a0d0266
4e50651
9c991b6
8ec01f3
ec2c137
8f11465
26a1b9c
015423b
5d819f3
dd44faa
fecd968
85fe94e
2c538c9
711c8c8
ac9ccd0
5fd6b6a
189a054
07f0730
c0fc86b
8485aac
897c4e7
ed77e23
6eb14ac
a951f43
eab2afe
12db2a7
a8a187b
e025c0f
583bbf7
a965f09
d8b6622
7b8a645
1641606
4c45090
98ecc0f
065c446
847d7f2
f216d64
607327b
7ad40de
8494afb
4590438
60d0092
7feb8fc
d85392a
d49205f
4840a1a
002e594
dcc42e9
0d9c95c
4dd73e6
e38e953
5ccf3ca
be3b8d2
3ea6868
c0142f5
1f0b01d
f03d635
5eb35d2
b1e99c8
d3c4cf8
24a4dae
d91eecc
e1be887
0ec723a
4b2d297
82f1eb9
af60398
bdbe931
723c27e
8508651
86d34a4
ab55758
ff2ba4c
dba8539
9a66f82
9e7df76
File filter
Filter by extension
Conversations
Jump to
Diff view
Diff view
There are no files selected for viewing
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -1,5 +1 @@ | ||
| <<<<<<< HEAD | ||
| 56765e8c1f6490e21312b46242ed78cb2dd46d35 | ||
| ======= | ||
| fccfc522864cf8bc172abe0cd58ae5581e2d44b9 | ||
| >>>>>>> upstream/main | ||
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -49,20 +49,12 @@ struct OffsetCalculator { | |
| #if defined(USE_ROCM) | ||
| if ((dims > 0) && (dims <= 2)) { | ||
| auto divmod = sizes_[0].divmod(linear_idx); | ||
| <<<<<<< HEAD | ||
| #pragma unroll | ||
|
Collaborator
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. if we merge this line, then next IFU this will not show as conflict.
Collaborator
Author
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. Sorry, I don't follow. I chose upstream change so next time we do IFU it won't show as merge conflict. If I choose HEAD, then we are picking local changes and may show as a merge conflict. |
||
| ======= | ||
| #pragma unroll | ||
| >>>>>>> upstream/main | ||
| for (int arg = 0; arg < NARGS; arg++) | ||
| offsets[arg] = divmod.mod * strides_[0][arg]; | ||
| if (dims >= 2) { | ||
| divmod = sizes_[1].divmod(divmod.div); | ||
| <<<<<<< HEAD | ||
| #pragma unroll | ||
| ======= | ||
| #pragma unroll | ||
| >>>>>>> upstream/main | ||
| for (int arg = 0; arg < NARGS; arg++) | ||
| offsets[arg] += divmod.mod * strides_[1][arg]; | ||
| } | ||
|
|
||
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -91,10 +91,6 @@ constexpr hipDataType HipDataTypeFor<c10::Float4_e2m1fn_x2>() { | |
| #if ROCM_VERSION >= 70000 | ||
| return HIP_R_4F_E2M1; | ||
| #else | ||
| <<<<<<< HEAD | ||
| // Return HIP_R_4F_E2M1 enum value for earlier ROCm version. | ||
|
Collaborator
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. Same comment as above. |
||
| ======= | ||
| >>>>>>> upstream/main | ||
| return static_cast<hipDataType>(33); | ||
| #endif | ||
| } | ||
|
|
||
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -362,18 +362,12 @@ inline at::MemoryFormat miopen_conv_suggest_memory_format(const at::Tensor& inpu | |
| } | ||
|
|
||
| // TODO: Remove PYTORCH_MIOPEN_SUGGEST_NHWC once ROCm officially supports NHWC in MIOpen | ||
| <<<<<<< HEAD | ||
|
||
| // See #64427 | ||
| static std::optional<bool> PYTORCH_MIOPEN_SUGGEST_NHWC = c10::utils::check_env("PYTORCH_MIOPEN_SUGGEST_NHWC"); | ||
| static bool suggest_nhwc = PYTORCH_MIOPEN_SUGGEST_NHWC && *PYTORCH_MIOPEN_SUGGEST_NHWC; | ||
| ======= | ||
| // See https://github.com/pytorch/pytorch/issues/64427. | ||
| // non static variable is used to be able to change environment variable in runtime for testing | ||
| // enabled by default for ROCm >= 7.0.0 with miopen 3.5 | ||
| int miopen_version = detail::getCUDAHooks().compiledWithMIOpen() ? detail::getCUDAHooks().versionMIOpen() : 0; | ||
| bool is_miopen_3_5 = miopen_version >= 30500; // ROCm 7.0 | ||
| bool suggest_nhwc = c10::utils::check_env("PYTORCH_MIOPEN_SUGGEST_NHWC").value_or(is_miopen_3_5); | ||
| >>>>>>> upstream/main | ||
|
|
||
| auto input_memory_format = input.suggest_memory_format(); | ||
| auto weight_memory_format = weight.suggest_memory_format(); | ||
|
|
@@ -383,23 +377,17 @@ inline at::MemoryFormat miopen_conv_suggest_memory_format(const at::Tensor& inpu | |
| (input_memory_format == at::MemoryFormat::ChannelsLast) || | ||
| (weight_memory_format == at::MemoryFormat::ChannelsLast) | ||
| ); | ||
| <<<<<<< HEAD | ||
| ======= | ||
| if (can_use_miopen_channels_last_2d) { | ||
| return at::MemoryFormat::ChannelsLast; | ||
| } | ||
| >>>>>>> upstream/main | ||
|
|
||
| bool can_use_miopen_channels_last_3d = suggest_nhwc && (weight_ndim == 5) && ( | ||
| (input_memory_format == at::MemoryFormat::ChannelsLast3d) || | ||
| (weight_memory_format == at::MemoryFormat::ChannelsLast3d) | ||
| ); | ||
| <<<<<<< HEAD | ||
| ======= | ||
| if (can_use_miopen_channels_last_3d) { | ||
| return at::MemoryFormat::ChannelsLast3d; | ||
| } | ||
| >>>>>>> upstream/main | ||
|
|
||
| return at::MemoryFormat::Contiguous; | ||
| } | ||
|
|
||
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -1406,15 +1406,8 @@ static inline at::MemoryFormat determine_backend_memory_format( | |
| case ConvBackend::Miopen: | ||
| case ConvBackend::MiopenDepthwise: | ||
| case ConvBackend::MiopenTranspose: | ||
| <<<<<<< HEAD | ||
|
||
| if (detail::getCUDAHooks().compiledWithMIOpen() && miopen_conv_use_channels_last(input, weight)) { | ||
| TORCH_INTERNAL_ASSERT((k == 4 || k == 5), | ||
| "Expected 4D or 5D input for miopen memory format selection in determine_backend_memory_format()"); | ||
| backend_memory_format = (k == 5) ? at::MemoryFormat::ChannelsLast3d : at::MemoryFormat::ChannelsLast; | ||
| ======= | ||
| if (detail::getCUDAHooks().compiledWithMIOpen()) { | ||
| backend_memory_format = miopen_conv_suggest_memory_format(input, weight); | ||
| >>>>>>> upstream/main | ||
| } | ||
| break; | ||
| case ConvBackend::Mkldnn: | ||
|
|
||
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
It will be updated to new branch
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Yes, this will be handled separately