diff --git a/WORKSPACE b/WORKSPACE
index e970ed993..88d764529 100644
--- a/WORKSPACE
+++ b/WORKSPACE
@@ -26,9 +26,9 @@ http_archive(
     ],
 )
 
-LLVM_COMMIT = "f2d32ddcec82c20582c6aa32558b82ca7c3d3c50"
+LLVM_COMMIT = "913f21ae5c460d6fdd73ac7663534e73f8e19044"
 
-LLVM_SHA256 = "f209ad8ddec7debd9fd20d8d0800d3c210eb6ed9b29630506fdbc8301e46d587"
+LLVM_SHA256 = "95b4859b7bc1bf4648f6718100dc19b7ffdcfad95f79c19491f36c2a58dc2104"
 
 http_archive(
     name = "llvm-raw",
diff --git a/build_tools/llvm_version.txt b/build_tools/llvm_version.txt
index 93ef0d055..a742d03f7 100644
--- a/build_tools/llvm_version.txt
+++ b/build_tools/llvm_version.txt
@@ -1,2 +1,2 @@
-f2d32ddcec82c20582c6aa32558b82ca7c3d3c50
+913f21ae5c460d6fdd73ac7663534e73f8e19044
 
diff --git a/gml_st/IR/gml_st_ops.cc b/gml_st/IR/gml_st_ops.cc
index 57210d920..8da248b5b 100644
--- a/gml_st/IR/gml_st_ops.cc
+++ b/gml_st/IR/gml_st_ops.cc
@@ -166,7 +166,7 @@ ParseResult FusionOp::parse(OpAsmParser &parser, OperationState &result) {
   if (parser.parseOptionalColonTypeList(result.types)) return failure();
 
   result.addAttribute(
-      "operand_segment_sizes",
+      "operandSegmentSizes",
       parser.getBuilder().getDenseI32ArrayAttr(
           {static_cast<int32_t>(numInputs), static_cast<int32_t>(numInits)}));
 
diff --git a/lhlo/IR/lhlo_ops.cc b/lhlo/IR/lhlo_ops.cc
index a59c565c9..cbf8c9df0 100644
--- a/lhlo/IR/lhlo_ops.cc
+++ b/lhlo/IR/lhlo_ops.cc
@@ -122,7 +122,6 @@ LogicalResult AbsOp::verify() {
 //===----------------------------------------------------------------------===//
 
 void CaseOp::getSuccessorRegions(std::optional<unsigned> index,
-                                 ArrayRef<Attribute> /*operands*/,
                                  SmallVectorImpl<RegionSuccessor>& regions) {
   // If the predecessor is the CaseOp, branch to all other branches.
   if (!index.has_value()) {
@@ -350,7 +349,6 @@ LogicalResult ReduceWindowOp::verify() {
 //===----------------------------------------------------------------------===//
 
 void WhileOp::getSuccessorRegions(std::optional<unsigned> index,
-                                  ArrayRef<Attribute> /*operands*/,
                                   SmallVectorImpl<RegionSuccessor>& regions) {
   // If the predecessor is the WhileOp or the body region, branch into the
   // cond region.
@@ -390,7 +388,6 @@ void FusionOp::build(OpBuilder& builder, OperationState& result,
 }
 
 void FusionOp::getSuccessorRegions(std::optional<unsigned> index,
-                                   ArrayRef<Attribute> /*operands*/,
                                    SmallVectorImpl<RegionSuccessor>& regions) {
   // If the predecessor is the fusion region, jump back to the parent op.
   if (index.has_value()) {
diff --git a/stablehlo/build_tools/llvm_version.txt b/stablehlo/build_tools/llvm_version.txt
index 63132f905..97da0f58e 100644
--- a/stablehlo/build_tools/llvm_version.txt
+++ b/stablehlo/build_tools/llvm_version.txt
@@ -1 +1 @@
-f2d32ddcec82c20582c6aa32558b82ca7c3d3c50
+913f21ae5c460d6fdd73ac7663534e73f8e19044
diff --git a/stablehlo/stablehlo/dialect/ChloOps.cpp b/stablehlo/stablehlo/dialect/ChloOps.cpp
index ee328bc20..0389353b5 100644
--- a/stablehlo/stablehlo/dialect/ChloOps.cpp
+++ b/stablehlo/stablehlo/dialect/ChloOps.cpp
@@ -416,8 +416,7 @@ LogicalResult BroadcastSelectOp::reifyReturnTypeShapes(
 //===----------------------------------------------------------------------===//
 
 void RankSpecializationClusterOp::getSuccessorRegions(
-    std::optional<unsigned> index, ArrayRef<Attribute> /*operands*/,
-    SmallVectorImpl<RegionSuccessor>& regions) {
+    std::optional<unsigned> index, SmallVectorImpl<RegionSuccessor>& regions) {
   // RankSpecializationClusterOp has unconditional control flows into the region
   // and back to the parent, so return the correct RegionSuccessor purely based
   // on the index being None or 0.
diff --git a/tests/Dialect/gml_st/simplify_dead_copy.mlir b/tests/Dialect/gml_st/simplify_dead_copy.mlir
index 48d295fe9..bec5a0822 100644
--- a/tests/Dialect/gml_st/simplify_dead_copy.mlir
+++ b/tests/Dialect/gml_st/simplify_dead_copy.mlir
@@ -183,7 +183,7 @@ func.func @target_is_alloc_with_loads_stores(%arg0: memref<8x8xf32>)
     backend_config = "",
     call_target_name = "foo",
     has_side_effect = false,
-    operand_segment_sizes = array<i32: 1, 1>
+    operandSegmentSizes = array<i32: 1, 1>
   } : (memref<8x8xf32>, memref<8x8xf32>) -> ()
 
   return %arg0 : memref<8x8xf32>
diff --git a/tests/Dialect/lhlo/ops.mlir b/tests/Dialect/lhlo/ops.mlir
index eb92b90d5..8be889f81 100644
--- a/tests/Dialect/lhlo/ops.mlir
+++ b/tests/Dialect/lhlo/ops.mlir
@@ -929,7 +929,7 @@ func.func @valid_custom_call(%arg0:memref<1xf32>, %arg1:memref<1xf32>) -> () {
     backend_config = "",
     call_target_name = "foo",
     has_side_effects = false,
-    operand_segment_sizes = array<i32: 2, 2>,
+    operandSegmentSizes = array<i32: 2, 2>,
     target_arg_mapping = #lmhlo.custom_call_target_arg_mapping<
       num_args = 4,
       num_results = 3,
@@ -948,7 +948,7 @@ func.func @invalid_custom_call(%arg0:memref<1xf32>, %arg1:memref<1xf32>) -> () {
     backend_config = "",
     call_target_name = "foo",
     has_side_effects = false,
-    operand_segment_sizes = array<i32: 2, 2>,
+    operandSegmentSizes = array<i32: 2, 2>,
     target_arg_mapping = #lmhlo.custom_call_target_arg_mapping<
       num_args = 4,
       num_results = 3,
@@ -967,7 +967,7 @@ func.func @invalid_custom_call(%arg0:memref<1xf32>, %arg1:memref<1xf32>) -> () {
     backend_config = "",
     call_target_name = "foo",
     has_side_effects = false,
-    operand_segment_sizes = array<i32: 2, 2>,
+    operandSegmentSizes = array<i32: 2, 2>,
     target_arg_mapping = #lmhlo.custom_call_target_arg_mapping<
       num_args = 4,
       num_results = 3,
@@ -986,7 +986,7 @@ func.func @invalid_custom_call(%arg0:memref<1xf32>, %arg1:memref<1xf32>) -> () {
     backend_config = "",
     call_target_name = "foo",
     has_side_effects = false,
-    operand_segment_sizes = array<i32: 2, 2>,
+    operandSegmentSizes = array<i32: 2, 2>,
     target_arg_mapping = #lmhlo.custom_call_target_arg_mapping<
       num_args = 4,
       num_results = 3,
@@ -1005,7 +1005,7 @@ func.func @invalid_custom_call(%arg0:memref<1xf32>, %arg1:memref<1xf32>) -> () {
     backend_config = "",
     call_target_name = "foo",
     has_side_effects = false,
-    operand_segment_sizes = array<i32: 2, 2>,
+    operandSegmentSizes = array<i32: 2, 2>,
     target_arg_mapping = #lmhlo.custom_call_target_arg_mapping<
       num_args = 4,
       num_results = 3,
@@ -1024,7 +1024,7 @@ func.func @invalid_custom_call(%arg0:memref<1xf32>, %arg1:memref<1xf32>) -> () {
     backend_config = "",
     call_target_name = "foo",
     has_side_effects = false,
-    operand_segment_sizes = array<i32: 2, 2>,
+    operandSegmentSizes = array<i32: 2, 2>,
     target_arg_mapping = #lmhlo.custom_call_target_arg_mapping<
       num_args = 4,
       num_results = 3,
@@ -1043,7 +1043,7 @@ func.func @invalid_custom_call(%arg0:memref<1xf32>, %arg1:memref<1xf32>) -> () {
     backend_config = "",
     call_target_name = "foo",
     has_side_effects = false,
-    operand_segment_sizes = array<i32: 2, 2>,
+    operandSegmentSizes = array<i32: 2, 2>,
     target_arg_mapping = #lmhlo.custom_call_target_arg_mapping<
       num_args = 4,
       num_results = 3,
@@ -1129,4 +1129,4 @@ func.func @recv_done(%arg0: !mhlo.token) {
     is_host_transfer = true
   } : (!mhlo.token) -> ()
   return
-}
\ No newline at end of file
+}
diff --git a/tests/Dialect/mhlo/hlo-legalize-to-lhlo.mlir b/tests/Dialect/mhlo/hlo-legalize-to-lhlo.mlir
index afae8cf8a..08866112d 100644
--- a/tests/Dialect/mhlo/hlo-legalize-to-lhlo.mlir
+++ b/tests/Dialect/mhlo/hlo-legalize-to-lhlo.mlir
@@ -620,7 +620,7 @@ func.func @transpose(%operand: tensor<2x2xf32>) -> tensor<2x2xf32> {
 // CHECK-SAME:([[ARG0:%.*]]: memref<2x2xf32>, [[ARG1:%.*]]: memref<2x3xf32>)
 func.func @custom_call(%arg0: tensor<2x2xf32>, %arg1: tensor<2x3xf32>) -> tensor<4x4xf16> {
   // CHECK: "lmhlo.custom_call"([[ARG0]], [[ARG1]], %{{.*}}) ({
-  // CHECK-NEXT: }) {backend_config = "", call_target_name = "foo", has_side_effect = false, operand_segment_sizes = array<i32: 2, 1>}
+  // CHECK-NEXT: }) {backend_config = "", call_target_name = "foo", has_side_effect = false, operandSegmentSizes = array<i32: 2, 1>}
   %result = "mhlo.custom_call"(%arg0, %arg1)
               {backend_config = "", call_target_name = "foo", has_side_effect = false}
               : (tensor<2x2xf32>, tensor<2x3xf32>) -> tensor<4x4xf16>
@@ -633,7 +633,7 @@ func.func @custom_call(%arg0: tensor<2x2xf32>, %arg1: tensor<2x3xf32>) -> tensor
 // CHECK-SAME:([[ARG0:%.*]]: memref<2x2xf32>, [[ARG1:%.*]]: memref<2x3xf32>)
 func.func @custom_call_multiout(%arg0: tensor<2x2xf32>, %arg1: tensor<2x3xf32>) -> tensor<4x4xf16> {
   // CHECK: "lmhlo.custom_call"([[ARG0]], [[ARG1]], %{{.*}}, %{{.*}}) ({
-  // CHECK-NEXT: }) {backend_config = "", call_target_name = "foo", has_side_effect = false, operand_segment_sizes = array<i32: 2, 2>}
+  // CHECK-NEXT: }) {backend_config = "", call_target_name = "foo", has_side_effect = false, operandSegmentSizes = array<i32: 2, 2>}
   %temp:2 = "mhlo.custom_call"(%arg0, %arg1)
                    {backend_config = "", call_target_name = "foo", has_side_effect = false}
                    : (tensor<2x2xf32>, tensor<2x3xf32>) -> (tensor<4x4xf16>, tensor<4x4xf16>)
diff --git a/tests/Dialect/mhlo/hlo-legalize-to-memref.mlir b/tests/Dialect/mhlo/hlo-legalize-to-memref.mlir
index bc8dc3b84..bbbc30d29 100644
--- a/tests/Dialect/mhlo/hlo-legalize-to-memref.mlir
+++ b/tests/Dialect/mhlo/hlo-legalize-to-memref.mlir
@@ -122,7 +122,7 @@ func.func @custom_call_multiple_inputs_outputs(%x: tensor<2xf32>,
 // CHECK-DAG: %[[O1:.*]] = memref.alloc() {{.*}} : memref<2xf32>
 // CHECK-DAG: %[[O2:.*]] = memref.alloc() {{.*}} : memref<5xi32>
 // CHECK: "lmhlo.custom_call"(%[[I0]], %[[I1]], %[[O0]], %[[O1]], %[[O2]]) ({
-// CHECK-NEXT: }) {backend_config = "", call_target_name = "foo", has_side_effect = false, operand_segment_sizes = array<i32: 2, 3>} : (memref<2xf32>, memref<5xi32>, memref<2xf32>, memref<2xf32>, memref<5xi32>) -> ()
+// CHECK-NEXT: }) {backend_config = "", call_target_name = "foo", has_side_effect = false, operandSegmentSizes = array<i32: 2, 3>} : (memref<2xf32>, memref<5xi32>, memref<2xf32>, memref<2xf32>, memref<5xi32>) -> ()
 // CHECK-DAG: %[[T0:.+]] = bufferization.to_tensor %[[O0]] : memref<2xf32>
 // CHECK-DAG: %[[T1:.+]] = bufferization.to_tensor %[[O1]] : memref<2xf32>
 // CHECK: return %[[T0]], %[[T1]] : tensor<2xf32>, tensor<2xf32>
@@ -149,5 +149,5 @@ func.func @custom_call_side_effect(%x: tensor<2xf32>,
 // CHECK-DAG: %[[I1:.+]] = bufferization.to_memref %[[ARG1]] : memref<5xi32>
 // CHECK-DAG: %[[ALLOC:.+]] = memref.alloc
 // CHECK: "lmhlo.custom_call"(%[[I0]], %[[I1]], %[[ALLOC]]) ({
-// CHECK-NEXT: }) {backend_config = "", call_target_name = "bar", has_side_effect = true, operand_segment_sizes = array<i32: 2, 1>, target_arg_mapping = #lmhlo.custom_call_target_arg_mapping<num_args = 3, num_results = 2, args_to_target_args = [0, 1], results_to_target_results = [1]>} : (memref<2xf32>, memref<5xi32>, memref<2xi32>)
+// CHECK-NEXT: }) {backend_config = "", call_target_name = "bar", has_side_effect = true, operandSegmentSizes = array<i32: 2, 1>, target_arg_mapping = #lmhlo.custom_call_target_arg_mapping<num_args = 3, num_results = 2, args_to_target_args = [0, 1], results_to_target_results = [1]>} : (memref<2xf32>, memref<5xi32>, memref<2xi32>)
 // CHECK: return %[[TOKEN]] : !mhlo.token