Review changes

Author: linuxrocks123

clang/include/clang/Basic/BuiltinsAMDGPU.def

@@ -63,8 +63,8 @@ BUILTIN(__builtin_amdgcn_grid_size_z, "Ui", "nc")
 BUILTIN(__builtin_amdgcn_mbcnt_hi, "UiUiUi", "nc")
 BUILTIN(__builtin_amdgcn_mbcnt_lo, "UiUiUi", "nc")
 
-BUILTIN(__builtin_amdgcn_bcnt032_lo, "UiUi", "nc")
-BUILTIN(__builtin_amdgcn_bcnt064_lo, "UiWUi", "nc")
+BUILTIN(__builtin_amdgcn_bcnt32_lo, "UiUi", "nc")
+BUILTIN(__builtin_amdgcn_bcnt64_lo, "UiWUi", "nc")
 
 TARGET_BUILTIN(__builtin_amdgcn_s_memtime, "WUi", "n", "s-memtime-inst")
 

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

@@ -2359,12 +2359,12 @@ def int_amdgcn_mbcnt_hi :
   DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
             [IntrNoMem]>;
 
-def int_amdgcn_bcnt032_lo :
-  ClangBuiltin<"__builtin_amdgcn_bcnt032_lo">,
+def int_amdgcn_bcnt32_lo :
+  ClangBuiltin<"__builtin_amdgcn_bcnt32_lo">,
   DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem]>;
 
-def int_amdgcn_bcnt064_lo :
-  ClangBuiltin<"__builtin_amdgcn_bcnt064_lo">,
+def int_amdgcn_bcnt64_lo :
+  ClangBuiltin<"__builtin_amdgcn_bcnt64_lo">,
   DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem]>;
 
 // llvm.amdgcn.ds.swizzle src offset

llvm/lib/Target/AMDGPU/AMDGPUCodeGenPrepare.cpp

@@ -95,7 +95,7 @@ static cl::opt<bool> DisableFDivExpand(
   cl::ReallyHidden,
   cl::init(false));
 
-// Disable processing of fdiv so we can better test the backend implementations.
+// Disable bitsin(typeof(x)) - popcnt(x) to s_bcnt0(x) transformation.
 static cl::opt<bool>
     DisableBcnt0("amdgpu-codegenprepare-disable-bcnt0",
                  cl::desc("Prevent transforming bitsin(typeof(x)) - "
@@ -2005,8 +2005,8 @@ bool AMDGPUCodeGenPrepareImpl::visitCtpop(IntrinsicInst &I) {
 
   IRBuilder<> Builder(MustBeSub);
   Instruction *TransformedIns =
-      Builder.CreateIntrinsic(BitWidth > 32 ? Intrinsic::amdgcn_bcnt064_lo
-                                            : Intrinsic::amdgcn_bcnt032_lo,
+      Builder.CreateIntrinsic(BitWidth > 32 ? Intrinsic::amdgcn_bcnt64_lo
+                                            : Intrinsic::amdgcn_bcnt32_lo,
                               {}, {I.getArgOperand(0)});
 
   if ((DestinationWidth = MustBeSub->getType()->getIntegerBitWidth()) !=

llvm/lib/Target/AMDGPU/SOPInstructions.td

@@ -265,10 +265,10 @@ def S_BREV_B64 : SOP1_64 <"s_brev_b64",
 
 let Defs = [SCC] in {
 def S_BCNT0_I32_B32 : SOP1_32 <"s_bcnt0_i32_b32",
-  [(set i32:$sdst, (UniformUnaryFrag<int_amdgcn_bcnt032_lo> i32:$src0))]
+  [(set i32:$sdst, (UniformUnaryFrag<int_amdgcn_bcnt32_lo> i32:$src0))]
 >;
 def S_BCNT0_I32_B64 : SOP1_32_64 <"s_bcnt0_i32_b64",
-  [(set i32:$sdst, (UniformUnaryFrag<int_amdgcn_bcnt064_lo> i64:$src0))]
+  [(set i32:$sdst, (UniformUnaryFrag<int_amdgcn_bcnt64_lo> i64:$src0))]
 >;
 def S_BCNT1_I32_B32 : SOP1_32 <"s_bcnt1_i32_b32",
   [(set i32:$sdst, (UniformUnaryFrag<ctpop> i32:$src0))]