diff --git a/llvm/lib/Target/RISCV/RISCVISelLowering.cpp b/llvm/lib/Target/RISCV/RISCVISelLowering.cpp
index 85ddcec64797..0ae6284a4371 100644
--- a/llvm/lib/Target/RISCV/RISCVISelLowering.cpp
+++ b/llvm/lib/Target/RISCV/RISCVISelLowering.cpp
@@ -7460,46 +7460,6 @@ SDValue RISCVTargetLowering::lowerKernArgParameterPtr(SelectionDAG &DAG,
   return DAG.getObjectPtrOffset(SL, BasePtr, TypeSize::Fixed(Offset));
 }
 
-SDValue RISCVTargetLowering::getFPExtOrFPRound(SelectionDAG &DAG,
-                                            SDValue Op,
-                                            const SDLoc &DL,
-                                            EVT VT) const {
-  return Op.getValueType().bitsLE(VT) ?
-      DAG.getNode(ISD::FP_EXTEND, DL, VT, Op) :
-    DAG.getNode(ISD::FP_ROUND, DL, VT, Op,
-                DAG.getTargetConstant(0, DL, MVT::i32));
-}
-
-SDValue RISCVTargetLowering::convertArgType(SelectionDAG &DAG, EVT VT, EVT MemVT,
-                                         const SDLoc &SL, SDValue Val,
-                                         bool Signed,
-                                         const ISD::InputArg *Arg) const {
-  // First, if it is a widened vector, narrow it.
-  if (VT.isVector() &&
-      VT.getVectorNumElements() != MemVT.getVectorNumElements()) {
-    EVT NarrowedVT =
-        EVT::getVectorVT(*DAG.getContext(), MemVT.getVectorElementType(),
-                         VT.getVectorNumElements());
-    Val = DAG.getNode(ISD::EXTRACT_SUBVECTOR, SL, NarrowedVT, Val,
-                      DAG.getConstant(0, SL, MVT::i32));
-  }
-
-  // Then convert the vector elements or scalar value.
-  if (Arg && (Arg->Flags.isSExt() || Arg->Flags.isZExt()) &&
-      VT.bitsLT(MemVT)) {
-    unsigned Opc = Arg->Flags.isZExt() ? ISD::AssertZext : ISD::AssertSext;
-    Val = DAG.getNode(Opc, SL, MemVT, Val, DAG.getValueType(VT));
-  }
-
-  if (MemVT.isFloatingPoint())
-    Val = getFPExtOrFPRound(DAG, Val, SL, VT);
-  else if (Signed)
-    Val = DAG.getSExtOrTrunc(Val, SL, VT);
-  else
-    Val = DAG.getZExtOrTrunc(Val, SL, VT);
-
-  return Val;
-}
 
 SDValue RISCVTargetLowering::lowerKernargMemParameter(
     SelectionDAG &DAG, EVT VT, EVT MemVT, const SDLoc &SL, SDValue Chain,
@@ -7530,7 +7490,7 @@ SDValue RISCVTargetLowering::lowerKernargMemParameter(
     SDValue ArgVal = DAG.getNode(ISD::TRUNCATE, SL, IntVT, Extract);
     ArgVal = DAG.getNode(ISD::BITCAST, SL, MemVT, ArgVal);
     // TODO: Support vector and half type.
-    ArgVal = convertArgType(DAG, VT, MemVT, SL, ArgVal, Signed, Arg);
+    //ArgVal = convertArgType(DAG, VT, MemVT, SL, ArgVal, Signed, Arg);
 
     return DAG.getMergeValues({ ArgVal, Load.getValue(1) }, SL);
   }
@@ -7540,9 +7500,8 @@ SDValue RISCVTargetLowering::lowerKernargMemParameter(
                              MachineMemOperand::MODereferenceable |
                                  MachineMemOperand::MOInvariant);
 
-  SDValue Val = convertArgType(DAG, VT, MemVT, SL, Load, Signed, Arg);
-  // return DAG.getMergeValues({ Load, Load.getValue(1) }, SL);
-  return DAG.getMergeValues({ Val, Load.getValue(1) }, SL);
+  // SDValue Val = convertArgType(DAG, VT, MemVT, SL, Load, Signed, Arg);
+  return DAG.getMergeValues({ Load, Load.getValue(1) }, SL);
 }
 
 // Returns the opcode of the target-specific SDNode that implements the 32-bit
diff --git a/llvm/lib/Target/RISCV/RISCVISelLowering.h b/llvm/lib/Target/RISCV/RISCVISelLowering.h
index 0965b7020ed1..71ffa4a2b0f0 100644
--- a/llvm/lib/Target/RISCV/RISCVISelLowering.h
+++ b/llvm/lib/Target/RISCV/RISCVISelLowering.h
@@ -722,13 +722,6 @@ private:
   SDValue expandUnalignedRVVStore(SDValue Op, SelectionDAG &DAG) const;
  SDValue lowerKernArgParameterPtr(SelectionDAG &DAG, const SDLoc &SL,
                                    SDValue Chain, uint64_t Offset) const;
-  SDValue getFPExtOrFPRound(SelectionDAG &DAG, SDValue Op,
-                                            const SDLoc &DL,
-                                            EVT VT) const;
-  SDValue convertArgType(SelectionDAG &DAG, EVT VT, EVT MemVT,
-                                         const SDLoc &SL, SDValue Val,
-                                         bool Signed,
-                                         const ISD::InputArg *Arg) const;
   SDValue lowerKernargMemParameter(SelectionDAG &DAG, EVT VT, EVT MemVT,
                                    const SDLoc &SL, SDValue Chain,
                                    uint64_t Offset, Align Alignment,
diff --git a/llvm/test/CodeGen/RISCV/VentusGPGPU/vecargtest.cl b/llvm/test/CodeGen/RISCV/VentusGPGPU/vecargtest.cl
deleted file mode 100644
index f40c5702f7bf..000000000000
--- a/llvm/test/CodeGen/RISCV/VentusGPGPU/vecargtest.cl
+++ /dev/null
@@ -1,19 +0,0 @@
-// RUN: clang -target riscv32 -mcpu=ventus-gpgpu < %s \
-// RUN:   | FileCheck -check-prefix=VENTUS %s
-
-__kernel void test_kernel(
-char2 c, uchar2 uc, short2 s, ushort2 us, int2 i, uint2 ui, float2 f,
-__global float2 *result)
-{
-    // there is no FileCheck now
-    result[0] = convert_float2(c);
-    result[1] = convert_float2(uc);
-    result[2] = convert_float2(s);
-    result[3] = convert_float2(us);
-    result[4] = convert_float2(i);
-    result[5] = convert_float2(ui);
-    result[6] = f;
-}
-
-
-