This PR flips a switch to enable PE + TE

[Krovatkin]

This PR flips a switch to enable PE + TE
next PR: #45397

[dr-ci]

💊 CI failures summary and remediations

As of commit 391de07 (more details on the Dr. CI page):

---

• 1/1 failures possibly* introduced in this PR
  • 1/1 non-CircleCI failure(s)

---

ci.pytorch.org: 1 failed

• Failed: pr/pytorch-linux-bionic-rocm3.7-py3.6

---

This comment was automatically generated by Dr. CI (expand for details).

Follow this link to opt-out of these comments for your Pull Requests.

Please report bugs/suggestions on the GitHub issue tracker or post in the (internal) Dr. CI Users group.

See how this bot performed.

This comment has been revised 10 times.

[ZolotukhinM]

🎆

[facebook-github-bot]

@Krovatkin has imported this pull request. If you are a Facebook employee, you can view this diff on Phabricator.

[facebook-github-bot]

@Krovatkin has imported this pull request. If you are a Facebook employee, you can view this diff on Phabricator.

[facebook-github-bot]

@Krovatkin has imported this pull request. If you are a Facebook employee, you can view this diff on Phabricator.

[facebook-github-bot]

@Krovatkin merged this pull request in dddb685.

[t-vi]

This breaks ROCm. In the short term, we could skip CI tests that fail (and leave things broken if real users like those functions) or leave the default to off for now when using ROCm.
@jeffdaily preferences?

[mruberry]

Unlanding. As @t-vi said this does break ROCm. Relevant snippet:

11:18:14 ======================================================================
11:18:14 ERROR: test_triplet_margin_with_distance_loss_cuda (__main__.TestNNDeviceTypeCUDA)
11:18:14 ----------------------------------------------------------------------
11:18:14 Traceback (most recent call last):
11:18:14   File "/var/lib/jenkins/.local/lib/python3.6/site-packages/torch/testing/_internal/common_utils.py", line 827, in wrapper
11:18:14     method(*args, **kwargs)
11:18:14   File "/var/lib/jenkins/.local/lib/python3.6/site-packages/torch/testing/_internal/common_device_type.py", line 273, in instantiated_test
11:18:14     result = test_fn(self, *args)
11:18:14   File "/var/lib/jenkins/.local/lib/python3.6/site-packages/torch/testing/_internal/common_device_type.py", line 651, in only_fn
11:18:14     return fn(self, device, *args, **kwargs)
11:18:14   File "test_nn.py", line 12810, in test_triplet_margin_with_distance_loss
11:18:14     a, p, n), (anchor, positive, negative)))
11:18:14   File "/var/lib/jenkins/.local/lib/python3.6/site-packages/torch/autograd/gradcheck.py", line 322, in gradcheck
11:18:14     func_out = func(*tupled_inputs)
11:18:14   File "test_nn.py", line 12810, in <lambda>
11:18:14     a, p, n), (anchor, positive, negative)))
11:18:14   File "/var/lib/jenkins/.local/lib/python3.6/site-packages/torch/nn/modules/module.py", line 727, in _call_impl
11:18:14     result = self.forward(*input, **kwargs)
11:18:14   File "/var/lib/jenkins/.local/lib/python3.6/site-packages/torch/testing/_internal/common_utils.py", line 129, in prof_meth_call
11:18:14     return prof_callable(meth_call, *args, **kwargs)
11:18:14   File "/var/lib/jenkins/.local/lib/python3.6/site-packages/torch/testing/_internal/common_utils.py", line 123, in prof_callable
11:18:14     return callable(*args, **kwargs)
11:18:14 RuntimeError: /tmp/comgr-b2ca31/input/CompileSource:16:12: error: unknown type name '__global__'
11:18:14 extern "C" __global__
11:18:14            ^
11:18:14 /tmp/comgr-b2ca31/input/CompileSource:19:13: error: use of undeclared identifier 'blockIdx'
11:18:14   if (512 * blockIdx.x + threadIdx.x<5 ? 1 : 0) {
11:18:14             ^
11:18:14 /tmp/comgr-b2ca31/input/CompileSource:19:26: error: use of undeclared identifier 'threadIdx'
11:18:14   if (512 * blockIdx.x + threadIdx.x<5 ? 1 : 0) {
11:18:14                          ^
11:18:14 /tmp/comgr-b2ca31/input/CompileSource:20:33: error: use of undeclared identifier 'blockIdx'
11:18:14     double v = __ldg(t0 + 512 * blockIdx.x + threadIdx.x);
11:18:14                                 ^
11:18:14 /tmp/comgr-b2ca31/input/CompileSource:20:46: error: use of undeclared identifier 'threadIdx'
11:18:14     double v = __ldg(t0 + 512 * blockIdx.x + threadIdx.x);
11:18:14                                              ^
11:18:14 /tmp/comgr-b2ca31/input/CompileSource:21:35: error: use of undeclared identifier 'blockIdx'
11:18:14     double v_1 = __ldg(t1 + 512 * blockIdx.x + threadIdx.x);
11:18:14                                   ^
11:18:14 /tmp/comgr-b2ca31/input/CompileSource:21:48: error: use of undeclared identifier 'threadIdx'
11:18:14     double v_1 = __ldg(t1 + 512 * blockIdx.x + threadIdx.x);
11:18:14                                                ^
11:18:14 /tmp/comgr-b2ca31/input/CompileSource:22:35: error: use of undeclared identifier 'blockIdx'
11:18:14     double v_2 = __ldg(t2 + 512 * blockIdx.x + threadIdx.x);
11:18:14                                   ^
11:18:14 /tmp/comgr-b2ca31/input/CompileSource:22:48: error: use of undeclared identifier 'threadIdx'
11:18:14     double v_2 = __ldg(t2 + 512 * blockIdx.x + threadIdx.x);
11:18:14                                                ^
11:18:14 /tmp/comgr-b2ca31/input/CompileSource:23:35: error: use of undeclared identifier 'blockIdx'
11:18:14     double v_3 = __ldg(t0 + 512 * blockIdx.x + threadIdx.x);
11:18:14                                   ^
11:18:14 /tmp/comgr-b2ca31/input/CompileSource:23:48: error: use of undeclared identifier 'threadIdx'
11:18:14     double v_3 = __ldg(t0 + 512 * blockIdx.x + threadIdx.x);
11:18:14                                                ^
11:18:14 /tmp/comgr-b2ca31/input/CompileSource:24:35: error: use of undeclared identifier 'blockIdx'
11:18:14     double v_4 = __ldg(t1 + 512 * blockIdx.x + threadIdx.x);
11:18:14                                   ^
11:18:14 /tmp/comgr-b2ca31/input/CompileSource:24:48: error: use of undeclared identifier 'threadIdx'
11:18:14     double v_4 = __ldg(t1 + 512 * blockIdx.x + threadIdx.x);
11:18:14                                                ^
11:18:14 /tmp/comgr-b2ca31/input/CompileSource:25:35: error: use of undeclared identifier 'blockIdx'
11:18:14     double v_5 = __ldg(t2 + 512 * blockIdx.x + threadIdx.x);
11:18:14                                   ^
11:18:14 /tmp/comgr-b2ca31/input/CompileSource:25:48: error: use of undeclared identifier 'threadIdx'
11:18:14     double v_5 = __ldg(t2 + 512 * blockIdx.x + threadIdx.x);
11:18:14                                                ^
11:18:14 /tmp/comgr-b2ca31/input/CompileSource:26:27: error: use of undeclared identifier 'blockIdx'
11:18:14     aten_clamp_flat[512 * blockIdx.x + threadIdx.x] = (v - (minimum(v_1,v_2))) + 1<0 ? 0 : (v_3 - (minimum(v_4,v_5))) + 1;
11:18:14                           ^
11:18:14 /tmp/comgr-b2ca31/input/CompileSource:26:40: error: use of undeclared identifier 'threadIdx'
11:18:14     aten_clamp_flat[512 * blockIdx.x + threadIdx.x] = (v - (minimum(v_1,v_2))) + 1<0 ? 0 : (v_3 - (minimum(v_4,v_5))) + 1;
11:18:14                                        ^
11:18:14 17 errors generated when compiling for gfx900.
11:18:14 Error: Failed to compile opencl source (from CL or HIP source to LLVM IR).
11:18:14 
11:18:14 nvrtc compilation failed: 
11:18:14 
11:18:14 #define NAN __int_as_float(0x7fffffff)
11:18:14 #define POS_INFINITY __int_as_float(0x7f800000)
11:18:14 #define NEG_INFINITY __int_as_float(0xff800000)
11:18:14 
11:18:14 template<typename T>
11:18:14 T maximum(T a, T b) {
11:18:14   return isnan(a) ? a : (a > b ? a : b);
11:18:14 }
11:18:14 
11:18:14 template<typename T>
11:18:14 T minimum(T a, T b) {
11:18:14   return isnan(a) ? a : (a < b ? a : b);
11:18:14 }
11:18:14 
11:18:14 extern "C" __global__
11:18:14 void func_1(double* t0, double* t1, double* t2, double* aten_clamp_flat) {
11:18:14 {
11:18:14   if (512 * blockIdx.x + threadIdx.x<5 ? 1 : 0) {
11:18:14     double v = __ldg(t0 + 512 * blockIdx.x + threadIdx.x);
11:18:14     double v_1 = __ldg(t1 + 512 * blockIdx.x + threadIdx.x);
11:18:14     double v_2 = __ldg(t2 + 512 * blockIdx.x + threadIdx.x);
11:18:14     double v_3 = __ldg(t0 + 512 * blockIdx.x + threadIdx.x);
11:18:14     double v_4 = __ldg(t1 + 512 * blockIdx.x + threadIdx.x);
11:18:14     double v_5 = __ldg(t2 + 512 * blockIdx.x + threadIdx.x);
11:18:14     aten_clamp_flat[512 * blockIdx.x + threadIdx.x] = (v - (minimum(v_1,v_2))) + 1<0 ? 0 : (v_3 - (minimum(v_4,v_5))) + 1;
11:18:14   }
11:18:14 }
11:18:14 }

[t-vi]

Just a quick shout that I'm working on fixing the codegen quirks for ROCm. In fact, I do have the failing test compiling passing, but I'll need a few hours before it is ready for a PR.

[t-vi]

With ROCm hopefully working with PE + TE, should we have another go at this?