[PATCH] D15475: AMDGPU/SI: Add llvm.amdgcn.mbcnt.* intrinsics

[Matt Arsenault via llvm-commits]

arsenm added inline comments.

================
Comment at: include/llvm/IR/IntrinsicsAMDGPU.td:150-151
@@ +149,4 @@
+
+def int_amdgcn_v_mbcnt_lo_u32_b32 :
+  GCCBuiltin<"__builtin_amdgcn_v_mbcnt_lo_u32_b32">,
+  Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
----------------
I don't think the type should be included in the name. For instructions that have multiple types, it is redundant with the llvm name mangling. For ones that don't work on multiple types, it isn't necessary information. I would have this be int_amdgcn_mbcnt_lo/int_amdgcn_mbcnt_hi

================
Comment at: test/CodeGen/AMDGPU/llvm.SI.tid.ll:16-33
@@ -15,2 +15,20 @@
 
+;GCN-LABEL: {{^}}mbcnt_intrinsics:
+;GCN: v_mbcnt_lo_u32_b32_e64 [[LO:v[0-9]+]], -1, 0
+;SI: v_mbcnt_hi_u32_b32_e32 {{v[0-9]+}}, -1, [[LO]]
+;VI: v_mbcnt_hi_u32_b32_e64 {{v[0-9]+}}, -1, [[LO]]
+
+define void @mbcnt_intrinsics(<16 x i8> addrspace(2)* inreg, <16 x i8> addrspace(2)* inreg, <32 x i8> addrspace(2)* inreg, i32 inreg) "ShaderType"="0" {
+main_body:
+  %lo = call i32 @llvm.amdgcn.v.mbcnt.lo.u32.b32(i32 -1, i32 0)
+  %hi = call i32 @llvm.amdgcn.v.mbcnt.hi.u32.b32(i32 -1, i32 %lo)
+  %4 = bitcast i32 %hi to float
+  call void @llvm.SI.export(i32 15, i32 1, i32 1, i32 0, i32 1, float %4, float %4, float %4, float %4)
+  ret void
+}
+
+declare i32 @llvm.amdgcn.v.mbcnt.lo.u32.b32(i32, i32) readnone
+
+declare i32 @llvm.amdgcn.v.mbcnt.hi.u32.b32(i32, i32) readnone
+
 declare i32 @llvm.SI.tid() readnone
----------------
The test should be in its own mbcnt file


http://reviews.llvm.org/D15475