diff --git a/build_tools/llvm_version.txt b/build_tools/llvm_version.txt
index 5d5aafef..7cd9bb4f 100644
--- a/build_tools/llvm_version.txt
+++ b/build_tools/llvm_version.txt
@@ -1 +1 @@
-6ae7b735dbd50eb7ade1573a86d037a2943e679c
+511ba45a47d6f9e48ad364181830c9fb974135b2
diff --git a/include/emitc/Conversion/EmitCCommon/GenericOpConversion.h b/include/emitc/Conversion/EmitCCommon/GenericOpConversion.h
index 44c90bdb..9ac6a734 100644
--- a/include/emitc/Conversion/EmitCCommon/GenericOpConversion.h
+++ b/include/emitc/Conversion/EmitCCommon/GenericOpConversion.h
@@ -18,7 +18,7 @@ namespace {
 using namespace mlir;
 using namespace mlir::emitc;
 
-/// Convert a common operation into an `emitc.call` operation.
+/// Convert a common operation into an `emitc.call_opaque` operation.
 template <typename SrcOp, typename Adaptor = typename SrcOp::Adaptor>
 class GenericOpConversion : public OpConversionPattern<SrcOp> {
   using OpConversionPattern<SrcOp>::OpConversionPattern;
@@ -57,9 +57,9 @@ class GenericOpConversion : public OpConversionPattern<SrcOp> {
       templateArgs = ArrayAttr::get(srcOp.getContext(), templateArguments);
     }
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(srcOp, srcOp.getType(), callee,
-                                               args, templateArgs,
-                                               adaptor.getOperands());
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(srcOp, srcOp.getType(),
+                                                     callee, args, templateArgs,
+                                                     adaptor.getOperands());
 
     return success();
   }
diff --git a/lib/Conversion/ArithToEmitC/ArithToEmitC.cpp b/lib/Conversion/ArithToEmitC/ArithToEmitC.cpp
index 43c1b466..e54ae1f7 100644
--- a/lib/Conversion/ArithToEmitC/ArithToEmitC.cpp
+++ b/lib/Conversion/ArithToEmitC/ArithToEmitC.cpp
@@ -25,7 +25,7 @@ using namespace mlir::emitc;
 
 namespace {
 
-// Convert `arith.index_cast` into an `emitc.call` operation.
+// Convert `arith.index_cast` into an `emitc.call_opaque` operation.
 class IndexCastOpConversion : public OpConversionPattern<arith::IndexCastOp> {
   using OpConversionPattern<arith::IndexCastOp>::OpConversionPattern;
 
@@ -43,7 +43,7 @@ class IndexCastOpConversion : public OpConversionPattern<arith::IndexCastOp> {
     Type resultType = indexCastOp.getResult().getType();
     ArrayAttr templateArgs = rewriter.getArrayAttr({TypeAttr::get(resultType)});
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(
         indexCastOp, indexCastOp.getType(), callee, args, templateArgs,
         adaptor.getOperands());
 
diff --git a/lib/Conversion/StablehloToEmitC/StablehloRegionOpsToEmitC.cpp b/lib/Conversion/StablehloToEmitC/StablehloRegionOpsToEmitC.cpp
index e8fb4a37..cb78e22f 100644
--- a/lib/Conversion/StablehloToEmitC/StablehloRegionOpsToEmitC.cpp
+++ b/lib/Conversion/StablehloToEmitC/StablehloRegionOpsToEmitC.cpp
@@ -181,9 +181,9 @@ struct ConvertStablehloRegionOpsToEmitCPass
 
     ArrayAttr templateArgs = ArrayAttr::get(ctx, templateArguments);
 
-    emitc::CallOp callOp = builder.create<emitc::CallOp>(
+    emitc::CallOpaqueOp callOpaqueOp = builder.create<emitc::CallOpaqueOp>(
         op.getLoc(), op.getResultTypes(), callee, args, templateArgs, operands);
-    op.replaceAllUsesWith(callOp);
+    op.replaceAllUsesWith(callOpaqueOp);
     op.erase();
     return success();
   }
@@ -217,9 +217,9 @@ struct ConvertStablehloRegionOpsToEmitCPass
     ArrayAttr templateArgs =
         ArrayAttr::get(ctx, {TypeAttr::get(op.getResult(0).getType())});
 
-    emitc::CallOp callOp = builder.create<emitc::CallOp>(
+    emitc::CallOpaqueOp callOpaqueOp = builder.create<emitc::CallOpaqueOp>(
         op.getLoc(), op.getType(0), callee, args, templateArgs, operands);
-    op.replaceAllUsesWith(callOp);
+    op.replaceAllUsesWith(callOpaqueOp);
     op.erase();
     return success();
   }
diff --git a/lib/Conversion/StablehloToEmitC/StablehloToEmitC.cpp b/lib/Conversion/StablehloToEmitC/StablehloToEmitC.cpp
index b7886e17..64b5f5d7 100644
--- a/lib/Conversion/StablehloToEmitC/StablehloToEmitC.cpp
+++ b/lib/Conversion/StablehloToEmitC/StablehloToEmitC.cpp
@@ -56,7 +56,8 @@ class ConstOpConversion : public OpRewritePattern<stablehlo::ConstantOp> {
   }
 };
 
-/// Convert `stablehlo.batch_norm_inference` into an `emitc.call` operation.
+/// Convert `stablehlo.batch_norm_inference` into an `emitc.call_opaque`
+/// operation.
 class BatchNormInferenceOpConversion
     : public OpConversionPattern<stablehlo::BatchNormInferenceOp> {
 
@@ -83,7 +84,7 @@ class BatchNormInferenceOpConversion
         {TypeAttr::get(batchNormInferenceOp.getResult().getType()),
          TypeAttr::get(adaptor.getScale().getType())});
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(
         batchNormInferenceOp, batchNormInferenceOp.getType(), callee, args,
         templateArgs, adaptor.getOperands());
 
@@ -91,7 +92,7 @@ class BatchNormInferenceOpConversion
   }
 };
 
-/// Convert `stablehlo.broadcast_in_dim` into an `emitc.call` operation.
+/// Convert `stablehlo.broadcast_in_dim` into an `emitc.call_opaque` operation.
 class BroadcastInDimOpConversion
     : public OpConversionPattern<stablehlo::BroadcastInDimOp> {
 
@@ -116,7 +117,7 @@ class BroadcastInDimOpConversion
     ArrayAttr templateArgs = rewriter.getArrayAttr(
         {TypeAttr::get(broadcastInDimOp.getResult().getType())});
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(
         broadcastInDimOp, broadcastInDimOp.getType(), callee, args,
         templateArgs, adaptor.getOperands());
 
@@ -124,7 +125,7 @@ class BroadcastInDimOpConversion
   }
 };
 
-/// Convert `stablehlo.concatenate` into an `emitc.call` operation.
+/// Convert `stablehlo.concatenate` into an `emitc.call_opaque` operation.
 class ConcatenateOpConversion
     : public OpConversionPattern<stablehlo::ConcatenateOp> {
 
@@ -144,7 +145,7 @@ class ConcatenateOpConversion
         {rewriter.getI64IntegerAttr(concatenateOp.getDimension()),
          TypeAttr::get(concatenateOp.getResult().getType())});
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(
         concatenateOp, concatenateOp.getType(), callee, args, templateArgs,
         adaptor.getOperands());
 
@@ -152,7 +153,7 @@ class ConcatenateOpConversion
   }
 };
 
-/// Convert `stablehlo.convolution` into an `emitc.call` operation.
+/// Convert `stablehlo.convolution` into an `emitc.call_opaque` operation.
 class ConvOpConversion : public OpConversionPattern<stablehlo::ConvolutionOp> {
 
 public:
@@ -206,15 +207,15 @@ class ConvOpConversion : public OpConversionPattern<stablehlo::ConvolutionOp> {
                                TypeAttr::get(adaptor.getLhs().getType()),
                                TypeAttr::get(adaptor.getRhs().getType())});
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(convOp, convOp.getType(), callee,
-                                               args, templateArgs,
-                                               adaptor.getOperands());
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(convOp, convOp.getType(),
+                                                     callee, args, templateArgs,
+                                                     adaptor.getOperands());
 
     return success();
   }
 };
 
-/// Convert `stablehlo.compare` into an `emitc.call` operation.
+/// Convert `stablehlo.compare` into an `emitc.call_opaque` operation.
 class CompareOpConversion : public OpConversionPattern<stablehlo::CompareOp> {
   using OpConversionPattern<stablehlo::CompareOp>::OpConversionPattern;
 
@@ -252,15 +253,15 @@ class CompareOpConversion : public OpConversionPattern<stablehlo::CompareOp> {
         {TypeAttr::get(elementType),
          emitc::OpaqueAttr::get(ctx, functionName.value())});
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(compareOp, compareOp.getType(),
-                                               callee, args, templateArgs,
-                                               adaptor.getOperands());
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(
+        compareOp, compareOp.getType(), callee, args, templateArgs,
+        adaptor.getOperands());
 
     return success();
   }
 };
 
-/// Convert `stablehlo.get_tuple_element` into an `emitc.call` operation.
+/// Convert `stablehlo.get_tuple_element` into an `emitc.call_opaque` operation.
 class GetTupleElementOpConversion
     : public OpConversionPattern<stablehlo::GetTupleElementOp> {
   using OpConversionPattern<stablehlo::GetTupleElementOp>::OpConversionPattern;
@@ -282,7 +283,7 @@ class GetTupleElementOpConversion
     ArrayAttr templateArgs = rewriter.getArrayAttr(
         {IntegerAttr::get(rewriter.getIntegerType(32), index)});
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(
         getTupleElementOp, getTupleElementOp.getType(), callee, args,
         templateArgs, adaptor.getOperands());
 
@@ -290,7 +291,7 @@ class GetTupleElementOpConversion
   }
 };
 
-/// Convert `stablehlo.slice` into an `emitc.call` operation.
+/// Convert `stablehlo.slice` into an `emitc.call_opaque` operation.
 class SliceOpConversion : public OpConversionPattern<stablehlo::SliceOp> {
   using OpConversionPattern<stablehlo::SliceOp>::OpConversionPattern;
 
@@ -316,15 +317,15 @@ class SliceOpConversion : public OpConversionPattern<stablehlo::SliceOp> {
     ArrayAttr templateArgs =
         rewriter.getArrayAttr({TypeAttr::get(sliceOp.getResult().getType())});
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(sliceOp, sliceOp.getType(),
-                                               callee, args, templateArgs,
-                                               adaptor.getOperands());
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(sliceOp, sliceOp.getType(),
+                                                     callee, args, templateArgs,
+                                                     adaptor.getOperands());
 
     return success();
   }
 };
 
-/// Convert `stablehlo.dynamic_slice` into an `emitc.call` operation.
+/// Convert `stablehlo.dynamic_slice` into an `emitc.call_opaque` operation.
 class DynamicSliceOpConversion
     : public OpConversionPattern<stablehlo::DynamicSliceOp> {
   using OpConversionPattern<stablehlo::DynamicSliceOp>::OpConversionPattern;
@@ -350,7 +351,7 @@ class DynamicSliceOpConversion
     ArrayAttr templateArgs = rewriter.getArrayAttr(
         {TypeAttr::get(dynamicSliceOp.getResult().getType())});
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(
         dynamicSliceOp, dynamicSliceOp.getType(), callee, args, templateArgs,
         adaptor.getOperands());
 
@@ -358,7 +359,8 @@ class DynamicSliceOpConversion
   }
 };
 
-/// Convert `stablehlo.dynamic_update_slice` into an `emitc.call` operation.
+/// Convert `stablehlo.dynamic_update_slice` into an `emitc.call_opaque`
+/// operation.
 class DynamicUpdateSliceOpConversion
     : public OpConversionPattern<stablehlo::DynamicUpdateSliceOp> {
   using OpConversionPattern<
@@ -381,7 +383,7 @@ class DynamicUpdateSliceOpConversion
     ArrayAttr templateArgs =
         rewriter.getArrayAttr({TypeAttr::get(adaptor.getUpdate().getType())});
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(
         dynamicUpdateSliceOp, dynamicUpdateSliceOp.getType(), callee, args,
         templateArgs, adaptor.getOperands());
 
@@ -389,7 +391,7 @@ class DynamicUpdateSliceOpConversion
   }
 };
 
-/// Convert `stablehlo.pad` into an `emitc.call` operation.
+/// Convert `stablehlo.pad` into an `emitc.call_opaque` operation.
 class PadOpConversion : public OpConversionPattern<stablehlo::PadOp> {
   using OpConversionPattern<stablehlo::PadOp>::OpConversionPattern;
 
@@ -415,15 +417,15 @@ class PadOpConversion : public OpConversionPattern<stablehlo::PadOp> {
     Type resultType = padOp.getResult().getType();
     ArrayAttr templateArgs = rewriter.getArrayAttr({TypeAttr::get(resultType)});
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(padOp, padOp.getType(), callee,
-                                               args, templateArgs,
-                                               adaptor.getOperands());
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(padOp, padOp.getType(),
+                                                     callee, args, templateArgs,
+                                                     adaptor.getOperands());
 
     return success();
   }
 };
 
-/// Convert `stablehlo.transpose` into an `emitc.call` operation.
+/// Convert `stablehlo.transpose` into an `emitc.call_opaque` operation.
 class TransposeOpConversion
     : public OpConversionPattern<stablehlo::TransposeOp> {
   using OpConversionPattern<stablehlo::TransposeOp>::OpConversionPattern;
@@ -447,7 +449,7 @@ class TransposeOpConversion
     Type resultType = transposeOp.getResult().getType();
     ArrayAttr templateArgs = rewriter.getArrayAttr({TypeAttr::get(resultType)});
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(
         transposeOp, transposeOp.getType(), callee, args, templateArgs,
         adaptor.getOperands());
 
@@ -455,7 +457,7 @@ class TransposeOpConversion
   }
 };
 
-/// Convert `stablehlo.rng` into an `emitc.call` operation.
+/// Convert `stablehlo.rng` into an `emitc.call_opaque` operation.
 class RngOpConversion : public OpConversionPattern<stablehlo::RngOp> {
 
 public:
@@ -478,9 +480,9 @@ class RngOpConversion : public OpConversionPattern<stablehlo::RngOp> {
     ArrayAttr templateArgs =
         rewriter.getArrayAttr({TypeAttr::get(rngOp.getType())});
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(rngOp, rngOp.getType(), callee,
-                                               args, templateArgs,
-                                               adaptor.getOperands());
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(rngOp, rngOp.getType(),
+                                                     callee, args, templateArgs,
+                                                     adaptor.getOperands());
 
     return success();
   }
diff --git a/lib/Conversion/TensorToEmitC/TensorToEmitC.cpp b/lib/Conversion/TensorToEmitC/TensorToEmitC.cpp
index e7b87be0..fec2b3a7 100644
--- a/lib/Conversion/TensorToEmitC/TensorToEmitC.cpp
+++ b/lib/Conversion/TensorToEmitC/TensorToEmitC.cpp
@@ -25,7 +25,7 @@ using namespace mlir::emitc;
 
 namespace {
 
-/// Convert `tensor.extract` into an `emitc.call` operation.
+/// Convert `tensor.extract` into an `emitc.call_opaque` operation.
 class ExtractOpConversion : public OpConversionPattern<tensor::ExtractOp> {
   using OpConversionPattern<tensor::ExtractOp>::OpConversionPattern;
 
@@ -47,7 +47,7 @@ class ExtractOpConversion : public OpConversionPattern<tensor::ExtractOp> {
     ArrayAttr args;
     ArrayAttr templateArgs;
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(
         indexCastOp, indexCastOp.getType(), callee, args, templateArgs,
         adaptor.getOperands());
 
@@ -55,7 +55,7 @@ class ExtractOpConversion : public OpConversionPattern<tensor::ExtractOp> {
   }
 };
 
-/// Convert `tensor.splat` into an `emitc.call` operation.
+/// Convert `tensor.splat` into an `emitc.call_opaque` operation.
 class SplatOpConversion : public OpConversionPattern<tensor::SplatOp> {
   using OpConversionPattern<tensor::SplatOp>::OpConversionPattern;
 
@@ -73,9 +73,9 @@ class SplatOpConversion : public OpConversionPattern<tensor::SplatOp> {
     Type resultType = splatOp.getResult().getType();
     ArrayAttr templateArgs = rewriter.getArrayAttr({TypeAttr::get(resultType)});
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(splatOp, splatOp.getType(),
-                                               callee, args, templateArgs,
-                                               adaptor.getOperands());
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(splatOp, splatOp.getType(),
+                                                     callee, args, templateArgs,
+                                                     adaptor.getOperands());
 
     return success();
   }
diff --git a/lib/Conversion/TosaToEmitC/TosaToEmitC.cpp b/lib/Conversion/TosaToEmitC/TosaToEmitC.cpp
index 92d31d01..600a74ba 100644
--- a/lib/Conversion/TosaToEmitC/TosaToEmitC.cpp
+++ b/lib/Conversion/TosaToEmitC/TosaToEmitC.cpp
@@ -42,7 +42,7 @@ SmallVector<Attribute, 2> indexSequence(int64_t n, MLIRContext *ctx) {
       }));
 }
 
-/// Convert `tosa.concat` into an `emitc.call` operation.
+/// Convert `tosa.concat` into an `emitc.call_opaque` operation.
 class ConcatOpConversion : public OpConversionPattern<tosa::ConcatOp> {
 
 public:
@@ -61,9 +61,9 @@ class ConcatOpConversion : public OpConversionPattern<tosa::ConcatOp> {
         rewriter.getArrayAttr({concatOp.getAxisAttr(),
                                TypeAttr::get(concatOp.getResult().getType())});
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(concatOp, concatOp.getType(),
-                                               callee, args, templateArgs,
-                                               adaptor.getOperands());
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(
+        concatOp, concatOp.getType(), callee, args, templateArgs,
+        adaptor.getOperands());
 
     return success();
   }
@@ -82,7 +82,7 @@ class ConstOpConversion : public OpRewritePattern<tosa::ConstOp> {
   }
 };
 
-/// Convert a common `tosa` convolution operation into an `emitc.call`
+/// Convert a common `tosa` convolution operation into an `emitc.call_opaque`
 /// operation.
 template <typename SrcOp, typename Adaptor = typename SrcOp::Adaptor>
 class GenericConvOpConversion : public OpConversionPattern<SrcOp> {
@@ -125,9 +125,9 @@ class GenericConvOpConversion : public OpConversionPattern<SrcOp> {
         rewriter.getArrayAttr({TypeAttr::get(convOp.getResult().getType())});
 
     // Create conv op.
-    auto emitcConvOp =
-        rewriter.create<emitc::CallOp>(convOp->getLoc(), convOp.getType(),
-                                       callee, args, templateArgs, operands);
+    auto emitcConvOp = rewriter.create<emitc::CallOpaqueOp>(
+        convOp->getLoc(), convOp.getType(), callee, args, templateArgs,
+        operands);
 
     auto output = emitcConvOp.getResult(0);
     auto tosaAddOp = rewriter.create<tosa::AddOp>(
@@ -141,7 +141,7 @@ class GenericConvOpConversion : public OpConversionPattern<SrcOp> {
   StringRef funcName;
 };
 
-/// Convert a common `tosa` pooling operation into an `emitc.call`
+/// Convert a common `tosa` pooling operation into an `emitc.call_opaque`
 /// operation.
 template <typename SrcOp, typename Adaptor = typename SrcOp::Adaptor>
 class GenericPoolOpConversion : public OpConversionPattern<SrcOp> {
@@ -171,9 +171,9 @@ class GenericPoolOpConversion : public OpConversionPattern<SrcOp> {
         rewriter.getArrayAttr({TypeAttr::get(poolOp.getResult().getType())});
 
     // Create pool op.
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(poolOp, poolOp.getType(), callee,
-                                               args, templateArgs,
-                                               adaptor.getOperands());
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(poolOp, poolOp.getType(),
+                                                     callee, args, templateArgs,
+                                                     adaptor.getOperands());
 
     return success();
   }
@@ -181,7 +181,7 @@ class GenericPoolOpConversion : public OpConversionPattern<SrcOp> {
   StringRef funcName;
 };
 
-/// Convert `tosa.fully_connected` into an `emitc.call` operation.
+/// Convert `tosa.fully_connected` into an `emitc.call_opaque` operation.
 class FullyConnectedOpConversion
     : public OpConversionPattern<tosa::FullyConnectedOp> {
   using OpConversionPattern<tosa::FullyConnectedOp>::OpConversionPattern;
@@ -208,14 +208,14 @@ class FullyConnectedOpConversion
     ArrayAttr templateArgs =
         ArrayAttr::get(fullyConnectedOp.getContext(), {TypeAttr::get(type)});
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(fullyConnectedOp, type, callee,
-                                               args, templateArgs,
-                                               adaptor.getOperands());
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(fullyConnectedOp, type,
+                                                     callee, args, templateArgs,
+                                                     adaptor.getOperands());
     return success();
   }
 };
 
-/// Convert `tosa.matmul` into an `emitc.call` operation.
+/// Convert `tosa.matmul` into an `emitc.call_opaque` operation.
 class MatMulOpConversion : public OpConversionPattern<tosa::MatMulOp> {
   using OpConversionPattern<tosa::MatMulOp>::OpConversionPattern;
 
@@ -238,14 +238,14 @@ class MatMulOpConversion : public OpConversionPattern<tosa::MatMulOp> {
     ArrayAttr args;
     ArrayAttr templateArgs;
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(matMulOp, matMulOp.getType(),
-                                               callee, args, templateArgs,
-                                               adaptor.getOperands());
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(
+        matMulOp, matMulOp.getType(), callee, args, templateArgs,
+        adaptor.getOperands());
     return success();
   }
 };
 
-/// Convert `tosa.clamp` into an `emitc.call` operation.
+/// Convert `tosa.clamp` into an `emitc.call_opaque` operation.
 class ClampOpConversion : public OpConversionPattern<tosa::ClampOp> {
   using OpConversionPattern<tosa::ClampOp>::OpConversionPattern;
 
@@ -289,15 +289,15 @@ class ClampOpConversion : public OpConversionPattern<tosa::ClampOp> {
     ArrayAttr args = rewriter.getArrayAttr(arguments);
     ArrayAttr templateArgs;
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(clampOp, clampOp.getType(),
-                                               callee, args, templateArgs,
-                                               adaptor.getOperands());
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(clampOp, clampOp.getType(),
+                                                     callee, args, templateArgs,
+                                                     adaptor.getOperands());
 
     return success();
   }
 };
 
-/// Convert `tosa.negate` into an `emitc.call` operation.
+/// Convert `tosa.negate` into an `emitc.call_opaque` operation.
 class NegateOpConversion : public OpConversionPattern<tosa::NegateOp> {
   using OpConversionPattern<tosa::NegateOp>::OpConversionPattern;
 
@@ -320,14 +320,14 @@ class NegateOpConversion : public OpConversionPattern<tosa::NegateOp> {
     ArrayAttr args;
     ArrayAttr templateArgs;
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(negateOp, negateOp.getType(),
-                                               callee, args, templateArgs,
-                                               adaptor.getOperands());
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(
+        negateOp, negateOp.getType(), callee, args, templateArgs,
+        adaptor.getOperands());
     return success();
   }
 };
 
-/// Convert `tosa.rescale` into an `emitc.call` operation.
+/// Convert `tosa.rescale` into an `emitc.call_opaque` operation.
 class RescaleOpConversion : public OpConversionPattern<tosa::RescaleOp> {
   using OpConversionPattern<tosa::RescaleOp>::OpConversionPattern;
 
@@ -360,15 +360,15 @@ class RescaleOpConversion : public OpConversionPattern<tosa::RescaleOp> {
         rewriter.getI32IntegerAttr(rescaleOp.getMultiplierAttr().size());
     ArrayAttr templateArgs = rewriter.getArrayAttr({resultType, arraySize});
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(rescaleOp, rescaleOp.getType(),
-                                               callee, args, templateArgs,
-                                               adaptor.getOperands());
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(
+        rescaleOp, rescaleOp.getType(), callee, args, templateArgs,
+        adaptor.getOperands());
 
     return success();
   }
 };
 
-/// Convert `tosa.rsqrt` into an `emitc.call` operation.
+/// Convert `tosa.rsqrt` into an `emitc.call_opaque` operation.
 class RsqrtOpConversion : public OpConversionPattern<tosa::RsqrtOp> {
   using OpConversionPattern<tosa::RsqrtOp>::OpConversionPattern;
 
@@ -387,7 +387,7 @@ class RsqrtOpConversion : public OpConversionPattern<tosa::RsqrtOp> {
     StringRef sqrtFuncName = "emitc::sqrt";
     StringAttr sqrtCallee = rewriter.getStringAttr(sqrtFuncName);
 
-    auto sqrtEmitCOp = rewriter.create<emitc::CallOp>(
+    auto sqrtEmitCOp = rewriter.create<emitc::CallOpaqueOp>(
         rsqrtOp.getLoc(), rsqrtOp.getType(), sqrtCallee, args, templateArgs,
         adaptor.getOperands());
 
@@ -395,7 +395,7 @@ class RsqrtOpConversion : public OpConversionPattern<tosa::RsqrtOp> {
     StringRef reciprocalFuncName = "emitc::tosa::reciprocal";
     StringAttr reciprocalCallee = rewriter.getStringAttr(reciprocalFuncName);
 
-    auto reciprocalOp = rewriter.create<emitc::CallOp>(
+    auto reciprocalOp = rewriter.create<emitc::CallOpaqueOp>(
         sqrtEmitCOp.getLoc(), rsqrtOp.getType(), reciprocalCallee, args,
         templateArgs, sqrtEmitCOp.getResults());
 
@@ -452,7 +452,7 @@ createBroadcastOpIfNeeded(SrcOp &srcOp, Adaptor adaptor,
       ArrayAttr templateBroadcastArgs =
           rewriter.getArrayAttr({TypeAttr::get(newBroadcastType)});
 
-      auto broadcastArg = rewriter.create<emitc::CallOp>(
+      auto broadcastArg = rewriter.create<emitc::CallOpaqueOp>(
           srcOp->getLoc(), newBroadcastType, broadcastCallee, broadcastArgs,
           templateBroadcastArgs, operand);
       // Replace the original operand with the result of the broadcast_in_dim
@@ -466,7 +466,7 @@ createBroadcastOpIfNeeded(SrcOp &srcOp, Adaptor adaptor,
   return broadcastedOperands;
 }
 
-/// Convert a common, broadcastable `tosa` operation into an `emitc.call`
+/// Convert a common, broadcastable `tosa` operation into an `emitc.call_opaque`
 /// operation.
 template <typename SrcOp, typename Adaptor = typename SrcOp::Adaptor>
 class CallOpBroadcastableConversion : public OpConversionPattern<SrcOp> {
@@ -509,7 +509,7 @@ class CallOpBroadcastableConversion : public OpConversionPattern<SrcOp> {
     SmallVector<Value, 2> broadcastedOperands =
         createBroadcastOpIfNeeded(srcOp, adaptor, rewriter);
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(
         srcOp, srcOp.getType(), callee, args, templateArgs,
         ValueRange({broadcastedOperands[0], broadcastedOperands[1]}));
 
@@ -523,7 +523,7 @@ class CallOpBroadcastableConversion : public OpConversionPattern<SrcOp> {
   bool explicitOperandTypes;
 };
 
-/// Convert `tosa.mul` into an `emitc.call` operation.
+/// Convert `tosa.mul` into an `emitc.call_opaque` operation.
 class MulOpConversion : public OpConversionPattern<tosa::MulOp> {
   using OpConversionPattern<tosa::MulOp>::OpConversionPattern;
 
@@ -555,7 +555,7 @@ class MulOpConversion : public OpConversionPattern<tosa::MulOp> {
     SmallVector<Value, 2> broadcastedOperands =
         createBroadcastOpIfNeeded(mulOp, adaptor, rewriter);
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(
         mulOp, mulOp.getType(), callee, args, templateArgs,
         ValueRange({broadcastedOperands[0], broadcastedOperands[1]}));
 
@@ -563,7 +563,7 @@ class MulOpConversion : public OpConversionPattern<tosa::MulOp> {
   }
 };
 
-/// Convert `tosa.arithmetic_right_shift` into an `emitc.call` operation.
+/// Convert `tosa.arithmetic_right_shift` into an `emitc.call_opaque` operation.
 class ArithmeticRightShiftOpConversion
     : public OpConversionPattern<tosa::ArithmeticRightShiftOp> {
   using OpConversionPattern<tosa::ArithmeticRightShiftOp>::OpConversionPattern;
@@ -595,7 +595,7 @@ class ArithmeticRightShiftOpConversion
     SmallVector<Value, 2> broadcastedOperands =
         createBroadcastOpIfNeeded(arithmeticRightShiftOp, adaptor, rewriter);
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(
         arithmeticRightShiftOp, arithmeticRightShiftOp.getType(), callee, args,
         templateArgs,
         ValueRange({broadcastedOperands[0], broadcastedOperands[1]}));
@@ -604,7 +604,7 @@ class ArithmeticRightShiftOpConversion
   }
 };
 
-/// Convert `tosa.select` into an `emitc.call` operation.
+/// Convert `tosa.select` into an `emitc.call_opaque` operation.
 class SelectOpConversion : public OpConversionPattern<tosa::SelectOp> {
   using OpConversionPattern<tosa::SelectOp>::OpConversionPattern;
 
@@ -625,15 +625,15 @@ class SelectOpConversion : public OpConversionPattern<tosa::SelectOp> {
     SmallVector<Value, 3> broadcastedOperands =
         createBroadcastOpIfNeeded(selectOp, adaptor, rewriter);
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(selectOp, selectOp.getType(),
-                                               callee, args, templateArgs,
-                                               broadcastedOperands);
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(
+        selectOp, selectOp.getType(), callee, args, templateArgs,
+        broadcastedOperands);
 
     return success();
   }
 };
 
-/// Convert `tosa.reduce_*` into an `emitc.call` operation.
+/// Convert `tosa.reduce_*` into an `emitc.call_opaque` operation.
 template <typename SrcOp, typename Adaptor = typename SrcOp::Adaptor>
 class ReduceOpConversion : public OpConversionPattern<SrcOp> {
   using OpConversionPattern<SrcOp>::OpConversionPattern;
@@ -680,7 +680,7 @@ class ReduceOpConversion : public OpConversionPattern<SrcOp> {
           rewriter.getArrayAttr({TypeAttr::get(newOutputType),
                                  TypeAttr::get(reduceOp.getInput().getType())});
 
-      auto emitcReduceOp = rewriter.create<emitc::CallOp>(
+      auto emitcReduceOp = rewriter.create<emitc::CallOpaqueOp>(
           reduceOp.getLoc(), newOutputType, callee, args, templateArgs,
           adaptor.getOperands());
 
@@ -699,9 +699,9 @@ class ReduceOpConversion : public OpConversionPattern<SrcOp> {
       ArrayAttr templateArgs =
           rewriter.getArrayAttr({TypeAttr::get(reduceOp.getType()),
                                  TypeAttr::get(reduceOp.getInput().getType())});
-      rewriter.replaceOpWithNewOp<emitc::CallOp>(reduceOp, reduceOp.getType(),
-                                                 callee, args, templateArgs,
-                                                 adaptor.getOperands());
+      rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(
+          reduceOp, reduceOp.getType(), callee, args, templateArgs,
+          adaptor.getOperands());
     }
 
     return success();
@@ -711,7 +711,7 @@ class ReduceOpConversion : public OpConversionPattern<SrcOp> {
   bool keepDims;
 };
 
-/// Convert `tosa.pad` into an `emitc.call` operation.
+/// Convert `tosa.pad` into an `emitc.call_opaque` operation.
 class PadOpConversion : public OpConversionPattern<tosa::PadOp> {
   using OpConversionPattern<tosa::PadOp>::OpConversionPattern;
 
@@ -736,15 +736,15 @@ class PadOpConversion : public OpConversionPattern<tosa::PadOp> {
     Type resultType = padOp.getOutput().getType();
     ArrayAttr templateArgs = rewriter.getArrayAttr({TypeAttr::get(resultType)});
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(padOp, padOp.getType(), callee,
-                                               args, templateArgs,
-                                               adaptor.getOperands());
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(padOp, padOp.getType(),
+                                                     callee, args, templateArgs,
+                                                     adaptor.getOperands());
 
     return success();
   }
 };
 
-/// Convert `tosa.slice` into an `emitc.call` operation.
+/// Convert `tosa.slice` into an `emitc.call_opaque` operation.
 class SliceOpConversion : public OpConversionPattern<tosa::SliceOp> {
   using OpConversionPattern<tosa::SliceOp>::OpConversionPattern;
 
@@ -769,15 +769,15 @@ class SliceOpConversion : public OpConversionPattern<tosa::SliceOp> {
     Type resultType = sliceOp.getOutput().getType();
     ArrayAttr templateArgs = rewriter.getArrayAttr({TypeAttr::get(resultType)});
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(sliceOp, sliceOp.getType(),
-                                               callee, args, templateArgs,
-                                               adaptor.getOperands());
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(sliceOp, sliceOp.getType(),
+                                                     callee, args, templateArgs,
+                                                     adaptor.getOperands());
 
     return success();
   }
 };
 
-/// Convert `tosa.tile` into an `emitc.call` operation.
+/// Convert `tosa.tile` into an `emitc.call_opaque` operation.
 class TileOpConversion : public OpConversionPattern<tosa::TileOp> {
   using OpConversionPattern<tosa::TileOp>::OpConversionPattern;
 
@@ -808,9 +808,9 @@ class TileOpConversion : public OpConversionPattern<tosa::TileOp> {
     Type resultType = tileOp.getOutput().getType();
     ArrayAttr templateArgs = rewriter.getArrayAttr({TypeAttr::get(resultType)});
 
-    rewriter.replaceOpWithNewOp<emitc::CallOp>(tileOp, tileOp.getType(), callee,
-                                               args, templateArgs,
-                                               adaptor.getOperands());
+    rewriter.replaceOpWithNewOp<emitc::CallOpaqueOp>(tileOp, tileOp.getType(),
+                                                     callee, args, templateArgs,
+                                                     adaptor.getOperands());
     return success();
   }
 };
diff --git a/test/Conversion/arith-to-emitc.mlir b/test/Conversion/arith-to-emitc.mlir
index 8d12bee4..c9fff13f 100644
--- a/test/Conversion/arith-to-emitc.mlir
+++ b/test/Conversion/arith-to-emitc.mlir
@@ -12,9 +12,9 @@ func.func @arith_index_cast(%arg0: tensor<index>, %arg1: tensor<2xi32>, %arg2: t
   return %1 : tensor<2xindex>
 }
 // CHECK-LABEL: func @arith_index_cast
-//  CHECK-NEXT: emitc.call "emitc::arith::index_cast"(%arg0) {template_args = [tensor<i32>]} : (tensor<index>) -> tensor<i32>
-//  CHECK-NEXT: emitc.call "emitc::arith::index_cast"(%arg1) {template_args = [tensor<2xindex>]} : (tensor<2xi32>) -> tensor<2xindex>
-//  CHECK-NEXT: emitc.call "emitc::arith::index_cast"(%arg2) {template_args = [tensor<2x2xindex>]} : (tensor<2x2xi32>) -> tensor<2x2xindex>
+//  CHECK-NEXT: emitc.call_opaque "emitc::arith::index_cast"(%arg0) {template_args = [tensor<i32>]} : (tensor<index>) -> tensor<i32>
+//  CHECK-NEXT: emitc.call_opaque "emitc::arith::index_cast"(%arg1) {template_args = [tensor<2xindex>]} : (tensor<2xi32>) -> tensor<2xindex>
+//  CHECK-NEXT: emitc.call_opaque "emitc::arith::index_cast"(%arg2) {template_args = [tensor<2x2xindex>]} : (tensor<2x2xi32>) -> tensor<2x2xindex>
 
 // CPP-LABEL: Tensor<size_t, 2> arith_index_cast(Tensor<size_t> v1, Tensor<int32_t, 2> v2, Tensor<int32_t, 2, 2> v3)
 //  CPP-NEXT: emitc::arith::index_cast<Tensor<int32_t>>(v1)
diff --git a/test/Conversion/stablehlo-to-emitc.mlir b/test/Conversion/stablehlo-to-emitc.mlir
index f22717d8..a673717e 100644
--- a/test/Conversion/stablehlo-to-emitc.mlir
+++ b/test/Conversion/stablehlo-to-emitc.mlir
@@ -17,91 +17,91 @@ func.func @stablehlo_constant(%arg0: tensor<2xi32>) -> tensor<2xi32> {
 
 
 func.func @float_abs(%arg0: tensor<2xf32>) -> tensor<2xf32> {
-  // CHECK: emitc.call "emitc::stablehlo::abs"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::abs"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
   %0 = "stablehlo.abs"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
   return %0 : tensor<2xf32>
 }
 
 func.func @stablehlo_ceil(%arg0: tensor<2xf32>) -> tensor<2xf32> {
-  // CHECK: emitc.call "emitc::stablehlo::ceil"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::ceil"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
   %0 = "stablehlo.ceil"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
   return %0 : tensor<2xf32>
 }
 
 func.func @stablehlo_convert(%arg0: tensor<ui32>) -> tensor<ui64> {
-  // CHECK: emitc.call "emitc::stablehlo::convert"(%arg0) {template_args = [tensor<ui64>]} : (tensor<ui32>) -> tensor<ui64>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::convert"(%arg0) {template_args = [tensor<ui64>]} : (tensor<ui32>) -> tensor<ui64>
   %0 = "stablehlo.convert"(%arg0) : (tensor<ui32>) -> tensor<ui64>
   return %0 : tensor<ui64>
 }
 
 func.func @stablehlo_cos(%arg0: tensor<2xf32>) -> tensor<2xf32> {
-  // CHECK: emitc.call "emitc::stablehlo::cos"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::cos"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
   %0 = "stablehlo.cosine"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
   return %0 : tensor<2xf32>
 }
 
 func.func @stablehlo_exponential(%arg0: tensor<2xf32>) -> tensor<2xf32> {
-  // CHECK: emitc.call "emitc::stablehlo::exponential"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::exponential"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
   %0 = "stablehlo.exponential"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
   return %0 : tensor<2xf32>
 }
 
 func.func @stablehlo_exponential_minus_one(%arg0: tensor<2xf32>) -> tensor<2xf32> {
-  // CHECK: emitc.call "emitc::stablehlo::exponential_minus_one"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::exponential_minus_one"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
   %0 = "stablehlo.exponential_minus_one"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
   return %0 : tensor<2xf32>
 }
 
 func.func @stablehlo_floor(%arg0: tensor<2xf32>) -> tensor<2xf32> {
-  // CHECK: emitc.call "emitc::stablehlo::floor"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::floor"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
   %0 = "stablehlo.floor"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
   return %0 : tensor<2xf32>
 }
 
 func.func @stablehlo_is_finite(%arg0: tensor<4xf32>) -> tensor<4xi1> {
-  // CHECK: emitc.call "emitc::stablehlo::is_finite"(%arg0) : (tensor<4xf32>) -> tensor<4xi1>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::is_finite"(%arg0) : (tensor<4xf32>) -> tensor<4xi1>
   %0 = "stablehlo.is_finite"(%arg0) : (tensor<4xf32>) -> tensor<4xi1>
   return %0 : tensor<4xi1>
 }
 
 func.func @stablehlo_log(%arg0: tensor<2xf32>) -> tensor<2xf32> {
-  // CHECK: emitc.call "emitc::stablehlo::log"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::log"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
   %0 = "stablehlo.log"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
   return %0 : tensor<2xf32>
 }
 
 func.func @stablehlo_log_plus_one(%arg0: tensor<2xf32>) -> tensor<2xf32> {
-  // CHECK: emitc.call "emitc::stablehlo::log_plus_one"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::log_plus_one"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
   %0 = "stablehlo.log_plus_one"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
   return %0 : tensor<2xf32>
 }
 
 func.func @stablehlo_negate(%arg0: tensor<2xf32>) -> tensor<2xf32> {
-  // CHECK: emitc.call "emitc::stablehlo::negate"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::negate"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
   %0 = "stablehlo.negate"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
   return %0 : tensor<2xf32>
 }
 
 func.func @stablehlo_round(%arg0: tensor<2xf32>) -> tensor<2xf32> {
-  // CHECK: emitc.call "emitc::stablehlo::round"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::round"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
   %0 = "stablehlo.round_nearest_afz"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
   return %0 : tensor<2xf32>
 }
 
 func.func @stablehlo_sine(%arg0: tensor<2xf32>) -> tensor<2xf32> {
-  // CHECK: emitc.call "emitc::stablehlo::sin"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::sin"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
   %0 = "stablehlo.sine"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
   return %0 : tensor<2xf32>
 }
 
 func.func @stablehlo_sqrt(%arg0: tensor<2xf32>) -> tensor<2xf32> {
-  // CHECK: emitc.call "emitc::stablehlo::sqrt"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::sqrt"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
   %0 = "stablehlo.sqrt"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
   return %0 : tensor<2xf32>
 }
 
 func.func @stablehlo_tanh(%arg0: tensor<2xf32>) -> tensor<2xf32> {
-  // CHECK: emitc.call "emitc::stablehlo::tanh"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::tanh"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
   %0 = "stablehlo.tanh"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
   return %0 : tensor<2xf32>
 }
@@ -110,67 +110,67 @@ func.func @stablehlo_tanh(%arg0: tensor<2xf32>) -> tensor<2xf32> {
 // Binary elementwise ops
 
 func.func @stablehlo_add_i64(%arg0: tensor<i64>) -> tensor<i64> {
-  // CHECK: emitc.call "emitc::stablehlo::add"(%arg0, %arg0) : (tensor<i64>, tensor<i64>) -> tensor<i64>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::add"(%arg0, %arg0) : (tensor<i64>, tensor<i64>) -> tensor<i64>
   %0 = stablehlo.add %arg0, %arg0 : tensor<i64>
   return %0 : tensor<i64>
 }
 
 func.func @stablehlo_add_f64(%arg0: tensor<f64>) -> tensor<f64> {
-  // CHECK: emitc.call "emitc::stablehlo::add"(%arg0, %arg0) : (tensor<f64>, tensor<f64>) -> tensor<f64>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::add"(%arg0, %arg0) : (tensor<f64>, tensor<f64>) -> tensor<f64>
   %0 = stablehlo.add %arg0, %arg0 : tensor<f64>
   return %0 : tensor<f64>
 }
 
 func.func @stablehlo_atan2(%arg0: tensor<2xf32>, %arg1: tensor<2xf32>) -> tensor<2xf32> {
-  // CHECK: emitc.call "emitc::stablehlo::atan2"
+  // CHECK: emitc.call_opaque "emitc::stablehlo::atan2"
   %0 = "stablehlo.atan2"(%arg0, %arg1) : (tensor<2xf32>, tensor<2xf32>) -> tensor<2xf32>
   return %0 : tensor<2xf32>
 }
 
 func.func @stablehlo_divide(%arg0: tensor<f32>, %arg1: tensor<f32>) -> tensor<f32> {
-  // CHECK: emitc.call "emitc::stablehlo::div"(%arg0, %arg1) : (tensor<f32>, tensor<f32>) -> tensor<f32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::div"(%arg0, %arg1) : (tensor<f32>, tensor<f32>) -> tensor<f32>
   %0 = "stablehlo.divide"(%arg0, %arg1) : (tensor<f32>, tensor<f32>) -> tensor<f32>
   return %0 : tensor<f32>
 }
 
 func.func @stablehlo_max(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32> {
-  // CHECK: emitc.call "emitc::stablehlo::max"(%arg0, %arg0) : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::max"(%arg0, %arg0) : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32>
   %0 = "stablehlo.maximum"(%arg0, %arg0) : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32>
   return %0 : tensor<4xf32>
 }
 
 func.func @stablehlo_min(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32> {
-  // CHECK: emitc.call "emitc::stablehlo::min"(%arg0, %arg0) : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::min"(%arg0, %arg0) : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32>
   %0 = "stablehlo.minimum"(%arg0, %arg0) : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32>
   return %0 : tensor<4xf32>
 }
 
 func.func @stablehlo_multiply(%arg0: tensor<f32>, %arg1: tensor<f32>) -> tensor<f32> {
-  // CHECK: emitc.call "emitc::stablehlo::mul"(%arg0, %arg1) : (tensor<f32>, tensor<f32>) -> tensor<f32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::mul"(%arg0, %arg1) : (tensor<f32>, tensor<f32>) -> tensor<f32>
   %0 = "stablehlo.multiply"(%arg0, %arg1) : (tensor<f32>, tensor<f32>) -> tensor<f32>
   return %0 : tensor<f32>
 }
 
 func.func @stablehlo_power(%arg0: tensor<f32>, %arg1: tensor<f32>) -> tensor<f32> {
-  // CHECK: emitc.call "emitc::stablehlo::pow"(%arg0, %arg1) : (tensor<f32>, tensor<f32>) -> tensor<f32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::pow"(%arg0, %arg1) : (tensor<f32>, tensor<f32>) -> tensor<f32>
   %0 = "stablehlo.power"(%arg0, %arg1) : (tensor<f32>, tensor<f32>) -> tensor<f32>
   return %0 : tensor<f32>
 }
 
 func.func @stablehlo_shift_left(%arg0: tensor<i32>, %arg1: tensor<i32>) -> tensor<i32> {
-  // CHECK: emitc.call "emitc::stablehlo::shift_left"(%arg0, %arg1) : (tensor<i32>, tensor<i32>) -> tensor<i32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::shift_left"(%arg0, %arg1) : (tensor<i32>, tensor<i32>) -> tensor<i32>
   %0 = "stablehlo.shift_left"(%arg0, %arg1) : (tensor<i32>, tensor<i32>) -> tensor<i32>
   return %0 : tensor<i32>
 }
 
 func.func @stablehlo_shift_right_logical(%arg0: tensor<i32>, %arg1: tensor<i32>) -> tensor<i32> {
-  // CHECK: emitc.call "emitc::stablehlo::shift_right_logical"(%arg0, %arg1) : (tensor<i32>, tensor<i32>) -> tensor<i32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::shift_right_logical"(%arg0, %arg1) : (tensor<i32>, tensor<i32>) -> tensor<i32>
   %0 = "stablehlo.shift_right_logical"(%arg0, %arg1) : (tensor<i32>, tensor<i32>) -> tensor<i32>
   return %0 : tensor<i32>
 }
 
 func.func @stablehlo_sub(%arg0: tensor<f32>) -> tensor<f32> {
-  // CHECK: emitc.call "emitc::stablehlo::sub"(%arg0, %arg0) : (tensor<f32>, tensor<f32>) -> tensor<f32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::sub"(%arg0, %arg0) : (tensor<f32>, tensor<f32>) -> tensor<f32>
   %0 = "stablehlo.subtract"(%arg0, %arg0) : (tensor<f32>, tensor<f32>) -> tensor<f32>
   return %0 : tensor<f32>
 }
@@ -179,13 +179,13 @@ func.func @stablehlo_sub(%arg0: tensor<f32>) -> tensor<f32> {
 // Binary logical elementwise ops
 
 func.func @stablehlo_or(%arg0: tensor<ui64>, %arg1: tensor<ui64>) -> tensor<ui64> {
-  // CHECK: emitc.call "emitc::stablehlo::logical_or"(%arg0, %arg1) : (tensor<ui64>, tensor<ui64>) -> tensor<ui64>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::logical_or"(%arg0, %arg1) : (tensor<ui64>, tensor<ui64>) -> tensor<ui64>
   %0 = "stablehlo.or"(%arg0, %arg1) : (tensor<ui64>, tensor<ui64>) -> tensor<ui64>
   return %0 : tensor<ui64>
 }
 
 func.func @stablehlo_xor(%arg0: tensor<ui64>, %arg1: tensor<ui64>) -> tensor<ui64> {
-  // CHECK: emitc.call "emitc::stablehlo::logical_xor"(%arg0, %arg1) : (tensor<ui64>, tensor<ui64>) -> tensor<ui64>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::logical_xor"(%arg0, %arg1) : (tensor<ui64>, tensor<ui64>) -> tensor<ui64>
   %0 = "stablehlo.xor"(%arg0, %arg1) : (tensor<ui64>, tensor<ui64>) -> tensor<ui64>
   return %0 : tensor<ui64>
 }
@@ -194,46 +194,46 @@ func.func @stablehlo_xor(%arg0: tensor<ui64>, %arg1: tensor<ui64>) -> tensor<ui6
 // Tuple ops
 
 func.func @stablehlo_tuple(%arg0: tensor<i32>, %arg1: tensor<ui64>) -> (tuple<tensor<i32>, tensor<ui64>, tensor<i32>, tensor<ui64>>) {
-  // CHECK: emitc.call "std::make_tuple"() : () -> tuple<>
+  // CHECK: emitc.call_opaque "std::make_tuple"() : () -> tuple<>
   %0 = "stablehlo.tuple"() : () -> tuple<>
-  // CHECK: emitc.call "std::make_tuple"(%arg0) : (tensor<i32>) -> tuple<tensor<i32>>
+  // CHECK: emitc.call_opaque "std::make_tuple"(%arg0) : (tensor<i32>) -> tuple<tensor<i32>>
   %1 = "stablehlo.tuple"(%arg0) : (tensor<i32>) -> tuple<tensor<i32>>
-  // CHECK: emitc.call "std::make_tuple"(%arg0, %arg1) : (tensor<i32>, tensor<ui64>) -> tuple<tensor<i32>, tensor<ui64>>
+  // CHECK: emitc.call_opaque "std::make_tuple"(%arg0, %arg1) : (tensor<i32>, tensor<ui64>) -> tuple<tensor<i32>, tensor<ui64>>
   %2 = "stablehlo.tuple"(%arg0, %arg1) : (tensor<i32>, tensor<ui64>) -> tuple<tensor<i32>, tensor<ui64>>
-  // CHECK: emitc.call "std::make_tuple"(%arg0, %arg1, %arg0, %arg1) : (tensor<i32>, tensor<ui64>, tensor<i32>, tensor<ui64>) -> tuple<tensor<i32>, tensor<ui64>, tensor<i32>, tensor<ui64>>
+  // CHECK: emitc.call_opaque "std::make_tuple"(%arg0, %arg1, %arg0, %arg1) : (tensor<i32>, tensor<ui64>, tensor<i32>, tensor<ui64>) -> tuple<tensor<i32>, tensor<ui64>, tensor<i32>, tensor<ui64>>
   %3 = "stablehlo.tuple"(%arg0, %arg1, %arg0, %arg1) : (tensor<i32>, tensor<ui64>, tensor<i32>, tensor<ui64>) -> tuple<tensor<i32>, tensor<ui64>, tensor<i32>, tensor<ui64>>
   return %3 : tuple<tensor<i32>, tensor<ui64>, tensor<i32>, tensor<ui64>>
 }
 
 func.func @stablehlo_tuple_nested(%arg0: tensor<i32>, %arg1: tensor<ui64>) -> tuple<tensor<i32>, tuple<tensor<i32>, tensor<ui64>>> {
-  // CHECK: emitc.call "std::make_tuple"(%arg0, %arg1) : (tensor<i32>, tensor<ui64>) -> tuple<tensor<i32>, tensor<ui64>>
+  // CHECK: emitc.call_opaque "std::make_tuple"(%arg0, %arg1) : (tensor<i32>, tensor<ui64>) -> tuple<tensor<i32>, tensor<ui64>>
   %0 = "stablehlo.tuple"(%arg0, %arg1) : (tensor<i32>, tensor<ui64>) -> tuple<tensor<i32>, tensor<ui64>>
-  // CHECK: emitc.call "std::make_tuple"(%arg0, %0) : (tensor<i32>, tuple<tensor<i32>, tensor<ui64>>) -> tuple<tensor<i32>, tuple<tensor<i32>, tensor<ui64>>>
+  // CHECK: emitc.call_opaque "std::make_tuple"(%arg0, %0) : (tensor<i32>, tuple<tensor<i32>, tensor<ui64>>) -> tuple<tensor<i32>, tuple<tensor<i32>, tensor<ui64>>>
   %1 = "stablehlo.tuple"(%arg0, %0) : (tensor<i32>, tuple<tensor<i32>, tensor<ui64>>) -> tuple<tensor<i32>, tuple<tensor<i32>, tensor<ui64>>>
   return %1 : tuple<tensor<i32>, tuple<tensor<i32>, tensor<ui64>>>
 }
 
 func.func @stablehlo_tuple_unpack(%arg0: tensor<i32>, %arg1: tensor<ui64>) -> (tuple<tensor<i32>, tensor<ui64>>, tensor<i32>) {
   %0 = call @stablehlo_tuple_nested(%arg0, %arg1) : (tensor<i32>, tensor<ui64>) -> tuple<tensor<i32>, tuple<tensor<i32>, tensor<ui64>>>
-  // CHECK: emitc.call "std::get"(%0) {template_args = [1 : i32]} : (tuple<tensor<i32>, tuple<tensor<i32>, tensor<ui64>>>) -> tuple<tensor<i32>, tensor<ui64>>
+  // CHECK: emitc.call_opaque "std::get"(%0) {template_args = [1 : i32]} : (tuple<tensor<i32>, tuple<tensor<i32>, tensor<ui64>>>) -> tuple<tensor<i32>, tensor<ui64>>
   %1 = "stablehlo.get_tuple_element"(%0) {index = 1 : i32} : (tuple<tensor<i32>, tuple<tensor<i32>, tensor<ui64>>>) -> tuple<tensor<i32>, tensor<ui64>>
-  // CHECK: emitc.call "std::get"(%1) {template_args = [0 : i32]} : (tuple<tensor<i32>, tensor<ui64>>) -> tensor<i32>
+  // CHECK: emitc.call_opaque "std::get"(%1) {template_args = [0 : i32]} : (tuple<tensor<i32>, tensor<ui64>>) -> tensor<i32>
   %2 = "stablehlo.get_tuple_element"(%1) {index = 0 : i32} : (tuple<tensor<i32>, tensor<ui64>>) -> tensor<i32>
   return %1, %2 : tuple<tensor<i32>, tensor<ui64>>, tensor<i32>
 }
 
 func.func @stablehlo_compare(%arg0: tensor<4xi32>, %arg1: tensor<4xi32>) -> tensor<4xi1> {
-  // CHECK: emitc.call "emitc::stablehlo::compare"(%arg0, %arg1) {template_args = [tensor<4xi32>, #emitc.opaque<"std::less">]} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::compare"(%arg0, %arg1) {template_args = [tensor<4xi32>, #emitc.opaque<"std::less">]} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1>
   %0 = "stablehlo.compare"(%arg0, %arg1) {comparison_direction = #stablehlo<comparison_direction LT>} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1>
-  // CHECK: emitc.call "emitc::stablehlo::compare"(%arg0, %arg1) {template_args = [tensor<4xi32>, #emitc.opaque<"std::less_equal">]} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::compare"(%arg0, %arg1) {template_args = [tensor<4xi32>, #emitc.opaque<"std::less_equal">]} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1>
   %1 = "stablehlo.compare"(%arg0, %arg1) {comparison_direction = #stablehlo<comparison_direction LE>} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1>
-  // CHECK: emitc.call "emitc::stablehlo::compare"(%arg0, %arg1) {template_args = [tensor<4xi32>, #emitc.opaque<"std::greater">]} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::compare"(%arg0, %arg1) {template_args = [tensor<4xi32>, #emitc.opaque<"std::greater">]} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1>
   %2 = "stablehlo.compare"(%arg0, %arg1) {comparison_direction = #stablehlo<comparison_direction GT>} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1>
-  // CHECK: emitc.call "emitc::stablehlo::compare"(%arg0, %arg1) {template_args = [tensor<4xi32>, #emitc.opaque<"std::greater_equal">]} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::compare"(%arg0, %arg1) {template_args = [tensor<4xi32>, #emitc.opaque<"std::greater_equal">]} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1>
   %3 = "stablehlo.compare"(%arg0, %arg1) {comparison_direction = #stablehlo<comparison_direction GE>} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1>
-  // CHECK: emitc.call "emitc::stablehlo::compare"(%arg0, %arg1) {template_args = [tensor<4xi32>, #emitc.opaque<"std::equal_to">]} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::compare"(%arg0, %arg1) {template_args = [tensor<4xi32>, #emitc.opaque<"std::equal_to">]} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1>
   %4 = "stablehlo.compare"(%arg0, %arg1) {comparison_direction = #stablehlo<comparison_direction EQ>} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1>
-  // CHECK: emitc.call "emitc::stablehlo::compare"(%arg0, %arg1) {template_args = [tensor<4xi32>, #emitc.opaque<"std::not_equal_to">]} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::compare"(%arg0, %arg1) {template_args = [tensor<4xi32>, #emitc.opaque<"std::not_equal_to">]} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1>
   %5 = "stablehlo.compare"(%arg0, %arg1) {comparison_direction = #stablehlo<comparison_direction NE>} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1>
 
   return %0 : tensor<4xi1>
@@ -243,9 +243,9 @@ func.func @stablehlo_compare(%arg0: tensor<4xi32>, %arg1: tensor<4xi32>) -> tens
 // Slice ops
 
 func.func @stablehlo_slice(%arg0: tensor<12xi32>, %arg1: tensor<8x7xi32>) -> tensor<4x3xi32> {
-  // CHECK: emitc.call "emitc::stablehlo::slice"(%arg0) {args = [0 : index, dense<0> : tensor<1xi64>, dense<1> : tensor<1xi64>, dense<1> : tensor<1xi64>], template_args = [tensor<1xi32>]} : (tensor<12xi32>) -> tensor<1xi32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::slice"(%arg0) {args = [0 : index, dense<0> : tensor<1xi64>, dense<1> : tensor<1xi64>, dense<1> : tensor<1xi64>], template_args = [tensor<1xi32>]} : (tensor<12xi32>) -> tensor<1xi32>
   %0 = "stablehlo.slice"(%arg0) {limit_indices = dense<1> : tensor<1xi64>, start_indices = dense<0> : tensor<1xi64>, strides = dense<1> : tensor<1xi64>} : (tensor<12xi32>) -> tensor<1xi32>
-  // CHECK: emitc.call "emitc::stablehlo::slice"(%arg1) {args = [0 : index, dense<0> : tensor<2xi64>, dense<[4, 3]> : tensor<2xi64>, dense<1> : tensor<2xi64>], template_args = [tensor<4x3xi32>]} : (tensor<8x7xi32>) -> tensor<4x3xi32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::slice"(%arg1) {args = [0 : index, dense<0> : tensor<2xi64>, dense<[4, 3]> : tensor<2xi64>, dense<1> : tensor<2xi64>], template_args = [tensor<4x3xi32>]} : (tensor<8x7xi32>) -> tensor<4x3xi32>
   %1 = "stablehlo.slice"(%arg1) {limit_indices = dense<[4, 3]> : tensor<2xi64>, start_indices = dense<0> : tensor<2xi64>, strides = dense<1> : tensor<2xi64>} : (tensor<8x7xi32>) -> tensor<4x3xi32>
   return %1 : tensor<4x3xi32>
 }
@@ -253,9 +253,9 @@ func.func @stablehlo_slice(%arg0: tensor<12xi32>, %arg1: tensor<8x7xi32>) -> ten
 func.func @stablehlo_dynamic_slice(%arg0: tensor<12xi32>, %arg1: tensor<8x7xi32>) -> () {
   %cst = "arith.constant"() {value = dense<1> : tensor<i64>} : () -> tensor<i64>
   %cst_0 = "arith.constant"() {value = dense<3> : tensor<i64>} : () -> tensor<i64>
-  // CHECK: emitc.call "emitc::stablehlo::dynamic_slice"(%arg0, %cst) {args = [0 : index, 1 : index, dense<4> : tensor<1xi64>], template_args = [tensor<4xi32>]} : (tensor<12xi32>, tensor<i64>) -> tensor<4xi32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::dynamic_slice"(%arg0, %cst) {args = [0 : index, 1 : index, dense<4> : tensor<1xi64>], template_args = [tensor<4xi32>]} : (tensor<12xi32>, tensor<i64>) -> tensor<4xi32>
   %0 = "stablehlo.dynamic_slice"(%arg0, %cst) {slice_sizes = dense<4> : tensor<1xi64>} : (tensor<12xi32>, tensor<i64>) -> tensor<4xi32>
-  // CHECK: emitc.call "emitc::stablehlo::dynamic_slice"(%arg1, %cst, %cst_0) {args = [0 : index, 1 : index, 2 : index, dense<[4, 2]> : tensor<2xi64>], template_args = [tensor<4x2xi32>]} : (tensor<8x7xi32>, tensor<i64>, tensor<i64>) -> tensor<4x2xi32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::dynamic_slice"(%arg1, %cst, %cst_0) {args = [0 : index, 1 : index, 2 : index, dense<[4, 2]> : tensor<2xi64>], template_args = [tensor<4x2xi32>]} : (tensor<8x7xi32>, tensor<i64>, tensor<i64>) -> tensor<4x2xi32>
   %1 = "stablehlo.dynamic_slice"(%arg1, %cst, %cst_0) {slice_sizes = dense<[4, 2]> : tensor<2xi64>} : (tensor<8x7xi32>, tensor<i64>, tensor<i64>) -> tensor<4x2xi32>
   return
 }
@@ -265,9 +265,9 @@ func.func @stablehlo_dynamic_update_slice(%arg0: tensor<12xi32>, %arg1: tensor<8
   %cst_0 = "arith.constant"() {value = dense<3> : tensor<i64>} : () -> tensor<i64>
   %cst_1 = "arith.constant"() {value = dense<1> : tensor<4xi32>} : () -> tensor<4xi32>
   %cst_2 = "arith.constant"() {value = dense<1> : tensor<2x4xi32>} : () -> tensor<2x4xi32>
-  // CHECK: emitc.call "emitc::stablehlo::dynamic_update_slice"(%arg0, %cst_1, %cst) {template_args = [tensor<4xi32>]}
+  // CHECK: emitc.call_opaque "emitc::stablehlo::dynamic_update_slice"(%arg0, %cst_1, %cst) {template_args = [tensor<4xi32>]}
   %0 = "stablehlo.dynamic_update_slice"(%arg0, %cst_1, %cst) : (tensor<12xi32>, tensor<4xi32>, tensor<i64>) -> tensor<12xi32>
-  // CHECK: emitc.call "emitc::stablehlo::dynamic_update_slice"(%arg1, %cst_2, %cst, %cst_0) {template_args = [tensor<2x4xi32>]}
+  // CHECK: emitc.call_opaque "emitc::stablehlo::dynamic_update_slice"(%arg1, %cst_2, %cst, %cst_0) {template_args = [tensor<2x4xi32>]}
   %1 = "stablehlo.dynamic_update_slice"(%arg1, %cst_2, %cst, %cst_0) : (tensor<8x7xi32>, tensor<2x4xi32>, tensor<i64>, tensor<i64>) -> tensor<8x7xi32>
   return
 }
@@ -276,37 +276,37 @@ func.func @stablehlo_dynamic_update_slice(%arg0: tensor<12xi32>, %arg1: tensor<8
 // Other ops
 
 func.func @stablehlo_batch_norm_inference(%arg0: tensor<4x2xf32>, %arg1: tensor<2xf32>, %arg2: tensor<2xf32>, %arg3: tensor<2xf32>, %arg4: tensor<2xf32>) -> tensor<4x2xf32> {
-  // CHECK: emitc.call "emitc::stablehlo::batch_norm_inference"(%arg0, %arg1, %arg2, %arg3, %arg4) {args = [0 : index, 1 : index, 2 : index, 3 : index, 4 : index, 1.000000e-03 : f32, 1], template_args = [tensor<4x2xf32>, tensor<2xf32>]} : (tensor<4x2xf32>, tensor<2xf32>, tensor<2xf32>, tensor<2xf32>, tensor<2xf32>) -> tensor<4x2xf32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::batch_norm_inference"(%arg0, %arg1, %arg2, %arg3, %arg4) {args = [0 : index, 1 : index, 2 : index, 3 : index, 4 : index, 1.000000e-03 : f32, 1], template_args = [tensor<4x2xf32>, tensor<2xf32>]} : (tensor<4x2xf32>, tensor<2xf32>, tensor<2xf32>, tensor<2xf32>, tensor<2xf32>) -> tensor<4x2xf32>
   %0 = "stablehlo.batch_norm_inference"(%arg0, %arg1, %arg2, %arg3, %arg4) {epsilon = 0.001 : f32, feature_index = 1 : i64} : (tensor<4x2xf32>, tensor<2xf32>, tensor<2xf32>, tensor<2xf32>, tensor<2xf32>) -> tensor<4x2xf32>
   return %0 : tensor<4x2xf32>
 }
 
 func.func @stablehlo_bitcast_convert(%arg0: tensor<ui32>) -> tensor<i32> {
-  // CHECK: emitc.call "emitc::stablehlo::bitcast_convert"(%arg0) {template_args = [tensor<i32>]} : (tensor<ui32>) -> tensor<i32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::bitcast_convert"(%arg0) {template_args = [tensor<i32>]} : (tensor<ui32>) -> tensor<i32>
   %0 = "stablehlo.bitcast_convert"(%arg0) : (tensor<ui32>) -> tensor<i32>
   return %0 : tensor<i32>
 }
 
 func.func @stablehlo_broadcast_in_dim(%arg0: tensor<i32>) -> tensor<3xi32> {
-  // CHECK: emitc.call "emitc::stablehlo::broadcast_in_dim"(%arg0) {args = [0 : index, dense<> : tensor<0xi64>], template_args = [tensor<3xi32>]} : (tensor<i32>) -> tensor<3xi32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::broadcast_in_dim"(%arg0) {args = [0 : index, dense<> : tensor<0xi64>], template_args = [tensor<3xi32>]} : (tensor<i32>) -> tensor<3xi32>
   %0 = "stablehlo.broadcast_in_dim"(%arg0) {broadcast_dimensions = dense<> : tensor<0xi64>} : (tensor<i32>) -> tensor<3xi32>
   return %0 : tensor<3xi32>
 }
 
 func.func @stablehlo_clamp(%arg0: tensor<2x1xf32>, %arg1: tensor<2x1xf32>, %arg2: tensor<2x1xf32>) -> tensor<2x1xf32> {
-  // CHECK: emitc.call "emitc::stablehlo::clamp"(%arg0, %arg1, %arg2) {template_args = [tensor<2x1xf32>, tensor<2x1xf32>, tensor<2x1xf32>]} : (tensor<2x1xf32>, tensor<2x1xf32>, tensor<2x1xf32>) -> tensor<2x1xf32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::clamp"(%arg0, %arg1, %arg2) {template_args = [tensor<2x1xf32>, tensor<2x1xf32>, tensor<2x1xf32>]} : (tensor<2x1xf32>, tensor<2x1xf32>, tensor<2x1xf32>) -> tensor<2x1xf32>
   %0 = "stablehlo.clamp"(%arg0, %arg1, %arg2) : (tensor<2x1xf32>, tensor<2x1xf32>, tensor<2x1xf32>) -> tensor<2x1xf32>
   return %0 : tensor<2x1xf32>
 }
 
 func.func @stablehlo_clamp_broadcast(%arg0: tensor<i32>, %arg1: tensor<4x2x1xi32>, %arg2: tensor<i32>) -> tensor<4x2x1xi32> {
-  // CHECK: emitc.call "emitc::stablehlo::clamp"(%arg0, %arg1, %arg2) {template_args = [tensor<i32>, tensor<4x2x1xi32>, tensor<i32>]} : (tensor<i32>, tensor<4x2x1xi32>, tensor<i32>) -> tensor<4x2x1xi32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::clamp"(%arg0, %arg1, %arg2) {template_args = [tensor<i32>, tensor<4x2x1xi32>, tensor<i32>]} : (tensor<i32>, tensor<4x2x1xi32>, tensor<i32>) -> tensor<4x2x1xi32>
   %0 = "stablehlo.clamp"(%arg0, %arg1, %arg2) : (tensor<i32>, tensor<4x2x1xi32>, tensor<i32>) -> tensor<4x2x1xi32>
   return %0 : tensor<4x2x1xi32>
 }
 
 func.func @stablehlo_concaternate(%arg0: tensor<1xf32>, %arg1: tensor<2xf32>) -> tensor<3xf32> {
-  // CHECK: emitc.call "emitc::stablehlo::concatenate"(%arg0, %arg1) {template_args = [0, tensor<3xf32>]} : (tensor<1xf32>, tensor<2xf32>) -> tensor<3xf32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::concatenate"(%arg0, %arg1) {template_args = [0, tensor<3xf32>]} : (tensor<1xf32>, tensor<2xf32>) -> tensor<3xf32>
   %0 = "stablehlo.concatenate"(%arg0, %arg1) {dimension = 0 : i64} : (tensor<1xf32>, tensor<2xf32>) -> tensor<3xf32>
   return %0 : tensor<3xf32>
 }
@@ -317,7 +317,7 @@ func.func @stablehlo_concaternate(%arg0: tensor<1xf32>, %arg1: tensor<2xf32>) ->
 // https://github.com/tensorflow/mlir-hlo/blob/2ee7bcc4b5d5a8afc7cbd414d10cdfeda97b7a54/tests/Dialect/stablehlo/hlo-legalize-to-lhlo.mlir#L475-L501
 func.func @stablehlo_conv(%arg0: tensor<3x2x4x3xf32>, %arg1 : tensor<2x2x3x4xf32>) -> tensor<2x1x2x3xf32> {
   %c0 = arith.constant 0 : index
-  // CHECK: emitc.call "emitc::stablehlo::convolution"(%arg1, %arg0)
+  // CHECK: emitc.call_opaque "emitc::stablehlo::convolution"(%arg1, %arg0)
   %out = "stablehlo.convolution"(%arg1, %arg0) {
     batch_group_count = 1 : i64,
     dimension_numbers = #stablehlo.conv<raw
@@ -340,13 +340,13 @@ func.func @stablehlo_conv(%arg0: tensor<3x2x4x3xf32>, %arg1 : tensor<2x2x3x4xf32
 }
 
 func.func @stablehlo_dot(%arg0: tensor<512x512xf32>) -> tensor<512x512xf32> {
-  // CHECK: emitc.call "emitc::stablehlo::dot"(%arg0, %arg0) {template_args = [tensor<512x512xf32>]} : (tensor<512x512xf32>, tensor<512x512xf32>) -> tensor<512x512xf32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::dot"(%arg0, %arg0) {template_args = [tensor<512x512xf32>]} : (tensor<512x512xf32>, tensor<512x512xf32>) -> tensor<512x512xf32>
   %0 = "stablehlo.dot"(%arg0, %arg0) : (tensor<512x512xf32>, tensor<512x512xf32>) -> tensor<512x512xf32>
   return %0 : tensor<512x512xf32>
 }
 
 func.func @stablehlo_pad(%arg0: tensor<2x3xf32>, %arg1: tensor<f32>) -> tensor<4x7xf32> {
-  // CHECK: emitc.call "emitc::stablehlo::pad"(%arg0, %arg1) {args = [0 : index, 1 : index, dense<-1> : tensor<2xi64>, dense<1> : tensor<2xi64>, dense<2> : tensor<2xi64>], template_args = [tensor<4x7xf32>]} : (tensor<2x3xf32>, tensor<f32>) -> tensor<4x7xf32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::pad"(%arg0, %arg1) {args = [0 : index, 1 : index, dense<-1> : tensor<2xi64>, dense<1> : tensor<2xi64>, dense<2> : tensor<2xi64>], template_args = [tensor<4x7xf32>]} : (tensor<2x3xf32>, tensor<f32>) -> tensor<4x7xf32>
   %0 = "stablehlo.pad"(%arg0, %arg1) {
     edge_padding_low = dense<-1> : tensor<2xi64>,
     edge_padding_high = dense<1> : tensor<2xi64>,
@@ -364,21 +364,21 @@ func.func @stablehlo_reduce(%arg0 : tensor<2x1000xf32>, %arg1 : tensor<f32>, %ar
   // CHECK: "emitc::stablehlo::max"(%arg0, %arg2) : (tensor<f32>, tensor<f32>) -> tensor<f32>
   // CHECK: "emitc::stablehlo::min"(%arg1, %arg3) : (tensor<i32>, tensor<i32>) -> tensor<i32>
   
-  // CHECK: emitc.call "emitc::stablehlo::reduce"(%arg0, %arg1) {args = [0 : index, 1 : index, dense<1> : tensor<1xi64>, @stablehlo_reduce_lambda_0], template_args = [tensor<2xf32>, 1]} : (tensor<2x1000xf32>, tensor<f32>) -> tensor<2xf32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::reduce"(%arg0, %arg1) {args = [0 : index, 1 : index, dense<1> : tensor<1xi64>, @stablehlo_reduce_lambda_0], template_args = [tensor<2xf32>, 1]} : (tensor<2x1000xf32>, tensor<f32>) -> tensor<2xf32>
   %0 = "stablehlo.reduce"(%arg0, %arg1) ({
     ^bb0(%arg4: tensor<f32>, %arg5: tensor<f32>):
       %1 = stablehlo.add %arg4, %arg5 : tensor<f32>
       "stablehlo.return"(%1) : (tensor<f32>) -> ()
     }) {dimensions = dense<1> : tensor<1xi64>} : (tensor<2x1000xf32>, tensor<f32>) -> tensor<2xf32>
   
-  // CHECK: emitc.call "emitc::stablehlo::reduce"(%arg2, %arg3) {args = [0 : index, 1 : index, dense<1> : tensor<1xi64>, @stablehlo_reduce_lambda_1], template_args = [tensor<2xi32>, 1]} : (tensor<2x1000xi32>, tensor<i32>) -> tensor<2xi32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::reduce"(%arg2, %arg3) {args = [0 : index, 1 : index, dense<1> : tensor<1xi64>, @stablehlo_reduce_lambda_1], template_args = [tensor<2xi32>, 1]} : (tensor<2x1000xi32>, tensor<i32>) -> tensor<2xi32>
   %1 = "stablehlo.reduce"(%arg2, %arg3) ({
     ^bb0(%arg4: tensor<i32>, %arg5: tensor<i32>):
       %2 = stablehlo.maximum %arg4, %arg5 : tensor<i32>
       "stablehlo.return"(%2) : (tensor<i32>) -> ()
     }) {dimensions = dense<1> : tensor<1xi64>} : (tensor<2x1000xi32>, tensor<i32>) -> tensor<2xi32>
   
-  // CHECK: emitc.call "emitc::stablehlo::reduce"(%arg0, %arg2, %arg1, %arg3) {args = [0 : index, 1 : index, 2 : index, 3 : index, dense<1> : tensor<1xi64>, @stablehlo_reduce_lambda_2], template_args = [tensor<2xf32>, tensor<2xi32>, 1]} : (tensor<2x1000xf32>, tensor<2x1000xi32>, tensor<f32>, tensor<i32>) -> (tensor<2xf32>, tensor<2xi32>)
+  // CHECK: emitc.call_opaque "emitc::stablehlo::reduce"(%arg0, %arg2, %arg1, %arg3) {args = [0 : index, 1 : index, 2 : index, 3 : index, dense<1> : tensor<1xi64>, @stablehlo_reduce_lambda_2], template_args = [tensor<2xf32>, tensor<2xi32>, 1]} : (tensor<2x1000xf32>, tensor<2x1000xi32>, tensor<f32>, tensor<i32>) -> (tensor<2xf32>, tensor<2xi32>)
   %2:2 = stablehlo.reduce(%arg0 init: %arg1), (%arg2 init: %arg3) across dimensions = [1] : (tensor<2x1000xf32>, tensor<2x1000xi32>, tensor<f32>, tensor<i32>) -> (tensor<2xf32>, tensor<2xi32>)
      reducer(%arg4: tensor<f32>, %arg5: tensor<f32>) (%arg6: tensor<i32>, %arg7: tensor<i32>)  {
       %2 = stablehlo.maximum %arg4, %arg5 : tensor<f32>
@@ -392,7 +392,7 @@ func.func @stablehlo_reduce(%arg0 : tensor<2x1000xf32>, %arg1 : tensor<f32>, %ar
 func.func @stablehlo_reduce_window(%arg0 : tensor<2x114x114x64xf32>, %arg1 : tensor<f32>) -> tensor<2x56x56x64xf32> {
   // CHECK: func @stablehlo_reduce_window_lambda_0(%arg0: tensor<f32>, %arg1: tensor<f32>)
   // CHECK: "emitc::stablehlo::max"
-  // CHECK: emitc.call "emitc::stablehlo::reduce_window"(%arg0, %arg1) {args = [0 : index, 1 : index, dense<[1, 3, 3, 1]> : tensor<4xi64>, dense<[1, 2, 2, 1]> : tensor<4xi64>, dense<1> : tensor<4xi64>, dense<1> : tensor<4xi64>, dense<0> : tensor<8xi64>, @stablehlo_reduce_window_lambda_0], template_args = [tensor<2x56x56x64xf32>]}
+  // CHECK: emitc.call_opaque "emitc::stablehlo::reduce_window"(%arg0, %arg1) {args = [0 : index, 1 : index, dense<[1, 3, 3, 1]> : tensor<4xi64>, dense<[1, 2, 2, 1]> : tensor<4xi64>, dense<1> : tensor<4xi64>, dense<1> : tensor<4xi64>, dense<0> : tensor<8xi64>, @stablehlo_reduce_window_lambda_0], template_args = [tensor<2x56x56x64xf32>]}
   %0 = "stablehlo.reduce_window"(%arg0, %arg1) ( {
     ^bb0(%arg2: tensor<f32>, %arg3: tensor<f32>):  // no predecessors
       %516 = stablehlo.maximum %arg2, %arg3 : tensor<f32>
@@ -403,13 +403,13 @@ func.func @stablehlo_reduce_window(%arg0 : tensor<2x114x114x64xf32>, %arg1 : ten
 }
 
 func.func @stablehlo_reshape(%arg0: tensor<12xf32>) -> tensor<2x3x2xf32> {
-  // CHECK: emitc.call "emitc::stablehlo::reshape"(%arg0) {template_args = [tensor<2x3x2xf32>]} : (tensor<12xf32>) -> tensor<2x3x2xf32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::reshape"(%arg0) {template_args = [tensor<2x3x2xf32>]} : (tensor<12xf32>) -> tensor<2x3x2xf32>
   %0 = "stablehlo.reshape"(%arg0) : (tensor<12xf32>) -> tensor<2x3x2xf32>
   return %0 : tensor<2x3x2xf32>
 }
 
 func.func @stablehlo_select(%arg0: tensor<2xf32>, %arg1: tensor<2xf32>, %arg2: tensor<2xi1>) -> tensor<2xf32> {
-  // CHECK: emitc.call "emitc::stablehlo::select"(%arg2, %arg0, %arg1) : (tensor<2xi1>, tensor<2xf32>, tensor<2xf32>) -> tensor<2xf32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::select"(%arg2, %arg0, %arg1) : (tensor<2xi1>, tensor<2xf32>, tensor<2xf32>) -> tensor<2xf32>
   %1 = "stablehlo.select"(%arg2, %arg0, %arg1) : (tensor<2xi1>, tensor<2xf32>, tensor<2xf32>) -> tensor<2xf32>
   return %1 : tensor<2xf32>
 }
@@ -420,7 +420,7 @@ func.func @select_scalar_pred(%arg0: tensor<i1>, %arg1: tensor<2x3xi32>, %arg2:
 }
 
 func.func @stablehlo_transpose(%arg0: tensor<2x3x4xf32>) -> tensor<4x3x2xf32> {
-  // CHECK: emitc.call "emitc::stablehlo::transpose"(%arg0) {args = [0 : index, dense<[2, 1, 0]> : tensor<3xi64>], template_args = [tensor<4x3x2xf32>]} : (tensor<2x3x4xf32>) -> tensor<4x3x2xf32>
+  // CHECK: emitc.call_opaque "emitc::stablehlo::transpose"(%arg0) {args = [0 : index, dense<[2, 1, 0]> : tensor<3xi64>], template_args = [tensor<4x3x2xf32>]} : (tensor<2x3x4xf32>) -> tensor<4x3x2xf32>
   %0 = "stablehlo.transpose"(%arg0) {permutation = dense<[2, 1, 0]> : tensor<3xi64>} : (tensor<2x3x4xf32>) -> tensor<4x3x2xf32>
   return %0 : tensor<4x3x2xf32>
 }
@@ -432,14 +432,14 @@ func.func @stablehlo_rng_uniform() -> () {
   %cst_0 = "arith.constant"() {value = dense<100> : tensor<i32>} : () -> tensor<i32>
   %cst_1 = "arith.constant"() {value = dense<2> : tensor<1xi64>} : () -> tensor<1xi64>
 
-  // TODO: Fix or drop test #: emitc.call "emitc::stablehlo::rng_uniform"(%cst, %cst_0, %cst_1) {template_args = [tensor<2xi32>]} : (tensor<i32>, tensor<i32>, tensor<1xi64>) -> tensor<2xi32>
+  // TODO: Fix or drop test #: emitc.call_opaque "emitc::stablehlo::rng_uniform"(%cst, %cst_0, %cst_1) {template_args = [tensor<2xi32>]} : (tensor<i32>, tensor<i32>, tensor<1xi64>) -> tensor<2xi32>
   //%0 = "stablehlo.rng"(%cst, %cst_0, %cst_1) {rng_distribution = #stablehlo.rng_distribution<UNIFORM>}: (tensor<i32>, tensor<i32>, tensor<1xi64>) -> tensor<2xi32>
   
   %cst_2 = "arith.constant"() {value = dense<-100.0> : tensor<f32>} : () -> tensor<f32>
   %cst_3 = "arith.constant"() {value = dense<100.0> : tensor<f32>} : () -> tensor<f32>
   %cst_4 = "arith.constant"() {value = dense<17> : tensor<1xi64>} : () -> tensor<1xi64>
 
-  // TODO: Fix or drop test #: emitc.call "emitc::stablehlo::rng_uniform"(%cst_2, %cst_3, %cst_4) {template_args = [tensor<17xf32>]} : (tensor<f32>, tensor<f32>, tensor<1xi64>) -> tensor<17xf32>
+  // TODO: Fix or drop test #: emitc.call_opaque "emitc::stablehlo::rng_uniform"(%cst_2, %cst_3, %cst_4) {template_args = [tensor<17xf32>]} : (tensor<f32>, tensor<f32>, tensor<1xi64>) -> tensor<17xf32>
   //%1 = "stablehlo.rng"(%cst_2, %cst_3, %cst_4) {rng_distribution = #stablehlo.rng_distribution<UNIFORM>}: (tensor<f32>, tensor<f32>, tensor<1xi64>) -> tensor<17xf32>
   return
 }
diff --git a/test/Conversion/tensor-to-emitc.mlir b/test/Conversion/tensor-to-emitc.mlir
index e10c91c0..7643dc0f 100644
--- a/test/Conversion/tensor-to-emitc.mlir
+++ b/test/Conversion/tensor-to-emitc.mlir
@@ -16,9 +16,9 @@ func.func @std_extract_element(%arg0: tensor<i32>, %arg1: tensor<2xi32> ) -> ()
 // CHECK-LABEL: func @std_extract_element
 //  CHECK-NEXT: constant 0 : index
 //  CHECK-NEXT: constant 1 : index
-//  CHECK-NEXT: emitc.call "emitc::tensor::extract"(%arg0) : (tensor<i32>) -> i32
-//  CHECK-NEXT: emitc.call "emitc::tensor::extract"(%arg1, %c0) : (tensor<2xi32>, index) -> i32
-//  CHECK-NEXT: emitc.call "emitc::tensor::extract"(%arg1, %c1) : (tensor<2xi32>, index) -> i32
+//  CHECK-NEXT: emitc.call_opaque "emitc::tensor::extract"(%arg0) : (tensor<i32>) -> i32
+//  CHECK-NEXT: emitc.call_opaque "emitc::tensor::extract"(%arg1, %c0) : (tensor<2xi32>, index) -> i32
+//  CHECK-NEXT: emitc.call_opaque "emitc::tensor::extract"(%arg1, %c1) : (tensor<2xi32>, index) -> i32
 
 // CPP-LABEL: void std_extract_element(Tensor<int32_t> v1, Tensor<int32_t, 2> v2)
 //  CPP-NEXT: size_t v3 = 0;
@@ -32,7 +32,7 @@ func.func @splat_op(%s : f32) -> tensor<8xf32> {
   return %t : tensor<8xf32>
 }
 // CHECK-LABEL: func @splat_op
-//  CHECK-NEXT: emitc.call "emitc::tensor::splat"(%arg0) {template_args = [tensor<8xf32>]} : (f32) -> tensor<8xf32>
+//  CHECK-NEXT: emitc.call_opaque "emitc::tensor::splat"(%arg0) {template_args = [tensor<8xf32>]} : (f32) -> tensor<8xf32>
 
 // CPP-LABEL: Tensor<float, 8> splat_op(float v1)
 //  CPP-NEXT: emitc::tensor::splat<Tensor<float, 8>>(v1)
diff --git a/test/Conversion/tosa-to-emitc.mlir b/test/Conversion/tosa-to-emitc.mlir
index aa15c83b..7dbb183d 100644
--- a/test/Conversion/tosa-to-emitc.mlir
+++ b/test/Conversion/tosa-to-emitc.mlir
@@ -16,87 +16,87 @@ func.func @test_const(%arg0 : index) -> tensor<4xi32> {
 // Unary elementwise ops
 
 func.func @test_abs(%arg0: tensor<13x21x3xf32>) -> tensor<13x21x3xf32> {
-  // CHECK: emitc.call "emitc::tosa::abs"(%arg0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
+  // CHECK: emitc.call_opaque "emitc::tosa::abs"(%arg0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   %0 = "tosa.abs"(%arg0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   return %0 : tensor<13x21x3xf32>
 }
 
 // CHECK-LABEL: cast
 func.func @test_cast(%arg0: tensor<13x21x3xi32>) -> tensor<13x21x3xf32> {
-  // CHECK: %0 = emitc.call "emitc::tosa::cast"(%arg0) {template_args = [tensor<13x21x3xf32>]} : (tensor<13x21x3xi32>) -> tensor<13x21x3xf32>
+  // CHECK: %0 = emitc.call_opaque "emitc::tosa::cast"(%arg0) {template_args = [tensor<13x21x3xf32>]} : (tensor<13x21x3xi32>) -> tensor<13x21x3xf32>
   %0 = "tosa.cast"(%arg0) : (tensor<13x21x3xi32>) -> tensor<13x21x3xf32>
   return %0 : tensor<13x21x3xf32>
 }
 
 func.func @test_ceil(%arg0: tensor<13x21x3xf32>) -> tensor<13x21x3xf32> {
-  // CHECK: emitc.call "emitc::tosa::ceil"(%arg0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
+  // CHECK: emitc.call_opaque "emitc::tosa::ceil"(%arg0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   %0 = "tosa.ceil"(%arg0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   return %0 : tensor<13x21x3xf32>
 }
 
 func.func @test_clamp0(%arg0: tensor<13x21x3xf32>) -> tensor<13x21x3xf32> {
-  // CHECK: %0 = emitc.call "emitc::tosa::clamp"(%arg0) {args = [0 : index, 0.000000e+00 : f32, 1.000000e+00 : f32]} : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
+  // CHECK: %0 = emitc.call_opaque "emitc::tosa::clamp"(%arg0) {args = [0 : index, 0.000000e+00 : f32, 1.000000e+00 : f32]} : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   %0 = "tosa.clamp"(%arg0) {min_fp = 0.0 : f32, max_fp = 1.0 : f32, min_int = -2 : i64, max_int = 2 : i64} : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   return %0 : tensor<13x21x3xf32>
 }
 
 func.func @test_clamp1(%arg0: tensor<13x21x3xi32>) -> tensor<13x21x3xi32> {
-  // CHECK: %0 = emitc.call "emitc::tosa::clamp"(%arg0) {args = [0 : index, -2 : i32, 2 : i32]} : (tensor<13x21x3xi32>) -> tensor<13x21x3xi32>
+  // CHECK: %0 = emitc.call_opaque "emitc::tosa::clamp"(%arg0) {args = [0 : index, -2 : i32, 2 : i32]} : (tensor<13x21x3xi32>) -> tensor<13x21x3xi32>
   %0 = "tosa.clamp"(%arg0) {min_fp = 0.0 : f32, max_fp = 1.0 : f32, min_int = -2 : i64, max_int = 2 : i64} : (tensor<13x21x3xi32>) -> tensor<13x21x3xi32>
   return %0 : tensor<13x21x3xi32>
 }
 
 func.func @test_clz(%arg0: tensor<13x21x3xi32>) -> tensor<13x21x3xi32> {
-  // CHECK: emitc.call "emitc::tosa::clz"(%arg0) : (tensor<13x21x3xi32>) -> tensor<13x21x3xi32>
+  // CHECK: emitc.call_opaque "emitc::tosa::clz"(%arg0) : (tensor<13x21x3xi32>) -> tensor<13x21x3xi32>
   %0 = "tosa.clz"(%arg0) : (tensor<13x21x3xi32>) -> tensor<13x21x3xi32>
   return %0 : tensor<13x21x3xi32>
 }
 
 func.func @test_exp(%arg0: tensor<13x21x3xf32>) -> tensor<13x21x3xf32> {
-  // CHECK: emitc.call "emitc::tosa::exp"(%arg0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
+  // CHECK: emitc.call_opaque "emitc::tosa::exp"(%arg0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   %0 = "tosa.exp"(%arg0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   return %0 : tensor<13x21x3xf32>
 }
 
 func.func @test_floor(%arg0: tensor<13x21x3xf32>) -> tensor<13x21x3xf32> {
-  // CHECK: emitc.call "emitc::tosa::floor"(%arg0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
+  // CHECK: emitc.call_opaque "emitc::tosa::floor"(%arg0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   %0 = "tosa.floor"(%arg0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   return %0 : tensor<13x21x3xf32>
 }
 
 func.func @test_log(%arg0: tensor<13x21x3xf32>) -> tensor<13x21x3xf32> {
-  // CHECK: emitc.call "emitc::tosa::log"(%arg0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
+  // CHECK: emitc.call_opaque "emitc::tosa::log"(%arg0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   %0 = "tosa.log"(%arg0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   return %0 : tensor<13x21x3xf32>
 }
 
 func.func @test_negate(%arg0: tensor<13x21x3xf32>) -> tensor<13x21x3xf32> {
-  // CHECK: %0 = emitc.call "emitc::tosa::negate"(%arg0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
+  // CHECK: %0 = emitc.call_opaque "emitc::tosa::negate"(%arg0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   %0 = "tosa.negate"(%arg0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   return %0 : tensor<13x21x3xf32>
 }
 
 func.func @test_reciprocal(%arg0: tensor<13x21x3xf32>) -> tensor<13x21x3xf32> {
-  // CHECK: %0 = emitc.call "emitc::tosa::reciprocal"(%arg0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
+  // CHECK: %0 = emitc.call_opaque "emitc::tosa::reciprocal"(%arg0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   %0 = "tosa.reciprocal"(%arg0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   return %0 : tensor<13x21x3xf32>
 }
 
 func.func @test_rescale(%arg0: tensor<13x21x3xui8>) -> tensor<13x21x3xi8> {
-  // CHECK: %0 = emitc.call "emitc::tosa::rescale"(%arg0) {args = [0 : index, 127 : i32, -1 : i32, array<i32: 1073741824>, array<i32: 30>, true, false, false], template_args = [tensor<13x21x3xi8>, 1 : i32]} : (tensor<13x21x3xui8>) -> tensor<13x21x3xi8>
+  // CHECK: %0 = emitc.call_opaque "emitc::tosa::rescale"(%arg0) {args = [0 : index, 127 : i32, -1 : i32, array<i32: 1073741824>, array<i32: 30>, true, false, false], template_args = [tensor<13x21x3xi8>, 1 : i32]} : (tensor<13x21x3xui8>) -> tensor<13x21x3xi8>
   %0 = "tosa.rescale"(%arg0) {double_round = false, input_zp = 127 : i32, multiplier = array<i32: 1073741824>, output_zp = -1 : i32, per_channel = false, scale32 = true, shift = array<i32: 30>} : (tensor<13x21x3xui8>) -> tensor<13x21x3xi8>
   return %0 : tensor<13x21x3xi8>
 }
 
 func.func @test_rsqrt(%arg0: tensor<13x21x3xf32>) -> tensor<13x21x3xf32> {
-  // CHECK: %0 = emitc.call "emitc::sqrt"(%arg0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
-  // CHECK: %1 = emitc.call "emitc::tosa::reciprocal"(%0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
+  // CHECK: %0 = emitc.call_opaque "emitc::sqrt"(%arg0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
+  // CHECK: %1 = emitc.call_opaque "emitc::tosa::reciprocal"(%0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   %0 = "tosa.rsqrt"(%arg0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   return %0 : tensor<13x21x3xf32>
 }
 
 func.func @test_tanh(%arg0: tensor<13x21x3xf32>) -> tensor<13x21x3xf32> {
-  // CHECK: %0 = emitc.call "emitc::tosa::tanh"(%arg0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
+  // CHECK: %0 = emitc.call_opaque "emitc::tosa::tanh"(%arg0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   %0 = "tosa.tanh"(%arg0) : (tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   return %0 : tensor<13x21x3xf32>
 }
@@ -104,123 +104,123 @@ func.func @test_tanh(%arg0: tensor<13x21x3xf32>) -> tensor<13x21x3xf32> {
 // Binary elementwise ops
 
 func.func @test_add(%arg0: tensor<13x21x1xf32>, %arg1: tensor<13x21x3xf32>) -> tensor<13x21x3xf32> {
-  // CHECK: emitc.call "emitc::broadcast_in_dim"(%arg0) {args = [0 : index, dense<[0, 1, 2]> : tensor<3xi64>], template_args = [tensor<13x21x3xf32>]} : (tensor<13x21x1xf32>) -> tensor<13x21x3xf32>
-  // CHECK: emitc.call "emitc::tosa::add"(%0, %arg1) : (tensor<13x21x3xf32>, tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
+  // CHECK: emitc.call_opaque "emitc::broadcast_in_dim"(%arg0) {args = [0 : index, dense<[0, 1, 2]> : tensor<3xi64>], template_args = [tensor<13x21x3xf32>]} : (tensor<13x21x1xf32>) -> tensor<13x21x3xf32>
+  // CHECK: emitc.call_opaque "emitc::tosa::add"(%0, %arg1) : (tensor<13x21x3xf32>, tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   %0 = "tosa.add"(%arg0, %arg1) : (tensor<13x21x1xf32>, tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   return %0 : tensor<13x21x3xf32>
 }
 
 // ArithmeticRightShiftOp: no broadcast
 func.func @test_arithmetic_right_shift1(%arg0: tensor<13x21x3xi32>, %arg1: tensor<13x21x3xi32>) -> tensor<13x21x3xi32> {
-  // CHECK: emitc.call "emitc::tosa::arithmetic_right_shift"(%arg0, %arg1) {args = [0 : index, 1 : index, false]} : (tensor<13x21x3xi32>, tensor<13x21x3xi32>) -> tensor<13x21x3xi32>
+  // CHECK: emitc.call_opaque "emitc::tosa::arithmetic_right_shift"(%arg0, %arg1) {args = [0 : index, 1 : index, false]} : (tensor<13x21x3xi32>, tensor<13x21x3xi32>) -> tensor<13x21x3xi32>
   %0 = "tosa.arithmetic_right_shift"(%arg0, %arg1) { round = false } : (tensor<13x21x3xi32>, tensor<13x21x3xi32>) -> tensor<13x21x3xi32>
   return %0 : tensor<13x21x3xi32>
 }
 
 // ArithmeticRightShiftOp: First operand needs to be broadcasted
 func.func @test_arithmetic_right_shift2(%arg0: tensor<13x21x1xi32>, %arg1: tensor<13x21x3xi32>) -> tensor<13x21x3xi32> {
-  // CHECK: emitc.call "emitc::broadcast_in_dim"(%arg0) {args = [0 : index, dense<[0, 1, 2]> : tensor<3xi64>], template_args = [tensor<13x21x3xi32>]} : (tensor<13x21x1xi32>) -> tensor<13x21x3xi32>
-  // CHECK: emitc.call "emitc::tosa::arithmetic_right_shift"(%0, %arg1) {args = [0 : index, 1 : index, true]} : (tensor<13x21x3xi32>, tensor<13x21x3xi32>) -> tensor<13x21x3xi32>
+  // CHECK: emitc.call_opaque "emitc::broadcast_in_dim"(%arg0) {args = [0 : index, dense<[0, 1, 2]> : tensor<3xi64>], template_args = [tensor<13x21x3xi32>]} : (tensor<13x21x1xi32>) -> tensor<13x21x3xi32>
+  // CHECK: emitc.call_opaque "emitc::tosa::arithmetic_right_shift"(%0, %arg1) {args = [0 : index, 1 : index, true]} : (tensor<13x21x3xi32>, tensor<13x21x3xi32>) -> tensor<13x21x3xi32>
   %0 = "tosa.arithmetic_right_shift"(%arg0, %arg1) { round = true } : (tensor<13x21x1xi32>, tensor<13x21x3xi32>) -> tensor<13x21x3xi32>
   return %0 : tensor<13x21x3xi32>
 }
 
 func.func @test_equal(%arg0: tensor<13x21x3xi32>, %arg1: tensor<13x21x3xi32>) -> tensor<13x21x3xi1> {
-  // CHECK: emitc.call "emitc::tosa::equal"(%arg0, %arg1) {template_args = [tensor<13x21x3xi1>]} : (tensor<13x21x3xi32>, tensor<13x21x3xi32>) -> tensor<13x21x3xi1>
+  // CHECK: emitc.call_opaque "emitc::tosa::equal"(%arg0, %arg1) {template_args = [tensor<13x21x3xi1>]} : (tensor<13x21x3xi32>, tensor<13x21x3xi32>) -> tensor<13x21x3xi1>
   %0 = "tosa.equal"(%arg0, %arg1) : (tensor<13x21x3xi32>, tensor<13x21x3xi32>) -> tensor<13x21x3xi1>
   return %0 : tensor<13x21x3xi1>
 }
 
 func.func @test_greater_equal(%arg0: tensor<13x21x3xf32>, %arg1: tensor<13x21x3xf32>) -> tensor<13x21x3xi1> {
-  // CHECK: emitc.call "emitc::tosa::greater_equal"(%arg0, %arg1) {template_args = [tensor<13x21x3xi1>]} : (tensor<13x21x3xf32>, tensor<13x21x3xf32>) -> tensor<13x21x3xi1>
+  // CHECK: emitc.call_opaque "emitc::tosa::greater_equal"(%arg0, %arg1) {template_args = [tensor<13x21x3xi1>]} : (tensor<13x21x3xf32>, tensor<13x21x3xf32>) -> tensor<13x21x3xi1>
   %0 = "tosa.greater_equal"(%arg0, %arg1) : (tensor<13x21x3xf32>, tensor<13x21x3xf32>) -> tensor<13x21x3xi1>
   return %0 : tensor<13x21x3xi1>
 }
 
 func.func @test_logical_left_shift(%arg0: tensor<13x21x1xi32>, %arg1: tensor<13x21x3xi32>) -> tensor<13x21x3xi32> {
-  // CHECK: emitc.call "emitc::broadcast_in_dim"(%arg0) {args = [0 : index, dense<[0, 1, 2]> : tensor<3xi64>], template_args = [tensor<13x21x3xi32>]} : (tensor<13x21x1xi32>) -> tensor<13x21x3xi32>
-  // CHECK: emitc.call "emitc::tosa::logical_left_shift"(%0, %arg1) : (tensor<13x21x3xi32>, tensor<13x21x3xi32>) -> tensor<13x21x3xi32>
+  // CHECK: emitc.call_opaque "emitc::broadcast_in_dim"(%arg0) {args = [0 : index, dense<[0, 1, 2]> : tensor<3xi64>], template_args = [tensor<13x21x3xi32>]} : (tensor<13x21x1xi32>) -> tensor<13x21x3xi32>
+  // CHECK: emitc.call_opaque "emitc::tosa::logical_left_shift"(%0, %arg1) : (tensor<13x21x3xi32>, tensor<13x21x3xi32>) -> tensor<13x21x3xi32>
   %0 = "tosa.logical_left_shift"(%arg0, %arg1) : (tensor<13x21x1xi32>, tensor<13x21x3xi32>) -> tensor<13x21x3xi32>
   return %0 : tensor<13x21x3xi32>
 }
 
 // MulOp: no broadcast
 func.func @test_mul10(%arg0: tensor<13x21x3xf32>, %arg1: tensor<13x21x3xf32>) -> tensor<13x21x3xf32> {
-  // CHECK: emitc.call "emitc::tosa::mul"(%arg0, %arg1) : (tensor<13x21x3xf32>, tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
+  // CHECK: emitc.call_opaque "emitc::tosa::mul"(%arg0, %arg1) : (tensor<13x21x3xf32>, tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   %0 = "tosa.mul"(%arg0, %arg1)  { shift = 0 : i8 } : (tensor<13x21x3xf32>, tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   return %0 : tensor<13x21x3xf32>
 }
 
 // MulOp: First operand needs to be broadcasted
 func.func @test_mul1(%arg0: tensor<13x1x3xi32>, %arg1: tensor<13x21x3xi32>) -> tensor<13x21x3xi32> {
-  // CHECK: emitc.call "emitc::broadcast_in_dim"(%arg0) {args = [0 : index, dense<[0, 1, 2]> : tensor<3xi64>], template_args = [tensor<13x21x3xi32>]} : (tensor<13x1x3xi32>) -> tensor<13x21x3xi32>
-  // CHECK: emitc.call "emitc::tosa::mul"(%0, %arg1) {args = [0 : index, 1 : index, 1 : i8]} : (tensor<13x21x3xi32>, tensor<13x21x3xi32>) -> tensor<13x21x3xi32>
+  // CHECK: emitc.call_opaque "emitc::broadcast_in_dim"(%arg0) {args = [0 : index, dense<[0, 1, 2]> : tensor<3xi64>], template_args = [tensor<13x21x3xi32>]} : (tensor<13x1x3xi32>) -> tensor<13x21x3xi32>
+  // CHECK: emitc.call_opaque "emitc::tosa::mul"(%0, %arg1) {args = [0 : index, 1 : index, 1 : i8]} : (tensor<13x21x3xi32>, tensor<13x21x3xi32>) -> tensor<13x21x3xi32>
   %0 = "tosa.mul"(%arg0, %arg1)  { shift = 1 : i8 } : (tensor<13x1x3xi32>, tensor<13x21x3xi32>) -> tensor<13x21x3xi32>
   return %0 : tensor<13x21x3xi32>
 }
 
 // MulOp: Second operand needs to be broadcasted
 func.func @test_mul2(%arg0: tensor<13x21x3xi32>, %arg1: tensor<13x1x3xi32>) -> tensor<13x21x3xi32> {
-  // CHECK: emitc.call "emitc::broadcast_in_dim"(%arg1) {args = [0 : index, dense<[0, 1, 2]> : tensor<3xi64>], template_args = [tensor<13x21x3xi32>]} : (tensor<13x1x3xi32>) -> tensor<13x21x3xi32>
-  // CHECK: emitc.call "emitc::tosa::mul"(%arg0, %0) {args = [0 : index, 1 : index, 1 : i8]} : (tensor<13x21x3xi32>, tensor<13x21x3xi32>) -> tensor<13x21x3xi32>
+  // CHECK: emitc.call_opaque "emitc::broadcast_in_dim"(%arg1) {args = [0 : index, dense<[0, 1, 2]> : tensor<3xi64>], template_args = [tensor<13x21x3xi32>]} : (tensor<13x1x3xi32>) -> tensor<13x21x3xi32>
+  // CHECK: emitc.call_opaque "emitc::tosa::mul"(%arg0, %0) {args = [0 : index, 1 : index, 1 : i8]} : (tensor<13x21x3xi32>, tensor<13x21x3xi32>) -> tensor<13x21x3xi32>
   %0 = "tosa.mul"(%arg0, %arg1)  { shift = 1 : i8 } : (tensor<13x21x3xi32>, tensor<13x1x3xi32>) -> tensor<13x21x3xi32>
   return %0 : tensor<13x21x3xi32>
 }
 
 // MulOp: Second operand needs to be broadcasted + expanded to two dimensions
 func.func @test_mul3(%arg0: tensor<21x3xi32>, %arg1: tensor<3xi32>) -> tensor<21x3xi32> {
-  // CHECK: emitc.call "emitc::broadcast_in_dim"(%arg1) {args = [0 : index, dense<1> : tensor<1xi64>], template_args = [tensor<21x3xi32>]} : (tensor<3xi32>) -> tensor<21x3xi32>
-  // CHECK: emitc.call "emitc::tosa::mul"(%arg0, %0) {args = [0 : index, 1 : index, 3 : i8]} : (tensor<21x3xi32>, tensor<21x3xi32>) -> tensor<21x3xi32>
+  // CHECK: emitc.call_opaque "emitc::broadcast_in_dim"(%arg1) {args = [0 : index, dense<1> : tensor<1xi64>], template_args = [tensor<21x3xi32>]} : (tensor<3xi32>) -> tensor<21x3xi32>
+  // CHECK: emitc.call_opaque "emitc::tosa::mul"(%arg0, %0) {args = [0 : index, 1 : index, 3 : i8]} : (tensor<21x3xi32>, tensor<21x3xi32>) -> tensor<21x3xi32>
   %0 = "tosa.mul"(%arg0, %arg1)  { shift = 3 : i8 } : (tensor<21x3xi32>, tensor<3xi32>) -> tensor<21x3xi32>
   return %0 : tensor<21x3xi32>
 }
 
 // MulOp: Second operand needs to be broadcasted + expanded to three dimensions
 func.func @test_mul4(%arg0: tensor<13x21x3xi32>, %arg1: tensor<3xi32>) -> tensor<13x21x3xi32> {
-  // CHECK: emitc.call "emitc::broadcast_in_dim"(%arg1) {args = [0 : index, dense<2> : tensor<1xi64>], template_args = [tensor<13x21x3xi32>]} : (tensor<3xi32>) -> tensor<13x21x3xi32>
-  // CHECK: emitc.call "emitc::tosa::mul"(%arg0, %0) {args = [0 : index, 1 : index, 1 : i8]} : (tensor<13x21x3xi32>, tensor<13x21x3xi32>) -> tensor<13x21x3xi32>
+  // CHECK: emitc.call_opaque "emitc::broadcast_in_dim"(%arg1) {args = [0 : index, dense<2> : tensor<1xi64>], template_args = [tensor<13x21x3xi32>]} : (tensor<3xi32>) -> tensor<13x21x3xi32>
+  // CHECK: emitc.call_opaque "emitc::tosa::mul"(%arg0, %0) {args = [0 : index, 1 : index, 1 : i8]} : (tensor<13x21x3xi32>, tensor<13x21x3xi32>) -> tensor<13x21x3xi32>
   %0 = "tosa.mul"(%arg0, %arg1)  { shift = 1 : i8 } : (tensor<13x21x3xi32>, tensor<3xi32>) -> tensor<13x21x3xi32>
   return %0 : tensor<13x21x3xi32>
 }
 
 // MulOp: Second two dimensional operand needs to be broadcasted + expanded to four dimensions
 func.func @test_mul5(%arg0: tensor<2x13x21x3xi32>, %arg1: tensor<21x3xi32>) -> tensor<2x13x21x3xi32> {
-  // CHECK: emitc.call "emitc::broadcast_in_dim"(%arg1) {args = [0 : index, dense<[2, 3]> : tensor<2xi64>], template_args = [tensor<2x13x21x3xi32>]} : (tensor<21x3xi32>) -> tensor<2x13x21x3xi32>
-  // CHECK: emitc.call "emitc::tosa::mul"(%arg0, %0) {args = [0 : index, 1 : index, 5 : i8]} : (tensor<2x13x21x3xi32>, tensor<2x13x21x3xi32>) -> tensor<2x13x21x3xi32>
+  // CHECK: emitc.call_opaque "emitc::broadcast_in_dim"(%arg1) {args = [0 : index, dense<[2, 3]> : tensor<2xi64>], template_args = [tensor<2x13x21x3xi32>]} : (tensor<21x3xi32>) -> tensor<2x13x21x3xi32>
+  // CHECK: emitc.call_opaque "emitc::tosa::mul"(%arg0, %0) {args = [0 : index, 1 : index, 5 : i8]} : (tensor<2x13x21x3xi32>, tensor<2x13x21x3xi32>) -> tensor<2x13x21x3xi32>
   %0 = "tosa.mul"(%arg0, %arg1)  { shift = 5 : i8 } : (tensor<2x13x21x3xi32>, tensor<21x3xi32>) -> tensor<2x13x21x3xi32>
   return %0 : tensor<2x13x21x3xi32>
 }
 
 func.func @test_maximum(%arg0: tensor<13x21x3xf32>, %arg1: tensor<13x21x1xf32>) -> tensor<13x21x3xf32> {
-  // CHECK: %0 = emitc.call "emitc::broadcast_in_dim"(%arg1) {args = [0 : index, dense<[0, 1, 2]> : tensor<3xi64>], template_args = [tensor<13x21x3xf32>]} : (tensor<13x21x1xf32>) -> tensor<13x21x3xf32>
-  // CHECK: %1 = emitc.call "emitc::tosa::maximum"(%arg0, %0) : (tensor<13x21x3xf32>, tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
+  // CHECK: %0 = emitc.call_opaque "emitc::broadcast_in_dim"(%arg1) {args = [0 : index, dense<[0, 1, 2]> : tensor<3xi64>], template_args = [tensor<13x21x3xf32>]} : (tensor<13x21x1xf32>) -> tensor<13x21x3xf32>
+  // CHECK: %1 = emitc.call_opaque "emitc::tosa::maximum"(%arg0, %0) : (tensor<13x21x3xf32>, tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   %0 = "tosa.maximum"(%arg0, %arg1) : (tensor<13x21x3xf32>, tensor<13x21x1xf32>) -> tensor<13x21x3xf32>
   return %0 : tensor<13x21x3xf32>
 }
 
 func.func @test_minimum(%arg0: tensor<13x21x3xf32>, %arg1: tensor<1x21x3xf32>) -> tensor<13x21x3xf32> {
-  // CHECK: %0 = emitc.call "emitc::broadcast_in_dim"(%arg1) {args = [0 : index, dense<[0, 1, 2]> : tensor<3xi64>], template_args = [tensor<13x21x3xf32>]} : (tensor<1x21x3xf32>) -> tensor<13x21x3xf32>
-  // CHECK: %1 = emitc.call "emitc::tosa::minimum"(%arg0, %0) : (tensor<13x21x3xf32>, tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
+  // CHECK: %0 = emitc.call_opaque "emitc::broadcast_in_dim"(%arg1) {args = [0 : index, dense<[0, 1, 2]> : tensor<3xi64>], template_args = [tensor<13x21x3xf32>]} : (tensor<1x21x3xf32>) -> tensor<13x21x3xf32>
+  // CHECK: %1 = emitc.call_opaque "emitc::tosa::minimum"(%arg0, %0) : (tensor<13x21x3xf32>, tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   %0 = "tosa.minimum"(%arg0, %arg1) : (tensor<13x21x3xf32>, tensor<1x21x3xf32>) -> tensor<13x21x3xf32>
   return %0 : tensor<13x21x3xf32>
 }
 
 func.func @test_sub(%arg0: tensor<13x21x1xf32>, %arg1: tensor<13x21x3xf32>) -> tensor<13x21x3xf32> {
-  // CHECK: emitc.call "emitc::broadcast_in_dim"(%arg0) {args = [0 : index, dense<[0, 1, 2]> : tensor<3xi64>], template_args = [tensor<13x21x3xf32>]} : (tensor<13x21x1xf32>) -> tensor<13x21x3xf32>
-  // CHECK: emitc.call "emitc::tosa::sub"(%0, %arg1) : (tensor<13x21x3xf32>, tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
+  // CHECK: emitc.call_opaque "emitc::broadcast_in_dim"(%arg0) {args = [0 : index, dense<[0, 1, 2]> : tensor<3xi64>], template_args = [tensor<13x21x3xf32>]} : (tensor<13x21x1xf32>) -> tensor<13x21x3xf32>
+  // CHECK: emitc.call_opaque "emitc::tosa::sub"(%0, %arg1) : (tensor<13x21x3xf32>, tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   %0 = "tosa.sub"(%arg0, %arg1) : (tensor<13x21x1xf32>, tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   return %0 : tensor<13x21x3xf32>
 }
 
 func.func @test_pow(%arg0: tensor<13x21x3xf32>, %arg1: tensor<13x21x1xf32>) -> tensor<13x21x3xf32> {
-  // CHECK: %0 = emitc.call "emitc::broadcast_in_dim"(%arg1) {args = [0 : index, dense<[0, 1, 2]> : tensor<3xi64>], template_args = [tensor<13x21x3xf32>]} : (tensor<13x21x1xf32>) -> tensor<13x21x3xf32>
-  // CHECK: %1 = emitc.call "emitc::tosa::pow"(%arg0, %0) : (tensor<13x21x3xf32>, tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
+  // CHECK: %0 = emitc.call_opaque "emitc::broadcast_in_dim"(%arg1) {args = [0 : index, dense<[0, 1, 2]> : tensor<3xi64>], template_args = [tensor<13x21x3xf32>]} : (tensor<13x21x1xf32>) -> tensor<13x21x3xf32>
+  // CHECK: %1 = emitc.call_opaque "emitc::tosa::pow"(%arg0, %0) : (tensor<13x21x3xf32>, tensor<13x21x3xf32>) -> tensor<13x21x3xf32>
   %0 = "tosa.pow"(%arg0, %arg1) : (tensor<13x21x3xf32>, tensor<13x21x1xf32>) -> tensor<13x21x3xf32>
   return %0 : tensor<13x21x3xf32>
 }
 
 func.func @test_table(%arg0: tensor<64xi16>, %arg1: tensor<513xi16>) -> tensor<64xi32> {
-  // CHECK: emitc.call "emitc::tosa::table"(%arg0, %arg1) : (tensor<64xi16>, tensor<513xi16>) -> tensor<64xi32>
+  // CHECK: emitc.call_opaque "emitc::tosa::table"(%arg0, %arg1) : (tensor<64xi16>, tensor<513xi16>) -> tensor<64xi32>
   %0 = "tosa.table"(%arg0, %arg1) : (tensor<64xi16>, tensor<513xi16>) -> tensor<64xi32>
   return %0 : tensor<64xi32>
 }
@@ -228,30 +228,30 @@ func.func @test_table(%arg0: tensor<64xi16>, %arg1: tensor<513xi16>) -> tensor<6
 // Ternary ops
 
 func.func @test_select(%arg0: tensor<12x6x3xi1>, %arg1: tensor<12x6x3xi32>, %arg2: tensor<12x6x3xi32>) -> tensor<12x6x3xi32> {
-  // CHECK: emitc.call "emitc::tosa::select"(%arg0, %arg1, %arg2) : (tensor<12x6x3xi1>, tensor<12x6x3xi32>, tensor<12x6x3xi32>) -> tensor<12x6x3xi32>
+  // CHECK: emitc.call_opaque "emitc::tosa::select"(%arg0, %arg1, %arg2) : (tensor<12x6x3xi1>, tensor<12x6x3xi32>, tensor<12x6x3xi32>) -> tensor<12x6x3xi32>
   %0 = "tosa.select"(%arg0, %arg1, %arg2) : (tensor<12x6x3xi1>, tensor<12x6x3xi32>, tensor<12x6x3xi32>) -> tensor<12x6x3xi32>
   return %0 : tensor<12x6x3xi32>
 }
 
 func.func @test_select_broadcast_condition(%arg0: tensor<12x6x1xi1>, %arg1: tensor<12x6x3xi32>, %arg2: tensor<12x6x3xi32>) -> tensor<12x6x3xi32> {
-  // CHECK: %0 = emitc.call "emitc::broadcast_in_dim"(%arg0) {args = [0 : index, dense<[0, 1, 2]> : tensor<3xi64>], template_args = [tensor<12x6x3xi1>]} : (tensor<12x6x1xi1>) -> tensor<12x6x3xi1>
-  // CHECK: %1 = emitc.call "emitc::tosa::select"(%0, %arg1, %arg2) : (tensor<12x6x3xi1>, tensor<12x6x3xi32>, tensor<12x6x3xi32>) -> tensor<12x6x3xi32>
+  // CHECK: %0 = emitc.call_opaque "emitc::broadcast_in_dim"(%arg0) {args = [0 : index, dense<[0, 1, 2]> : tensor<3xi64>], template_args = [tensor<12x6x3xi1>]} : (tensor<12x6x1xi1>) -> tensor<12x6x3xi1>
+  // CHECK: %1 = emitc.call_opaque "emitc::tosa::select"(%0, %arg1, %arg2) : (tensor<12x6x3xi1>, tensor<12x6x3xi32>, tensor<12x6x3xi32>) -> tensor<12x6x3xi32>
   %0 = "tosa.select"(%arg0, %arg1, %arg2) : (tensor<12x6x1xi1>, tensor<12x6x3xi32>, tensor<12x6x3xi32>) -> tensor<12x6x3xi32>
   return %0 : tensor<12x6x3xi32>
 }
 
 func.func @test_select_broadcast_input(%arg0: tensor<12x6x3xi1>, %arg1: tensor<12x1x3xi32>, %arg2: tensor<12x6x3xi32>) -> tensor<12x6x3xi32> {
-  // CHECK: %0 = emitc.call "emitc::broadcast_in_dim"(%arg1) {args = [0 : index, dense<[0, 1, 2]> : tensor<3xi64>], template_args = [tensor<12x6x3xi32>]} : (tensor<12x1x3xi32>) -> tensor<12x6x3xi32>
-  // CHECK: %1 = emitc.call "emitc::tosa::select"(%arg0, %0, %arg2) : (tensor<12x6x3xi1>, tensor<12x6x3xi32>, tensor<12x6x3xi32>) -> tensor<12x6x3xi32>
+  // CHECK: %0 = emitc.call_opaque "emitc::broadcast_in_dim"(%arg1) {args = [0 : index, dense<[0, 1, 2]> : tensor<3xi64>], template_args = [tensor<12x6x3xi32>]} : (tensor<12x1x3xi32>) -> tensor<12x6x3xi32>
+  // CHECK: %1 = emitc.call_opaque "emitc::tosa::select"(%arg0, %0, %arg2) : (tensor<12x6x3xi1>, tensor<12x6x3xi32>, tensor<12x6x3xi32>) -> tensor<12x6x3xi32>
   %0 = "tosa.select"(%arg0, %arg1, %arg2) : (tensor<12x6x3xi1>, tensor<12x1x3xi32>, tensor<12x6x3xi32>) -> tensor<12x6x3xi32>
   return %0 : tensor<12x6x3xi32>
 }
 
 func.func @test_select_broadcast_all_elements(%arg0: tensor<12x6x1xi1>, %arg1: tensor<12x1x3xi32>, %arg2: tensor<12x1x3xi32>) -> tensor<12x6x3xi32> {
-  // CHECK: %0 = emitc.call "emitc::broadcast_in_dim"(%arg0) {args = [0 : index, dense<[0, 1, 2]> : tensor<3xi64>], template_args = [tensor<12x6x3xi1>]} : (tensor<12x6x1xi1>) -> tensor<12x6x3xi1>
-  // CHECK: %1 = emitc.call "emitc::broadcast_in_dim"(%arg1) {args = [0 : index, dense<[0, 1, 2]> : tensor<3xi64>], template_args = [tensor<12x6x3xi32>]} : (tensor<12x1x3xi32>) -> tensor<12x6x3xi32>
-  // CHECK: %2 = emitc.call "emitc::broadcast_in_dim"(%arg2) {args = [0 : index, dense<[0, 1, 2]> : tensor<3xi64>], template_args = [tensor<12x6x3xi32>]} : (tensor<12x1x3xi32>) -> tensor<12x6x3xi32>
-  // CHECK: %3 = emitc.call "emitc::tosa::select"(%0, %1, %2) : (tensor<12x6x3xi1>, tensor<12x6x3xi32>, tensor<12x6x3xi32>) -> tensor<12x6x3xi32>
+  // CHECK: %0 = emitc.call_opaque "emitc::broadcast_in_dim"(%arg0) {args = [0 : index, dense<[0, 1, 2]> : tensor<3xi64>], template_args = [tensor<12x6x3xi1>]} : (tensor<12x6x1xi1>) -> tensor<12x6x3xi1>
+  // CHECK: %1 = emitc.call_opaque "emitc::broadcast_in_dim"(%arg1) {args = [0 : index, dense<[0, 1, 2]> : tensor<3xi64>], template_args = [tensor<12x6x3xi32>]} : (tensor<12x1x3xi32>) -> tensor<12x6x3xi32>
+  // CHECK: %2 = emitc.call_opaque "emitc::broadcast_in_dim"(%arg2) {args = [0 : index, dense<[0, 1, 2]> : tensor<3xi64>], template_args = [tensor<12x6x3xi32>]} : (tensor<12x1x3xi32>) -> tensor<12x6x3xi32>
+  // CHECK: %3 = emitc.call_opaque "emitc::tosa::select"(%0, %1, %2) : (tensor<12x6x3xi1>, tensor<12x6x3xi32>, tensor<12x6x3xi32>) -> tensor<12x6x3xi32>
   %0 = "tosa.select"(%arg0, %arg1, %arg2) : (tensor<12x6x1xi1>, tensor<12x1x3xi32>, tensor<12x1x3xi32>) -> tensor<12x6x3xi32>
   return %0 : tensor<12x6x3xi32>
 }
@@ -259,59 +259,59 @@ func.func @test_select_broadcast_all_elements(%arg0: tensor<12x6x1xi1>, %arg1: t
 // Other ops
 
 func.func @test_concat(%arg0: tensor<13x21x3xf32>, %arg1: tensor<13x21x3xf32>, %arg2: tensor<13x21x3xf32>, %arg3: tensor<13x21x3xf32>) -> tensor<52x21x3xf32> {
-  // CHECK: %0 = emitc.call "emitc::tosa::concat"(%arg0, %arg1, %arg2, %arg3) {template_args = [0 : i32, tensor<52x21x3xf32>]} : (tensor<13x21x3xf32>, tensor<13x21x3xf32>, tensor<13x21x3xf32>, tensor<13x21x3xf32>) -> tensor<52x21x3xf32>
+  // CHECK: %0 = emitc.call_opaque "emitc::tosa::concat"(%arg0, %arg1, %arg2, %arg3) {template_args = [0 : i32, tensor<52x21x3xf32>]} : (tensor<13x21x3xf32>, tensor<13x21x3xf32>, tensor<13x21x3xf32>, tensor<13x21x3xf32>) -> tensor<52x21x3xf32>
   %0 = "tosa.concat"(%arg0, %arg1, %arg2, %arg3) {axis = 0 : i32} : (tensor<13x21x3xf32>, tensor<13x21x3xf32>, tensor<13x21x3xf32>, tensor<13x21x3xf32>) -> tensor<52x21x3xf32>
   return %0 : tensor<52x21x3xf32>
 }
 
 func.func @test_conv2d(%arg0: tensor<1x4x4x4xf32>, %arg1: tensor<8x1x1x4xf32>, %arg2: tensor<8xf32>) -> tensor<1x4x4x8xf32> {
-    // CHECK: %0 = emitc.call "emitc::tosa::conv2d"(%arg0, %arg1) {args = [0 : index, 1 : index, dense<0> : tensor<4xi64>, dense<1> : tensor<2xi64>, dense<1> : tensor<2xi64>], template_args = [tensor<1x4x4x8xf32>]} : (tensor<1x4x4x4xf32>, tensor<8x1x1x4xf32>) -> tensor<1x4x4x8xf32>
-    // CHECK: %1 = emitc.call "emitc::broadcast_in_dim"(%arg2) {args = [0 : index, dense<3> : tensor<1xi64>], template_args = [tensor<1x4x4x8xf32>]} : (tensor<8xf32>) -> tensor<1x4x4x8xf32>
-    // CHECK: %2 = emitc.call "emitc::tosa::add"(%0, %1) : (tensor<1x4x4x8xf32>, tensor<1x4x4x8xf32>) -> tensor<1x4x4x8xf32>
+    // CHECK: %0 = emitc.call_opaque "emitc::tosa::conv2d"(%arg0, %arg1) {args = [0 : index, 1 : index, dense<0> : tensor<4xi64>, dense<1> : tensor<2xi64>, dense<1> : tensor<2xi64>], template_args = [tensor<1x4x4x8xf32>]} : (tensor<1x4x4x4xf32>, tensor<8x1x1x4xf32>) -> tensor<1x4x4x8xf32>
+    // CHECK: %1 = emitc.call_opaque "emitc::broadcast_in_dim"(%arg2) {args = [0 : index, dense<3> : tensor<1xi64>], template_args = [tensor<1x4x4x8xf32>]} : (tensor<8xf32>) -> tensor<1x4x4x8xf32>
+    // CHECK: %2 = emitc.call_opaque "emitc::tosa::add"(%0, %1) : (tensor<1x4x4x8xf32>, tensor<1x4x4x8xf32>) -> tensor<1x4x4x8xf32>
     %0 = "tosa.conv2d"(%arg0, %arg1, %arg2) {dilation = array<i64: 1, 1>, pad = array<i64: 0, 0, 0, 0>, stride = array<i64: 1, 1>} : (tensor<1x4x4x4xf32>, tensor<8x1x1x4xf32>, tensor<8xf32>) -> tensor<1x4x4x8xf32>
     return %0 : tensor<1x4x4x8xf32>
 }
 
 func.func @test_depthwise_conv2d(%arg0: tensor<1x4x5x2xf32>, %arg1: tensor<2x2x2x2xf32>, %arg2: tensor<4xf32>) -> tensor<1x3x4x4xf32> {
-    // CHECK: %0 = emitc.call "emitc::tosa::depthwise_conv2d"(%arg0, %arg1) {args = [0 : index, 1 : index, dense<0> : tensor<4xi64>, dense<1> : tensor<2xi64>, dense<1> : tensor<2xi64>], template_args = [tensor<1x3x4x4xf32>]} : (tensor<1x4x5x2xf32>, tensor<2x2x2x2xf32>) -> tensor<1x3x4x4xf32>
-    // CHECK: %1 = emitc.call "emitc::broadcast_in_dim"(%arg2) {args = [0 : index, dense<3> : tensor<1xi64>], template_args = [tensor<1x3x4x4xf32>]} : (tensor<4xf32>) -> tensor<1x3x4x4xf32>
-    // CHECK: %2 = emitc.call "emitc::tosa::add"(%0, %1) : (tensor<1x3x4x4xf32>, tensor<1x3x4x4xf32>) -> tensor<1x3x4x4xf32>
+    // CHECK: %0 = emitc.call_opaque "emitc::tosa::depthwise_conv2d"(%arg0, %arg1) {args = [0 : index, 1 : index, dense<0> : tensor<4xi64>, dense<1> : tensor<2xi64>, dense<1> : tensor<2xi64>], template_args = [tensor<1x3x4x4xf32>]} : (tensor<1x4x5x2xf32>, tensor<2x2x2x2xf32>) -> tensor<1x3x4x4xf32>
+    // CHECK: %1 = emitc.call_opaque "emitc::broadcast_in_dim"(%arg2) {args = [0 : index, dense<3> : tensor<1xi64>], template_args = [tensor<1x3x4x4xf32>]} : (tensor<4xf32>) -> tensor<1x3x4x4xf32>
+    // CHECK: %2 = emitc.call_opaque "emitc::tosa::add"(%0, %1) : (tensor<1x3x4x4xf32>, tensor<1x3x4x4xf32>) -> tensor<1x3x4x4xf32>
     %0 = "tosa.depthwise_conv2d"(%arg0, %arg1, %arg2) {dilation = array<i64: 1, 1>, pad = array<i64: 0, 0, 0, 0>, stride = array<i64: 1, 1>} : (tensor<1x4x5x2xf32>, tensor<2x2x2x2xf32>, tensor<4xf32>) -> tensor<1x3x4x4xf32>
     return %0 : tensor<1x3x4x4xf32>
 }
 
 func.func @test_max_pool2d(%arg0: tensor<1x32x32x8xf32>) -> tensor<1x32x32x8xf32> {
-  // CHECK: emitc.call "emitc::tosa::max_pool2d"(%arg0) {args = [0 : index, dense<0> : tensor<4xi64>, dense<1> : tensor<2xi64>, dense<1> : tensor<2xi64>], template_args = [tensor<1x32x32x8xf32>]} : (tensor<1x32x32x8xf32>) -> tensor<1x32x32x8xf32>
+  // CHECK: emitc.call_opaque "emitc::tosa::max_pool2d"(%arg0) {args = [0 : index, dense<0> : tensor<4xi64>, dense<1> : tensor<2xi64>, dense<1> : tensor<2xi64>], template_args = [tensor<1x32x32x8xf32>]} : (tensor<1x32x32x8xf32>) -> tensor<1x32x32x8xf32>
   %0 = "tosa.max_pool2d"(%arg0) {kernel = array<i64: 1, 1>, pad = array<i64: 0, 0, 0, 0>, stride = array<i64: 1, 1>} : (tensor<1x32x32x8xf32>) -> tensor<1x32x32x8xf32>
   return %0 : tensor<1x32x32x8xf32>
 }
 
 func.func @test_avg_pool2d(%arg0: tensor<1x32x32x8xf32>) -> tensor<1x32x32x8xf32> {
-  // CHECK: emitc.call "emitc::tosa::avg_pool2d"(%arg0) {args = [0 : index, dense<[0, 1, 0, 1]> : tensor<4xi64>, dense<1> : tensor<2xi64>, dense<2> : tensor<2xi64>], template_args = [tensor<1x32x32x8xf32>]} : (tensor<1x32x32x8xf32>) -> tensor<1x32x32x8xf32>
+  // CHECK: emitc.call_opaque "emitc::tosa::avg_pool2d"(%arg0) {args = [0 : index, dense<[0, 1, 0, 1]> : tensor<4xi64>, dense<1> : tensor<2xi64>, dense<2> : tensor<2xi64>], template_args = [tensor<1x32x32x8xf32>]} : (tensor<1x32x32x8xf32>) -> tensor<1x32x32x8xf32>
   %0 = "tosa.avg_pool2d"(%arg0) {acc_type = f32, kernel = array<i64: 2, 2>, pad = array<i64: 0, 1, 0, 1>, stride = array<i64: 1, 1>} : (tensor<1x32x32x8xf32>) -> tensor<1x32x32x8xf32>
   return %0 : tensor<1x32x32x8xf32>
 }
 
 func.func @test_fully_connected(%arg0: tensor<14x19xf32>, %arg1: tensor<19x28xf32>, %arg2: tensor<28xf32>) -> tensor<14x28xf32> {
-  // CHECK: emitc.call "emitc::tosa::fully_connected"(%arg0, %arg1, %arg2) {template_args = [tensor<14x28xf32>]} : (tensor<14x19xf32>, tensor<19x28xf32>, tensor<28xf32>) -> tensor<14x28xf32>
+  // CHECK: emitc.call_opaque "emitc::tosa::fully_connected"(%arg0, %arg1, %arg2) {template_args = [tensor<14x28xf32>]} : (tensor<14x19xf32>, tensor<19x28xf32>, tensor<28xf32>) -> tensor<14x28xf32>
   %0 = "tosa.fully_connected"(%arg0, %arg1, %arg2) : (tensor<14x19xf32>, tensor<19x28xf32>, tensor<28xf32>) -> tensor<14x28xf32>
   return %0 : tensor<14x28xf32>
 }
 
 func.func @test_gather(%arg0: tensor<3x4x5xf32>, %arg1: tensor<3x6xi32>) -> tensor<3x6x5xf32> {
-  // CHECK: emitc.call "emitc::tosa::gather"(%arg0, %arg1) {template_args = [tensor<3x6x5xf32>]} : (tensor<3x4x5xf32>, tensor<3x6xi32>) -> tensor<3x6x5xf32>
+  // CHECK: emitc.call_opaque "emitc::tosa::gather"(%arg0, %arg1) {template_args = [tensor<3x6x5xf32>]} : (tensor<3x4x5xf32>, tensor<3x6xi32>) -> tensor<3x6x5xf32>
   %0 = "tosa.gather"(%arg0, %arg1) : (tensor<3x4x5xf32>, tensor<3x6xi32>) -> tensor<3x6x5xf32>
   return %0 : tensor<3x6x5xf32>
 }
 
 func.func @test_matmul(%arg0: tensor<1x14x19xf32>, %arg1: tensor<1x19x28xf32>) -> tensor<1x14x28xf32> {
-  // CHECK: emitc.call "emitc::tosa::matmul"(%arg0, %arg1) : (tensor<1x14x19xf32>, tensor<1x19x28xf32>) -> tensor<1x14x28xf32>
+  // CHECK: emitc.call_opaque "emitc::tosa::matmul"(%arg0, %arg1) : (tensor<1x14x19xf32>, tensor<1x19x28xf32>) -> tensor<1x14x28xf32>
   %0 = "tosa.matmul"(%arg0, %arg1) : (tensor<1x14x19xf32>, tensor<1x19x28xf32>) -> tensor<1x14x28xf32>
   return %0 : tensor<1x14x28xf32>
 }
 
 func.func @test_argmax(%arg0: tensor<13x21x3xi32>) -> tensor<21x3xi32> {
-  // CHECK: %0 = emitc.call "emitc::tosa::argmax"(%arg0) {args = [0 : index, 0 : i32], template_args = [tensor<21x3xi32>, tensor<13x21x3xi32>]} : (tensor<13x21x3xi32>) -> tensor<21x3xi32>
+  // CHECK: %0 = emitc.call_opaque "emitc::tosa::argmax"(%arg0) {args = [0 : index, 0 : i32], template_args = [tensor<21x3xi32>, tensor<13x21x3xi32>]} : (tensor<13x21x3xi32>) -> tensor<21x3xi32>
   %0 = "tosa.argmax"(%arg0) {axis = 0 : i32} : (tensor<13x21x3xi32>) -> tensor<21x3xi32>
   return %0 : tensor<21x3xi32>
 }
@@ -319,55 +319,55 @@ func.func @test_argmax(%arg0: tensor<13x21x3xi32>) -> tensor<21x3xi32> {
 // Reduce ops
 
 func.func @test_reduce_all(%arg0: tensor<13x21x3xi1>) -> tensor<1x21x3xi1> {
-  // CHECK: %0 = emitc.call "emitc::tosa::reduce_all"(%arg0) {args = [0 : index, 0 : i32], template_args = [tensor<21x3xi1>, tensor<13x21x3xi1>]} : (tensor<13x21x3xi1>) -> tensor<21x3xi1>
-  // CHECK: %1 = emitc.call "emitc::tosa::reshape"(%0) {template_args = [tensor<1x21x3xi1>]} : (tensor<21x3xi1>) -> tensor<1x21x3xi1>
+  // CHECK: %0 = emitc.call_opaque "emitc::tosa::reduce_all"(%arg0) {args = [0 : index, 0 : i32], template_args = [tensor<21x3xi1>, tensor<13x21x3xi1>]} : (tensor<13x21x3xi1>) -> tensor<21x3xi1>
+  // CHECK: %1 = emitc.call_opaque "emitc::tosa::reshape"(%0) {template_args = [tensor<1x21x3xi1>]} : (tensor<21x3xi1>) -> tensor<1x21x3xi1>
   %0 = "tosa.reduce_all"(%arg0) {axis = 0 : i32} : (tensor<13x21x3xi1>) -> tensor<1x21x3xi1>
   return %0 : tensor<1x21x3xi1>
 }
 
 func.func @test_reduce_any(%arg0: tensor<13x21x3xi1>) -> tensor<13x1x3xi1> {
-  // CHECK: %0 = emitc.call "emitc::tosa::reduce_any"(%arg0) {args = [0 : index, 1 : i32], template_args = [tensor<13x3xi1>, tensor<13x21x3xi1>]} : (tensor<13x21x3xi1>) -> tensor<13x3xi1>
-  // %1 = emitc.call "emitc::tosa::reshape"(%0) {template_args = [tensor<13x1x3xi1>]} : (tensor<13x3xi1>) -> tensor<13x1x3xi1>
+  // CHECK: %0 = emitc.call_opaque "emitc::tosa::reduce_any"(%arg0) {args = [0 : index, 1 : i32], template_args = [tensor<13x3xi1>, tensor<13x21x3xi1>]} : (tensor<13x21x3xi1>) -> tensor<13x3xi1>
+  // %1 = emitc.call_opaque "emitc::tosa::reshape"(%0) {template_args = [tensor<13x1x3xi1>]} : (tensor<13x3xi1>) -> tensor<13x1x3xi1>
   %0 = "tosa.reduce_any"(%arg0) {axis = 1 : i32} : (tensor<13x21x3xi1>) -> tensor<13x1x3xi1>
   return %0 : tensor<13x1x3xi1>
 }
 
 func.func @test_reduce_max(%arg0: tensor<13x21x3xf32>) -> tensor<1x21x3xf32> {
-  // CHECK: %0 = emitc.call "emitc::tosa::reduce_max"(%arg0) {args = [0 : index, 0 : i32], template_args = [tensor<21x3xf32>, tensor<13x21x3xf32>]} : (tensor<13x21x3xf32>) -> tensor<21x3xf32>
-  // CHECK: %1 = emitc.call "emitc::tosa::reshape"(%0) {template_args = [tensor<1x21x3xf32>]} : (tensor<21x3xf32>) -> tensor<1x21x3xf32>
+  // CHECK: %0 = emitc.call_opaque "emitc::tosa::reduce_max"(%arg0) {args = [0 : index, 0 : i32], template_args = [tensor<21x3xf32>, tensor<13x21x3xf32>]} : (tensor<13x21x3xf32>) -> tensor<21x3xf32>
+  // CHECK: %1 = emitc.call_opaque "emitc::tosa::reshape"(%0) {template_args = [tensor<1x21x3xf32>]} : (tensor<21x3xf32>) -> tensor<1x21x3xf32>
   %0 = "tosa.reduce_max"(%arg0) {axis = 0 : i32} : (tensor<13x21x3xf32>) -> tensor<1x21x3xf32>
   return %0 : tensor<1x21x3xf32>
 }
 
 func.func @test_reduce_min(%arg0: tensor<13x21x3xf32>) -> tensor<13x1x3xf32> {
-  // CHECK: %0 = emitc.call "emitc::tosa::reduce_min"(%arg0) {args = [0 : index, 1 : i32], template_args = [tensor<13x3xf32>, tensor<13x21x3xf32>]} : (tensor<13x21x3xf32>) -> tensor<13x3xf32>
-  // CHECK: %1 = emitc.call "emitc::tosa::reshape"(%0) {template_args = [tensor<13x1x3xf32>]} : (tensor<13x3xf32>) -> tensor<13x1x3xf32>
+  // CHECK: %0 = emitc.call_opaque "emitc::tosa::reduce_min"(%arg0) {args = [0 : index, 1 : i32], template_args = [tensor<13x3xf32>, tensor<13x21x3xf32>]} : (tensor<13x21x3xf32>) -> tensor<13x3xf32>
+  // CHECK: %1 = emitc.call_opaque "emitc::tosa::reshape"(%0) {template_args = [tensor<13x1x3xf32>]} : (tensor<13x3xf32>) -> tensor<13x1x3xf32>
   %0 = "tosa.reduce_min"(%arg0) {axis = 1 : i32} : (tensor<13x21x3xf32>) -> tensor<13x1x3xf32>
   return %0 : tensor<13x1x3xf32>
 }
 
 func.func @test_reduce_prod(%arg0: tensor<13x21x3xf32>) -> tensor<1x21x3xf32> {
-  // CHECK: %0 = emitc.call "emitc::tosa::reduce_prod"(%arg0) {args = [0 : index, 0 : i32], template_args = [tensor<21x3xf32>, tensor<13x21x3xf32>]} : (tensor<13x21x3xf32>) -> tensor<21x3xf32>
-  // CHECK: %1 = emitc.call "emitc::tosa::reshape"(%0) {template_args = [tensor<1x21x3xf32>]} : (tensor<21x3xf32>) -> tensor<1x21x3xf32>
+  // CHECK: %0 = emitc.call_opaque "emitc::tosa::reduce_prod"(%arg0) {args = [0 : index, 0 : i32], template_args = [tensor<21x3xf32>, tensor<13x21x3xf32>]} : (tensor<13x21x3xf32>) -> tensor<21x3xf32>
+  // CHECK: %1 = emitc.call_opaque "emitc::tosa::reshape"(%0) {template_args = [tensor<1x21x3xf32>]} : (tensor<21x3xf32>) -> tensor<1x21x3xf32>
   %0 = "tosa.reduce_prod"(%arg0) {axis = 0 : i32} : (tensor<13x21x3xf32>) -> tensor<1x21x3xf32>
   return %0 : tensor<1x21x3xf32>
 }
 
 func.func @test_reduce_sum(%arg0: tensor<13x21x3xf32>) -> tensor<13x1x3xf32> {
-  // CHECK: %0 = emitc.call "emitc::tosa::reduce_sum"(%arg0) {args = [0 : index, 1 : i32], template_args = [tensor<13x3xf32>, tensor<13x21x3xf32>]} : (tensor<13x21x3xf32>) -> tensor<13x3xf32>
-  // CHECK: %1 = emitc.call "emitc::tosa::reshape"(%0) {template_args = [tensor<13x1x3xf32>]} : (tensor<13x3xf32>) -> tensor<13x1x3xf32>
+  // CHECK: %0 = emitc.call_opaque "emitc::tosa::reduce_sum"(%arg0) {args = [0 : index, 1 : i32], template_args = [tensor<13x3xf32>, tensor<13x21x3xf32>]} : (tensor<13x21x3xf32>) -> tensor<13x3xf32>
+  // CHECK: %1 = emitc.call_opaque "emitc::tosa::reshape"(%0) {template_args = [tensor<13x1x3xf32>]} : (tensor<13x3xf32>) -> tensor<13x1x3xf32>
   %0 = "tosa.reduce_sum"(%arg0) {axis = 1 : i32} : (tensor<13x21x3xf32>) -> tensor<13x1x3xf32>
   return %0 : tensor<13x1x3xf32>
 }
 
 func.func @test_slice(%arg0: tensor<13x21x3xf32>) -> tensor<4x11x1xf32> {
-  // CHECK: %0 = emitc.call "emitc::tosa::slice"(%arg0) {args = [0 : index, array<i64: 6, 8, 0>, array<i64: 4, 11, 1>], template_args = [tensor<4x11x1xf32>]} : (tensor<13x21x3xf32>) -> tensor<4x11x1xf32>
+  // CHECK: %0 = emitc.call_opaque "emitc::tosa::slice"(%arg0) {args = [0 : index, array<i64: 6, 8, 0>, array<i64: 4, 11, 1>], template_args = [tensor<4x11x1xf32>]} : (tensor<13x21x3xf32>) -> tensor<4x11x1xf32>
   %0 = "tosa.slice"(%arg0) {start = array<i64: 6, 8, 0>, size = array<i64: 4, 11, 1>} : (tensor<13x21x3xf32>) -> tensor<4x11x1xf32>
   return %0 : tensor<4x11x1xf32>
 }
 
 func.func @test_pad(%arg0: tensor<2x3xf32>, %arg1: tensor<2x2xi32>) -> tensor<3x6xf32> {
-  // CHECK: %0 = emitc.call "emitc::tosa::pad"(%arg0, %arg1) {template_args = [tensor<3x6xf32>]} : (tensor<2x3xf32>, tensor<2x2xi32>) -> tensor<3x6xf32>
+  // CHECK: %0 = emitc.call_opaque "emitc::tosa::pad"(%arg0, %arg1) {template_args = [tensor<3x6xf32>]} : (tensor<2x3xf32>, tensor<2x2xi32>) -> tensor<3x6xf32>
   %0 = "tosa.pad"(%arg0, %arg1) : (tensor<2x3xf32>, tensor<2x2xi32>) -> tensor<3x6xf32>
   return %0 : tensor<3x6xf32>
 }
@@ -375,19 +375,19 @@ func.func @test_pad(%arg0: tensor<2x3xf32>, %arg1: tensor<2x2xi32>) -> tensor<3x
 func.func @test_pad_explicit_value(%arg0: tensor<2x3xf32>, %arg1: tensor<2x2xi32>) -> tensor<3x6xf32> {
   // CHECK: %0 = "emitc.constant"() <{value = dense<3.140000e+00> : tensor<f32>}> : () -> tensor<f32>
   %0 = "tosa.const"() {value = dense<3.14> : tensor<f32>} : () -> tensor<f32>
-  // CHECK-NEXT: %1 = emitc.call "emitc::tosa::pad"(%arg0, %arg1, %0) {template_args = [tensor<3x6xf32>]} : (tensor<2x3xf32>, tensor<2x2xi32>, tensor<f32>) -> tensor<3x6xf32>
+  // CHECK-NEXT: %1 = emitc.call_opaque "emitc::tosa::pad"(%arg0, %arg1, %0) {template_args = [tensor<3x6xf32>]} : (tensor<2x3xf32>, tensor<2x2xi32>, tensor<f32>) -> tensor<3x6xf32>
   %1 = "tosa.pad"(%arg0, %arg1, %0) : (tensor<2x3xf32>, tensor<2x2xi32>, tensor<f32>) -> tensor<3x6xf32>
   return %1 : tensor<3x6xf32>
 }
 
 func.func @test_reshape(%arg0: tensor<13x21x3xf32>) -> tensor<1x819xf32> {
-  // CHECK: %0 = emitc.call "emitc::tosa::reshape"(%arg0) {template_args = [tensor<1x819xf32>]} : (tensor<13x21x3xf32>) -> tensor<1x819xf32>
+  // CHECK: %0 = emitc.call_opaque "emitc::tosa::reshape"(%arg0) {template_args = [tensor<1x819xf32>]} : (tensor<13x21x3xf32>) -> tensor<1x819xf32>
   %0 = "tosa.reshape"(%arg0) {new_shape = array<i64: 1, 819>} : (tensor<13x21x3xf32>) -> tensor<1x819xf32>
   return %0 : tensor<1x819xf32>
 }
 
 func.func @test_tile(%arg0: tensor<1x3x1x4xf32>) -> tensor<2x3x3x8xf32> {
-  // CHECK: %0 = emitc.call "emitc::tosa::tile"(%arg0) {args = [0 : index, array<i64: 2, 1, 3, 2>], template_args = [tensor<2x3x3x8xf32>]} : (tensor<1x3x1x4xf32>) -> tensor<2x3x3x8xf32>
+  // CHECK: %0 = emitc.call_opaque "emitc::tosa::tile"(%arg0) {args = [0 : index, array<i64: 2, 1, 3, 2>], template_args = [tensor<2x3x3x8xf32>]} : (tensor<1x3x1x4xf32>) -> tensor<2x3x3x8xf32>
   %0 = "tosa.tile"(%arg0) {multiples = array<i64: 2, 1, 3, 2>} : (tensor<1x3x1x4xf32>) -> tensor<2x3x3x8xf32>
   return %0 : tensor<2x3x3x8xf32>
 }
@@ -395,7 +395,7 @@ func.func @test_tile(%arg0: tensor<1x3x1x4xf32>) -> tensor<2x3x3x8xf32> {
 func.func @test_transpose(%arg0: tensor<13x21x3xf32>) -> tensor<3x13x21xf32> {
   // CHECK: %0 = "emitc.constant"() <{value = dense<[2, 0, 1]> : tensor<3xi32>}> : () -> tensor<3xi32>
   %0 = "tosa.const"() {value = dense<[2, 0, 1]> : tensor<3xi32>} : () -> tensor<3xi32>
-  // CHECK-NEXT: %1 = emitc.call "emitc::tosa::transpose"(%arg0, %0) {template_args = [tensor<3x13x21xf32>]} : (tensor<13x21x3xf32>, tensor<3xi32>) -> tensor<3x13x21xf32>
+  // CHECK-NEXT: %1 = emitc.call_opaque "emitc::tosa::transpose"(%arg0, %0) {template_args = [tensor<3x13x21xf32>]} : (tensor<13x21x3xf32>, tensor<3xi32>) -> tensor<3x13x21xf32>
   %1 = "tosa.transpose"(%arg0, %0) : (tensor<13x21x3xf32>, tensor<3xi32>) -> tensor<3x13x21xf32>
   return %1 : tensor<3x13x21xf32>
 }
diff --git a/third_party/stablehlo b/third_party/stablehlo
index 5e41e674..42387b0f 160000
--- a/third_party/stablehlo
+++ b/third_party/stablehlo
@@ -1 +1 @@
-Subproject commit 5e41e674af78da676652459c2dcf6a0d76e59ddb
+Subproject commit 42387b0fecd4a2782b0b12f806dfc617e05c33f7