1 /* Copyright 2019 The TensorFlow Authors. All Rights Reserved.
2
3 Licensed under the Apache License, Version 2.0 (the "License");
4 you may not use this file except in compliance with the License.
5 You may obtain a copy of the License at
6
7 http://www.apache.org/licenses/LICENSE-2.0
8
9 Unless required by applicable law or agreed to in writing, software
10 distributed under the License is distributed on an "AS IS" BASIS,
11 WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12 See the License for the specific language governing permissions and
13 limitations under the License.
14 ==============================================================================*/
15
16 #include "tensorflow/compiler/xla/debug_options_flags.h"
17 #include "tensorflow/compiler/xla/test.h"
18 #include "tensorflow/compiler/xla/tests/hlo_test_base.h"
19 #include "tensorflow/compiler/xla/tests/test_macros.h"
20
21 namespace xla {
22 namespace {
23
24 class PtxasBugTest : public HloTestBase {};
25
26 // Checks for a bug in ptxas, tracked as Google bug 120501638, and nvidia bug
27 // 2459377. We never received an explanation of what exactly was going wrong
28 // here in ptxas. Known-bad in ptxas 10.0.145, known-good in ptxas 10.0.249.
TEST_F(PtxasBugTest,DoIt)29 TEST_F(PtxasBugTest, DoIt) {
30 const char* const kModuleStr = R"(
31 HloModule test
32
33 add_F32.14 {
34 lhs.15 = f32[] parameter(0)
35 rhs.16 = f32[] parameter(1)
36 ROOT add.17 = f32[] add(lhs.15, rhs.16)
37 }
38
39 ENTRY testcase {
40 arg0.1 = f32[2,5,2]{2,1,0} parameter(0)
41 reshape.2 = f32[2,5,2]{2,1,0} reshape(arg0.1)
42 constant.3 = f32[] constant(0)
43 pad.4 = f32[2,6,2]{2,1,0} pad(reshape.2, constant.3), padding=0_0x0_1x0_0
44 reshape.5 = f32[2,3,2,2]{3,2,1,0} reshape(pad.4)
45 transpose.6 = f32[2,2,3,2]{3,0,2,1} transpose(reshape.5), dimensions={2,0,1,3}
46 reshape.7 = f32[4,3,2]{2,1,0} reshape(transpose.6)
47 reshape.8 = f32[4,1,3,2]{3,2,1,0} reshape(reshape.7)
48 transpose.9 = f32[4,2,1,3]{1,3,2,0} transpose(reshape.8), dimensions={0,3,1,2}
49 convert.10 = f32[4,2,1,3]{1,3,2,0} convert(transpose.9)
50 constant.12 = f32[] constant(0)
51 pad.13 = f32[4,2,1,3]{3,2,1,0} pad(convert.10, constant.12), padding=0_0x0_0x0_0x0_0
52 constant.11 = f32[] constant(0)
53 reduce-window.18 = f32[4,2,1,3]{3,2,1,0} reduce-window(pad.13, constant.11),
54 window={size=1x1x1x1}, to_apply=add_F32.14
55 constant.19 = f32[] constant(1)
56 broadcast.20 = f32[4,2,1,3]{3,2,1,0} broadcast(constant.19), dimensions={}
57 divide.21 = f32[4,2,1,3]{3,2,1,0} divide(reduce-window.18, broadcast.20)
58 convert.22 = f32[4,2,1,3]{3,2,1,0} convert(divide.21)
59 transpose.23 = f32[4,1,3,2]{2,1,3,0} transpose(convert.22), dimensions={0,2,3,1}
60 reshape.24 = f32[4,3,2]{2,1,0} reshape(transpose.23)
61 reshape.25 = f32[2,2,3,2]{3,2,1,0} reshape(reshape.24)
62 transpose.26 = f32[2,3,2,2]{3,1,0,2} transpose(reshape.25), dimensions={1,2,0,3}
63 reshape.27 = f32[2,6,2]{2,1,0} reshape(transpose.26)
64 slice.28 = f32[2,5,2]{2,1,0} slice(reshape.27), slice={[0:2], [0:5], [0:2]}
65 reshape.29 = f32[2,5,2]{2,1,0} reshape(slice.28)
66 tuple.30 = (f32[2,5,2]{2,1,0}) tuple(reshape.29)
67 ROOT get-tuple-element.31 = f32[2,5,2]{2,1,0} get-tuple-element(tuple.30), index=0
68 })";
69
70 // Create a module with the true-default flags, not the default-for-testing
71 // flags. In particular, true-default flags enable unrolling, whereas for
72 // testing we disable unrolling, and this bug doesn't trigger without
73 // unrolling.
74 HloModuleConfig config;
75 config.set_debug_options(DefaultDebugOptionsIgnoringFlags());
76 TF_ASSERT_OK_AND_ASSIGN(auto module,
77 ParseAndReturnVerifiedModule(kModuleStr, config));
78 EXPECT_TRUE(RunAndCompare(std::move(module), ErrorSpec{0.01, 0.01}));
79 }
80
81 } // anonymous namespace
82 } // namespace xla
83