llvmbot wrote:

<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang

@llvm/pr-subscribers-llvm-ir

Author: None (DeanSturtevant1)

<details>
<summary>Changes</summary>



---

Patch is 142.46 KiB, truncated to 20.00 KiB below, full version: 
https://github.com/llvm/llvm-project/pull/147726.diff


44 Files Affected:

- (modified) clang/include/clang/Basic/BuiltinsAMDGPU.def (+1) 
- (modified) clang/test/CodeGenOpenCL/amdgpu-features.cl (+1-1) 
- (modified) clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl (+38) 
- (modified) clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl (+5) 
- (modified) lldb/source/Host/windows/MainLoopWindows.cpp (+25-8) 
- (modified) lldb/unittests/Host/MainLoopTest.cpp (+140-1) 
- (modified) llvm/include/llvm/Analysis/TargetTransformInfo.h (+7-6) 
- (modified) llvm/include/llvm/Analysis/TargetTransformInfoImpl.h (+3-2) 
- (modified) llvm/include/llvm/CodeGen/SDPatternMatch.h (+28-14) 
- (modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+6) 
- (modified) llvm/lib/Analysis/TargetTransformInfo.cpp (+3-2) 
- (modified) llvm/lib/IR/Verifier.cpp (+27-3) 
- (modified) llvm/lib/Target/AArch64/AArch64TargetTransformInfo.cpp (+4-4) 
- (modified) llvm/lib/Target/AArch64/AArch64TargetTransformInfo.h (+3-2) 
- (modified) llvm/lib/Target/AMDGPU/AMDGPU.td (+13) 
- (modified) llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp (+1) 
- (modified) llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp (+4-4) 
- (modified) llvm/lib/Target/AMDGPU/GCNSubtarget.h (+3) 
- (modified) llvm/lib/Target/AMDGPU/VOP1Instructions.td (+25-7) 
- (modified) llvm/lib/Target/AMDGPU/VOPInstructions.td (+6-6) 
- (modified) llvm/lib/TargetParser/TargetParser.cpp (+1) 
- (modified) llvm/lib/Transforms/Scalar/EarlyCSE.cpp (+9-6) 
- (added) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.fp8.e5m3.ll (+50) 
- (modified) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.fp8.ll (+267) 
- (modified) llvm/test/CodeGen/RISCV/rvv/vp-splice-mask-vectors.ll (+43-41) 
- (modified) llvm/test/CodeGen/RISCV/rvv/vp-splice.ll (+107-105) 
- (modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s (+9) 
- (modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1.s (+9) 
- (modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1-fake16.s (+45) 
- (modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1.s (+45) 
- (modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16-fake16.s 
(+12) 
- (modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16.s (+12) 
- (modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8-fake16.s (+12) 
- (modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8.s (+12) 
- (modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1.txt (+9) 
- (modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1_dpp16.txt (+6) 
- (modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1.txt 
(+45) 
- (modified) 
llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1_dpp16.txt (+9) 
- (modified) 
llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1_dpp8.txt (+9) 
- (modified) llvm/test/Transforms/EarlyCSE/AArch64/intrinsics.ll (+17-1) 
- (modified) llvm/test/Verifier/invalid-vp-intrinsics.ll (+33) 
- (modified) llvm/unittests/CodeGen/SelectionDAGPatternMatchTest.cpp (+33) 
- (modified) utils/bazel/llvm-project-overlay/llvm/BUILD.bazel (+1) 
- (modified) utils/bazel/llvm-project-overlay/mlir/BUILD.bazel (+37) 


``````````diff
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def 
b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index fb358297a5eed..a5ee8013adff6 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -429,6 +429,7 @@ 
TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_32x32x32_fp8_fp8, "V16fV2iV4iV16fiIiI
 
 TARGET_BUILTIN(__builtin_amdgcn_cvt_f32_bf8, "fiIi", "nc", 
"fp8-conversion-insts")
 TARGET_BUILTIN(__builtin_amdgcn_cvt_f32_fp8, "fiIi", "nc", 
"fp8-conversion-insts")
+TARGET_BUILTIN(__builtin_amdgcn_cvt_f32_fp8_e5m3, "fiIi", "nc", 
"fp8e5m3-insts")
 TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_f32_bf8, "V2fiIb", "nc", 
"fp8-conversion-insts")
 TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_f32_fp8, "V2fiIb", "nc", 
"fp8-conversion-insts")
 TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_bf8_f32, "iffiIb", "nc", 
"fp8-conversion-insts")
diff --git a/clang/test/CodeGenOpenCL/amdgpu-features.cl 
b/clang/test/CodeGenOpenCL/amdgpu-features.cl
index dc7a83002b7f1..77d2414230cf2 100644
--- a/clang/test/CodeGenOpenCL/amdgpu-features.cl
+++ b/clang/test/CodeGenOpenCL/amdgpu-features.cl
@@ -108,7 +108,7 @@
 // GFX1153: 
"target-features"="+16-bit-insts,+atomic-fadd-rtn-insts,+ci-insts,+dl-insts,+dot10-insts,+dot12-insts,+dot5-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx8-insts,+gfx9-insts,+wavefrontsize32"
 // GFX1200: 
"target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+ci-insts,+dl-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+fp8-conversion-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx8-insts,+gfx9-insts,+wavefrontsize32"
 // GFX1201: 
"target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+ci-insts,+dl-insts,+dot10-insts,+dot11-insts,+dot12-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+fp8-conversion-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx8-insts,+gfx9-insts,+wavefrontsize32"
-// GFX1250: 
"target-features"="+16-bit-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+bitop3-insts,+ci-insts,+dl-insts,+dot7-insts,+dot8-insts,+dpp,+fp8-conversion-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx8-insts,+gfx9-insts,+permlane16-swap,+prng-inst,+setprio-inc-wg-inst,+transpose-load-f4f6-insts,+wavefrontsize32"
+// GFX1250: 
"target-features"="+16-bit-insts,+ashr-pk-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-buffer-pk-add-bf16-inst,+atomic-ds-pk-add-16-insts,+atomic-fadd-rtn-insts,+atomic-flat-pk-add-16-insts,+atomic-global-pk-add-bf16-inst,+bitop3-insts,+ci-insts,+dl-insts,+dot7-insts,+dot8-insts,+dpp,+fp8-conversion-insts,+fp8e5m3-insts,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx12-insts,+gfx1250-insts,+gfx8-insts,+gfx9-insts,+permlane16-swap,+prng-inst,+setprio-inc-wg-inst,+transpose-load-f4f6-insts,+wavefrontsize32"
 
 // GFX1103-W64: 
"target-features"="+16-bit-insts,+atomic-fadd-rtn-insts,+ci-insts,+dl-insts,+dot10-insts,+dot12-insts,+dot5-insts,+dot7-insts,+dot8-insts,+dot9-insts,+dpp,+gfx10-3-insts,+gfx10-insts,+gfx11-insts,+gfx8-insts,+gfx9-insts,+wavefrontsize64"
 
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl 
b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl
index 4d4afedae3658..421099d3876e3 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl
@@ -139,3 +139,41 @@ void test_cvt_pk_f16_bf8(global half2* out, short a)
 {
   out[0] = __builtin_amdgcn_cvt_pk_f16_bf8(a);
 }
+
+// CHECK-LABEL: @test_cvt_f32_fp8_e5m3(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, 
addrspace(5)
+// CHECK-NEXT:    [[A_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT:    [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) 
[[OUT_ADDR]] to ptr
+// CHECK-NEXT:    [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) 
[[A_ADDR]] to ptr
+// CHECK-NEXT:    store ptr addrspace(1) [[OUT:%.*]], ptr [[OUT_ADDR_ASCAST]], 
align 8
+// CHECK-NEXT:    store i32 [[A:%.*]], ptr [[A_ADDR_ASCAST]], align 4
+// CHECK-NEXT:    [[TMP0:%.*]] = load i32, ptr [[A_ADDR_ASCAST]], align 4
+// CHECK-NEXT:    [[TMP1:%.*]] = call float @llvm.amdgcn.cvt.f32.fp8.e5m3(i32 
[[TMP0]], i32 0)
+// CHECK-NEXT:    [[CONV:%.*]] = fptosi float [[TMP1]] to i32
+// CHECK-NEXT:    [[TMP2:%.*]] = load ptr addrspace(1), ptr 
[[OUT_ADDR_ASCAST]], align 8
+// CHECK-NEXT:    store i32 [[CONV]], ptr addrspace(1) [[TMP2]], align 4
+// CHECK-NEXT:    [[TMP3:%.*]] = load i32, ptr [[A_ADDR_ASCAST]], align 4
+// CHECK-NEXT:    [[TMP4:%.*]] = call float @llvm.amdgcn.cvt.f32.fp8.e5m3(i32 
[[TMP3]], i32 1)
+// CHECK-NEXT:    [[CONV1:%.*]] = fptosi float [[TMP4]] to i32
+// CHECK-NEXT:    [[TMP5:%.*]] = load ptr addrspace(1), ptr 
[[OUT_ADDR_ASCAST]], align 8
+// CHECK-NEXT:    store i32 [[CONV1]], ptr addrspace(1) [[TMP5]], align 4
+// CHECK-NEXT:    [[TMP6:%.*]] = load i32, ptr [[A_ADDR_ASCAST]], align 4
+// CHECK-NEXT:    [[TMP7:%.*]] = call float @llvm.amdgcn.cvt.f32.fp8.e5m3(i32 
[[TMP6]], i32 2)
+// CHECK-NEXT:    [[CONV2:%.*]] = fptosi float [[TMP7]] to i32
+// CHECK-NEXT:    [[TMP8:%.*]] = load ptr addrspace(1), ptr 
[[OUT_ADDR_ASCAST]], align 8
+// CHECK-NEXT:    store i32 [[CONV2]], ptr addrspace(1) [[TMP8]], align 4
+// CHECK-NEXT:    [[TMP9:%.*]] = load i32, ptr [[A_ADDR_ASCAST]], align 4
+// CHECK-NEXT:    [[TMP10:%.*]] = call float @llvm.amdgcn.cvt.f32.fp8.e5m3(i32 
[[TMP9]], i32 3)
+// CHECK-NEXT:    [[CONV3:%.*]] = fptosi float [[TMP10]] to i32
+// CHECK-NEXT:    [[TMP11:%.*]] = load ptr addrspace(1), ptr 
[[OUT_ADDR_ASCAST]], align 8
+// CHECK-NEXT:    store i32 [[CONV3]], ptr addrspace(1) [[TMP11]], align 4
+// CHECK-NEXT:    ret void
+//
+void test_cvt_f32_fp8_e5m3(global int* out, int a)
+{
+  *out = __builtin_amdgcn_cvt_f32_fp8_e5m3(a, 0);
+  *out = __builtin_amdgcn_cvt_f32_fp8_e5m3(a, 1);
+  *out = __builtin_amdgcn_cvt_f32_fp8_e5m3(a, 2);
+  *out = __builtin_amdgcn_cvt_f32_fp8_e5m3(a, 3);
+}
diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl 
b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl
index 3ba0d50e79031..7494c4f984353 100644
--- a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl
@@ -27,3 +27,8 @@ void test_amdgcn_tensor_load_store(v4i sg0, v8i sg1, v4i sg2, 
v4i sg3, int cpol)
   __builtin_amdgcn_tensor_store_from_lds(sg0, sg1, sg2, sg3, cpol); // 
expected-error {{'__builtin_amdgcn_tensor_store_from_lds' must be a constant 
integer}}
   __builtin_amdgcn_tensor_store_from_lds_d2(sg0, sg1, cpol); // expected-error 
{{'__builtin_amdgcn_tensor_store_from_lds_d2' must be a constant integer}}
 }
+
+void test_cvt_f32_fp8_e5m3(global int* out, int a)
+{
+  *out = __builtin_amdgcn_cvt_f32_fp8_e5m3(a, a); // expected-error 
{{'__builtin_amdgcn_cvt_f32_fp8_e5m3' must be a constant integer}}
+}
diff --git a/lldb/source/Host/windows/MainLoopWindows.cpp 
b/lldb/source/Host/windows/MainLoopWindows.cpp
index b3322e8b3ae59..a1de895c0ba98 100644
--- a/lldb/source/Host/windows/MainLoopWindows.cpp
+++ b/lldb/source/Host/windows/MainLoopWindows.cpp
@@ -12,16 +12,16 @@
 #include "lldb/Host/windows/windows.h"
 #include "lldb/Utility/Status.h"
 #include "llvm/Config/llvm-config.h"
-#include "llvm/Support/Casting.h"
 #include "llvm/Support/WindowsError.h"
 #include <algorithm>
 #include <cassert>
-#include <cerrno>
-#include <csignal>
 #include <ctime>
 #include <io.h>
+#include <synchapi.h>
 #include <thread>
 #include <vector>
+#include <winbase.h>
+#include <winerror.h>
 #include <winsock2.h>
 
 using namespace lldb;
@@ -42,11 +42,12 @@ namespace {
 class PipeEvent : public MainLoopWindows::IOEvent {
 public:
   explicit PipeEvent(HANDLE handle)
-      : IOEvent(CreateEventW(NULL, /*bManualReset=*/FALSE,
+      : IOEvent(CreateEventW(NULL, /*bManualReset=*/TRUE,
                              /*bInitialState=*/FALSE, NULL)),
-        m_handle(handle), m_ready(CreateEventW(NULL, /*bManualReset=*/FALSE,
+        m_handle(handle), m_ready(CreateEventW(NULL, /*bManualReset=*/TRUE,
                                                /*bInitialState=*/FALSE, NULL)) 
{
     assert(m_event && m_ready);
+    m_monitor_thread = std::thread(&PipeEvent::Monitor, this);
   }
 
   ~PipeEvent() override {
@@ -65,15 +66,27 @@ class PipeEvent : public MainLoopWindows::IOEvent {
   }
 
   void WillPoll() override {
-    if (!m_monitor_thread.joinable())
-      m_monitor_thread = std::thread(&PipeEvent::Monitor, this);
+    if (WaitForSingleObject(m_event, /*dwMilliseconds=*/0) != WAIT_TIMEOUT) {
+      // The thread has already signalled that the data is available. No need
+      // for further polling until we consume that event.
+      return;
+    }
+    if (WaitForSingleObject(m_ready, /*dwMilliseconds=*/0) != WAIT_TIMEOUT) {
+      // The thread is already waiting for data to become available.
+      return;
+    }
+    // Start waiting.
+    SetEvent(m_ready);
   }
 
-  void Disarm() override { SetEvent(m_ready); }
+  void Disarm() override { ResetEvent(m_event); }
 
   /// Monitors the handle performing a zero byte read to determine when data is
   /// avaiable.
   void Monitor() {
+    // Wait until the MainLoop tells us to start.
+    WaitForSingleObject(m_ready, INFINITE);
+
     do {
       char buf[1];
       DWORD bytes_read = 0;
@@ -110,7 +123,11 @@ class PipeEvent : public MainLoopWindows::IOEvent {
         continue;
       }
 
+      // Notify that data is available on the pipe. It's important to set this
+      // before clearing m_ready to avoid a race with WillPoll.
       SetEvent(m_event);
+      // Stop polling until we're told to resume.
+      ResetEvent(m_ready);
 
       // Wait until the current read is consumed before doing the next read.
       WaitForSingleObject(m_ready, INFINITE);
diff --git a/lldb/unittests/Host/MainLoopTest.cpp 
b/lldb/unittests/Host/MainLoopTest.cpp
index 502028ae1a343..30585d12fe81d 100644
--- a/lldb/unittests/Host/MainLoopTest.cpp
+++ b/lldb/unittests/Host/MainLoopTest.cpp
@@ -10,6 +10,7 @@
 #include "TestingSupport/SubsystemRAII.h"
 #include "lldb/Host/ConnectionFileDescriptor.h"
 #include "lldb/Host/FileSystem.h"
+#include "lldb/Host/MainLoopBase.h"
 #include "lldb/Host/PseudoTerminal.h"
 #include "lldb/Host/common/TCPSocket.h"
 #include "llvm/Config/llvm-config.h" // for LLVM_ON_UNIX
@@ -64,7 +65,7 @@ class MainLoopTest : public testing::Test {
 };
 } // namespace
 
-TEST_F(MainLoopTest, ReadObject) {
+TEST_F(MainLoopTest, ReadSocketObject) {
   char X = 'X';
   size_t len = sizeof(X);
   ASSERT_TRUE(socketpair[0]->Write(&X, len).Success());
@@ -101,6 +102,144 @@ TEST_F(MainLoopTest, ReadPipeObject) {
   ASSERT_EQ(1u, callback_count);
 }
 
+TEST_F(MainLoopTest, MultipleReadsPipeObject) {
+  Pipe pipe;
+
+  ASSERT_TRUE(pipe.CreateNew().Success());
+
+  MainLoop loop;
+
+  std::future<void> async_writer = std::async(std::launch::async, [&] {
+    for (int i = 0; i < 5; ++i) {
+      std::this_thread::sleep_for(std::chrono::milliseconds(500));
+      char X = 'X';
+      size_t len = sizeof(X);
+      ASSERT_THAT_EXPECTED(pipe.Write(&X, len), llvm::HasValue(1));
+    }
+  });
+
+  Status error;
+  lldb::FileSP file = std::make_shared<NativeFile>(
+      pipe.GetReadFileDescriptor(), File::eOpenOptionReadOnly, false);
+  auto handle = loop.RegisterReadObject(
+      file,
+      [&](MainLoopBase &loop) {
+        callback_count++;
+        if (callback_count == 5)
+          loop.RequestTermination();
+
+        // Read some data to ensure the handle is not in a readable state.
+        char buf[1024] = {0};
+        size_t len = sizeof(buf);
+        ASSERT_THAT_ERROR(file->Read(buf, len).ToError(), llvm::Succeeded());
+        EXPECT_EQ(len, 1);
+        EXPECT_EQ(buf[0], 'X');
+      },
+      error);
+  ASSERT_TRUE(error.Success());
+  ASSERT_TRUE(handle);
+  ASSERT_TRUE(loop.Run().Success());
+  ASSERT_EQ(5u, callback_count);
+  async_writer.wait();
+}
+
+TEST_F(MainLoopTest, PipeDelayBetweenRegisterAndRun) {
+  Pipe pipe;
+
+  ASSERT_TRUE(pipe.CreateNew().Success());
+
+  MainLoop loop;
+
+  Status error;
+  lldb::FileSP file = std::make_shared<NativeFile>(
+      pipe.GetReadFileDescriptor(), File::eOpenOptionReadOnly, false);
+  auto handle = loop.RegisterReadObject(
+      file,
+      [&](MainLoopBase &loop) {
+        callback_count++;
+
+        // Read some data to ensure the handle is not in a readable state.
+        char buf[1024] = {0};
+        size_t len = sizeof(buf);
+        ASSERT_THAT_ERROR(file->Read(buf, len).ToError(), llvm::Succeeded());
+        EXPECT_EQ(len, 2);
+        EXPECT_EQ(buf[0], 'X');
+        EXPECT_EQ(buf[1], 'X');
+      },
+      error);
+  auto cb = [&](MainLoopBase &) {
+    callback_count++;
+    char X = 'X';
+    size_t len = sizeof(X);
+    // Write twice and ensure we coalesce into a single read.
+    ASSERT_THAT_EXPECTED(pipe.Write(&X, len), llvm::HasValue(1));
+    ASSERT_THAT_EXPECTED(pipe.Write(&X, len), llvm::HasValue(1));
+  };
+  // Add a write that triggers a read events.
+  loop.AddCallback(cb, std::chrono::milliseconds(500));
+  loop.AddCallback([](MainLoopBase &loop) { loop.RequestTermination(); },
+                   std::chrono::milliseconds(1000));
+  ASSERT_TRUE(error.Success());
+  ASSERT_TRUE(handle);
+
+  // Write between RegisterReadObject / Run should NOT invoke the callback.
+  cb(loop);
+  ASSERT_EQ(1u, callback_count);
+
+  ASSERT_TRUE(loop.Run().Success());
+  ASSERT_EQ(4u, callback_count);
+}
+
+TEST_F(MainLoopTest, NoSelfTriggersDuringPipeHandler) {
+  Pipe pipe;
+
+  ASSERT_TRUE(pipe.CreateNew().Success());
+
+  MainLoop loop;
+
+  Status error;
+  lldb::FileSP file = std::make_shared<NativeFile>(
+      pipe.GetReadFileDescriptor(), File::eOpenOptionReadOnly, false);
+  auto handle = loop.RegisterReadObject(
+      file,
+      [&](MainLoopBase &lop) {
+        callback_count++;
+
+        char X = 'Y';
+        size_t len = sizeof(X);
+        // writes / reads during the handler callback should NOT trigger 
itself.
+        ASSERT_THAT_EXPECTED(pipe.Write(&X, len), llvm::HasValue(1));
+
+        char buf[1024] = {0};
+        len = sizeof(buf);
+        ASSERT_THAT_ERROR(file->Read(buf, len).ToError(), llvm::Succeeded());
+        EXPECT_EQ(len, 2);
+        EXPECT_EQ(buf[0], 'X');
+        EXPECT_EQ(buf[1], 'Y');
+
+        if (callback_count == 2)
+          loop.RequestTermination();
+      },
+      error);
+  // Add a write that triggers a read event.
+  loop.AddPendingCallback([&](MainLoopBase &) {
+    char X = 'X';
+    size_t len = sizeof(X);
+    ASSERT_THAT_EXPECTED(pipe.Write(&X, len), llvm::HasValue(1));
+  });
+  loop.AddCallback(
+      [&](MainLoopBase &) {
+        char X = 'X';
+        size_t len = sizeof(X);
+        ASSERT_THAT_EXPECTED(pipe.Write(&X, len), llvm::HasValue(1));
+      },
+      std::chrono::milliseconds(500));
+  ASSERT_TRUE(error.Success());
+  ASSERT_TRUE(handle);
+  ASSERT_TRUE(loop.Run().Success());
+  ASSERT_EQ(2u, callback_count);
+}
+
 TEST_F(MainLoopTest, NoSpuriousPipeReads) {
   Pipe pipe;
 
diff --git a/llvm/include/llvm/Analysis/TargetTransformInfo.h 
b/llvm/include/llvm/Analysis/TargetTransformInfo.h
index c43870392361d..98b793aace7a3 100644
--- a/llvm/include/llvm/Analysis/TargetTransformInfo.h
+++ b/llvm/include/llvm/Analysis/TargetTransformInfo.h
@@ -1702,12 +1702,13 @@ class TargetTransformInfo {
   /// unordered-atomic memory intrinsic.
   LLVM_ABI unsigned getAtomicMemIntrinsicMaxElementSize() const;
 
-  /// \returns A value which is the result of the given memory intrinsic.  New
-  /// instructions may be created to extract the result from the given 
intrinsic
-  /// memory operation.  Returns nullptr if the target cannot create a result
-  /// from the given intrinsic.
-  LLVM_ABI Value *getOrCreateResultFromMemIntrinsic(IntrinsicInst *Inst,
-                                                    Type *ExpectedType) const;
+  /// \returns A value which is the result of the given memory intrinsic. If \p
+  /// CanCreate is true, new instructions may be created to extract the result
+  /// from the given intrinsic memory operation. Returns nullptr if the target
+  /// cannot create a result from the given intrinsic.
+  LLVM_ABI Value *
+  getOrCreateResultFromMemIntrinsic(IntrinsicInst *Inst, Type *ExpectedType,
+                                    bool CanCreate = true) const;
 
   /// \returns The type to use in a loop expansion of a memcpy call.
   LLVM_ABI Type *getMemcpyLoopLoweringType(
diff --git a/llvm/include/llvm/Analysis/TargetTransformInfoImpl.h 
b/llvm/include/llvm/Analysis/TargetTransformInfoImpl.h
index 12f87226c5f57..ddc8a5eaffa94 100644
--- a/llvm/include/llvm/Analysis/TargetTransformInfoImpl.h
+++ b/llvm/include/llvm/Analysis/TargetTransformInfoImpl.h
@@ -983,8 +983,9 @@ class TargetTransformInfoImplBase {
     return 0;
   }
 
-  virtual Value *getOrCreateResultFromMemIntrinsic(IntrinsicInst *Inst,
-                                                   Type *ExpectedType) const {
+  virtual Value *
+  getOrCreateResultFromMemIntrinsic(IntrinsicInst *Inst, Type *ExpectedType,
+                                    bool CanCreate = true) const {
     return nullptr;
   }
 
diff --git a/llvm/include/llvm/CodeGen/SDPatternMatch.h 
b/llvm/include/llvm/CodeGen/SDPatternMatch.h
index 7c5cdbbeb0ca8..2967532226197 100644
--- a/llvm/include/llvm/CodeGen/SDPatternMatch.h
+++ b/llvm/include/llvm/CodeGen/SDPatternMatch.h
@@ -655,6 +655,21 @@ struct MaxMin_match {
 
   template <typename MatchContext>
   bool match(const MatchContext &Ctx, SDValue N) {
+    auto MatchMinMax = [&](SDValue L, SDValue R, SDValue TrueValue,
+                           SDValue FalseValue, ISD::CondCode CC) {
+      if ((TrueValue != L || FalseValue != R) &&
+          (TrueValue != R || FalseValue != L))
+        return false;
+
+      ISD::CondCode Cond =
+          TrueValue == L ? CC : getSetCCInverse(CC, L.getValueType());
+      if (!Pred_t::match(Cond))
+        return false;
+
+      return (LHS.match(Ctx, L) && RHS.match(Ctx, R)) ||
+             (Commutable && LHS.match(Ctx, R) && RHS.match(Ctx, L));
+    };
+
     if (sd_context_match(N, Ctx, m_Opc(ISD::SELECT)) ||
         sd_context_match(N, Ctx, m_Opc(ISD::VSELECT))) {
       EffectiveOperands<ExcludeChain> EO_SELECT(N, Ctx);
@@ -670,23 +685,22 @@ struct MaxMin_match {
         SDValue R = Cond->getOperand(EO_SETCC.FirstIndex + 1);
         auto *CondNode =
             cast<CondCodeSDNode>(Cond->getOperand(EO_SETCC.FirstIndex + 2));
-
-        if ((TrueValue != L || FalseValue != R) &&
-            (TrueValue != R || FalseValue != L)) {
-          return false;
-        }
-
-        ISD::CondCode Cond =
-            TrueValue == L ? CondNode->get()
-                           : getSetCCInverse(CondNode->get(), 
L.getValueType());
-        if (!Pred_t::match(Cond)) {
-          return false;
-        }
-        return (LHS.match(Ctx, L) && RHS.match(Ctx, R)) ||
-               (Commutable && LHS.match(Ctx, R) && RHS.match(Ctx, L));
+        return MatchMinMax(L, R, TrueValue, FalseValue, CondNode->get());
       }
     }
 
+    if (sd_context_match(N, Ctx, m_Opc(ISD::SELECT_CC))) {
+      EffectiveOperands<ExcludeChain> EO_SELECT(N, Ctx);
+      assert(EO_SELECT.Size == 5);
+      SDValue L = N->getOperand(EO_SELECT.FirstIndex);
+      SDValue R = N->getOperand(EO_SELECT.FirstIndex + 1);
+      SDValue TrueValue = N->getOperand(EO_SELECT.FirstIndex + 2);
+      SDValue FalseValue = N->getOperand(EO_SELECT.FirstIndex + 3);
+      auto *CondNode =
+          cast<CondCodeSDNode>(N->getOperand(EO_SELECT.FirstIndex + 4));
+      return MatchMinMax(L, R, TrueValue, FalseValue, CondNode->get());
+    }
+
     return false;
   }
 };
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td 
b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index f7a9b65854696..16885f331e9dd 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@...
[truncated]

``````````

</details>


https://github.com/llvm/llvm-project/pull/147726
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits

Reply via email to