[NVPTX] Avoid dots in global names

[asavonic]

Summary

It seems that ptxas cannot parse them:
ptxas fatal: Parsing error near '.2': syntax error

[tra]

Normally such replacement is done with nvptx-assign-valid-global-names pass, so we could pipe the tests through opt -nvptx-assign-valid-global-names which will ensure that tests don't have to know that dot is special for NVPTX. The downside is that it would have to be done for all llc tests. Perhaps we should consider moving the nvptx-assign-valid-global-names pass to the back-end.

Not using dots only fixes the issues we have now and would still leave us open to test breaks if a new NVPTX-invalid name is added by accident. I think running nvptx-assign-valid-global-names in llc may be a better solution.

[asavonic]

In D123041#3427128, @tra wrote:

Normally such replacement is done with nvptx-assign-valid-global-names pass, so we could pipe the tests through opt -nvptx-assign-valid-global-names which will ensure that tests don't have to know that dot is special for NVPTX. The downside is that it would have to be done for all llc tests. Perhaps we should consider moving the nvptx-assign-valid-global-names pass to the back-end.

Not using dots only fixes the issues we have now and would still leave us open to test breaks if a new NVPTX-invalid name is added by accident. I think running nvptx-assign-valid-global-names in llc may be a better solution.

NVPTXAssignValidGlobalNames should already run in llc, but it can only rename globals with local linkage. Globals with external linkage are not renamed, and I'm not sure what is the rationale here.

[tra]

LGTM. Automatic renaming should be addressed separately, anyways.

llvm/test/CodeGen/NVPTX/b52037.ll

 %int3 = type { i32, i32, i32 }
 %struct.spam.2 = type { %struct.foo.3, i16*, float, float, i32, float }
 %struct.foo.3 = type <{ %float4*, %float4*, %float4*, i32*, i32*, i32, i32, float }>
 %struct.zot = type { %struct.bar, [8 x i8], %struct.foo, [12 x i8] }
 %struct.bar = type { i32 (...)** }
 %struct.foo = type <{ i16*, %float4, %int3, i32, %float3, [4 x i8], i64, i32, i8, [3 x i8], i32 }>
 
 @global = external local_unnamed_addr addrspace(4) externally_initialized global [27 x %char3], align 1
-@global.1 = linkonce_odr unnamed_addr constant { [3 x i8*] } { [3 x i8*] [i8* inttoptr (i64 16 to i8*), i8* null, i8* null] }, align 8
+@global_1 = linkonce_odr unnamed_addr constant { [3 x i8*] } { [3 x i8*] [i8* inttoptr (i64 16 to i8*), i8* null, i8* null] }, align 8
 
 ; Function Attrs: argmemonly mustprogress nofree nounwind willreturn
 declare void @llvm.memcpy.p0i8.p0i8.i64(i8* noalias nocapture writeonly, i8* noalias nocapture readonly, i64, i1 immarg) #0
 
 declare %float4 @snork(float) local_unnamed_addr
 
-declare %float3 @bar.2(float, float) local_unnamed_addr
+declare %float3 @bar_2(float, float) local_unnamed_addr
 
 declare %float3 @zot() local_unnamed_addr
 
 declare %int3 @hoge(i32, i32, i32) local_unnamed_addr
 
 declare i64 @foo() local_unnamed_addr
 
 define void @barney(%struct.spam.2* nocapture readonly %arg) local_unnamed_addr {
 bb:
   tail call void asm sideeffect "// KEEP", ""() #1
   %tmp = alloca %struct.zot, align 16
   %tmp4 = getelementptr inbounds %struct.spam.2, %struct.spam.2* %arg, i64 0, i32 1
   %tmp5 = load i16*, i16** %tmp4, align 8
   %tmp6 = bitcast %struct.zot* %tmp to i8*
   %tmp9 = getelementptr inbounds %struct.zot, %struct.zot* %tmp, i64 0, i32 2, i32 1
   %0 = bitcast %float4* %tmp9 to i16**
   store i16* %tmp5, i16** %0, align 8
   %tmp10 = getelementptr inbounds %struct.zot, %struct.zot* %tmp, i64 0, i32 0, i32 0
-  store i32 (...)** bitcast (i8** getelementptr inbounds ({ [3 x i8*] }, { [3 x i8*] }* @global.1, i64 0, inrange i32 0, i64 3) to i32 (...)**), i32 (...)*** %tmp10, align 16
+  store i32 (...)** bitcast (i8** getelementptr inbounds ({ [3 x i8*] }, { [3 x i8*] }* @global_1, i64 0, inrange i32 0, i64 3) to i32 (...)**), i32 (...)*** %tmp10, align 16
   %tmp34 = getelementptr %struct.spam.2, %struct.spam.2* %arg, i64 0, i32 0, i32 0
   %tmp.i1 = tail call i64 @foo()
   %tmp44.i16 = getelementptr inbounds i16, i16* %tmp5, i64 undef
   %tmp45.i17 = load i16, i16* %tmp44.i16, align 2
   %tmp47.i18 = icmp eq i16 %tmp45.i17, -1
   br i1 %tmp47.i18, label %bb14, label %bb49.i.lr.ph
 
 bb49.i.lr.ph:                                     ; preds = %bb
bb49.i.lr.ph:                                     ; preds = %bb
   %tmp61.i.peel = load i32*, i32** %tmp60.i.peel, align 16
   %tmp10.i.i.peel = add nsw i8 %extract.t, -1
   store i8 %tmp10.i.i.peel, i8* %3, align 4
   %tmp13.i.i.peel = tail call %float3 @zot() #1
   %tmp15.i.i.peel = extractvalue %float3 %tmp13.i.i.peel, 0
   %tmp22.i.i.peel = fsub contract float %tmp19.i.i, %tmp15.i.i.peel
   %tmp17.i.i.peel = extractvalue %float3 %tmp13.i.i.peel, 2
   %tmp27.i.i.peel = fsub contract float %tmp24.i.i, %tmp17.i.i.peel
-  %tmp28.i.i.peel = tail call %float3 @bar.2(float %tmp22.i.i.peel, float %tmp27.i.i.peel) #1
+  %tmp28.i.i.peel = tail call %float3 @bar_2(float %tmp22.i.i.peel, float %tmp27.i.i.peel) #1
   %tmp28.i.elt.i.peel = extractvalue %float3 %tmp28.i.i.peel, 0
   store float %tmp28.i.elt.i.peel, float* %tmp59.i, align 16
   %tmp28.i.elt2.i.peel = extractvalue %float3 %tmp28.i.i.peel, 1
   store float %tmp28.i.elt2.i.peel, float* %6, align 4
   %tmp28.i.elt4.i.peel = extractvalue %float3 %tmp28.i.i.peel, 2
   store float %tmp28.i.elt4.i.peel, float* %7, align 8
   %tmp38.i.i.peel = zext i8 %tmp10.i.i.peel to i64
   %tmp39.i5.i.peel = getelementptr inbounds [27 x %char3], [27 x %char3] addrspace(4)* @global, i64 0, i64 %tmp38.i.i.peel
bb49.i:                                           ; preds = %bb49.i, %bb49.i.lr.ph.peel.newph
   %tmp45.i20.off11 = phi i8 [ %extract.t22.peel, %bb49.i.lr.ph.peel.newph ], [ %extract.t22, %bb49.i ]
   %tmp10.i.i = add nsw i8 %tmp45.i20.off11, -1
   store i8 %tmp10.i.i, i8* %3, align 4
   %tmp13.i.i = tail call %float3 @zot() #1
   %tmp15.i.i = extractvalue %float3 %tmp13.i.i, 0
   %tmp22.i.i = fsub contract float %tmp19.i.i, %tmp15.i.i
   %tmp17.i.i = extractvalue %float3 %tmp13.i.i, 2
   %tmp27.i.i = fsub contract float %tmp24.i.i, %tmp17.i.i
-  %tmp28.i.i = tail call %float3 @bar.2(float %tmp22.i.i, float %tmp27.i.i) #1
+  %tmp28.i.i = tail call %float3 @bar_2(float %tmp22.i.i, float %tmp27.i.i) #1
   %tmp28.i.elt.i = extractvalue %float3 %tmp28.i.i, 0
   store float %tmp28.i.elt.i, float* %tmp59.i, align 16
   %tmp28.i.elt2.i = extractvalue %float3 %tmp28.i.i, 1
   store float %tmp28.i.elt2.i, float* %6, align 4
   %tmp28.i.elt4.i = extractvalue %float3 %tmp28.i.i, 2
   store float %tmp28.i.elt4.i, float* %7, align 8
   %tmp38.i.i = zext i8 %tmp10.i.i to i64
   %tmp39.i5.i = getelementptr inbounds [27 x %char3], [27 x %char3] addrspace(4)* @global, i64 0, i64 %tmp38.i.i

llvm/test/CodeGen/NVPTX/barrier.ll

 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_30 -mattr=+ptx60 | FileCheck %s
 
 declare void @llvm.nvvm.bar.warp.sync(i32)
 declare void @llvm.nvvm.barrier.sync(i32)
 declare void @llvm.nvvm.barrier.sync.cnt(i32, i32)
 
-; CHECK-LABEL: .func{{.*}}barrier.sync
-define void @barrier.sync(i32 %id, i32 %cnt) {
-  ; CHECK: ld.param.u32   [[ID:%r[0-9]+]], [barrier.sync_param_0];
-  ; CHECK: ld.param.u32   [[CNT:%r[0-9]+]], [barrier.sync_param_1];
+; CHECK-LABEL: .func{{.*}}barrier_sync
+define void @barrier_sync(i32 %id, i32 %cnt) {
+  ; CHECK: ld.param.u32   [[ID:%r[0-9]+]], [barrier_sync_param_0];
+  ; CHECK: ld.param.u32   [[CNT:%r[0-9]+]], [barrier_sync_param_1];
 
   ; CHECK:  barrier.sync [[ID]], [[CNT]];
   call void @llvm.nvvm.barrier.sync.cnt(i32 %id, i32 %cnt)
   ; CHECK:  barrier.sync [[ID]], 2;
   call void @llvm.nvvm.barrier.sync.cnt(i32 %id, i32 2)
   ; CHECK:  barrier.sync 3, [[CNT]];
   call void @llvm.nvvm.barrier.sync.cnt(i32 3, i32 %cnt)
   ; CHECK:  barrier.sync 4, 5;

llvm/test/CodeGen/NVPTX/ld-st-addrrspace.py

addrspace_id = {
     ".const": 4,
     ".local": 5,
     ".param": 101
 }
 
 
 def gen_load_tests():
   load_template = """
-define ${type} @ld${_volatile}${_space}.${ptx_type}(${type} addrspace(${asid})* %ptr) {
+define ${type} @${testname}(${type} addrspace(${asid})* %ptr) {
+; CHECK: ${testname}
 ; CHECK_P32: ld${_volatile}${_volatile_as}.${ptx_type} %${ptx_reg}{{[0-9]+}}, [%r{{[0-9]+}}]
 ; CHECK_P64: ld${_volatile}${_volatile_as}.${ptx_type} %${ptx_reg}{{[0-9]+}}, [%rd{{[0-9]+}}]
 ; CHECK: ret
   %p = ${generic_ptr}
   %a = load ${volatile} ${type}, ${type}* %p
   ret ${type} %a
 }
 """
params = {
         "_volatile": ".volatile" if volatile else "",
         "_volatile_as": volatile_as,
         "_space": space,
         "ptx_reg": llvm_type_to_ptx_reg[op_type],
         "ptx_type": llvm_type_to_ptx_type[op_type],
         "asid": addrspace_id[space],
     }
 
+    testname = \
+      Template("ld_${_volatile}${_space}.${ptx_type}").substitute(params)
+    params["testname"] = testname.replace(".", "_")
+
     # LLVM does not accept "addrspacecast Type* addrspace(0) to Type*", so we
     # need to avoid it for generic pointer tests.
     if space:
       generic_ptr_template = ("addrspacecast ${type} addrspace(${asid})* %ptr "
                               "to ${type}*")
     else:
       generic_ptr_template = "select i1 true, ${type}* %ptr, ${type}* %ptr"
     params["generic_ptr"] = Template(generic_ptr_template).substitute(params)

llvm/test/CodeGen/NVPTX/match.ll

 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_70 -mattr=+ptx60 | FileCheck %s
 
 declare i32 @llvm.nvvm.match.any.sync.i32(i32, i32)
 declare i32 @llvm.nvvm.match.any.sync.i64(i32, i64)
 
-; CHECK-LABEL: .func{{.*}}match.any.sync.i32
-define i32 @match.any.sync.i32(i32 %mask, i32 %value) {
-  ; CHECK: ld.param.u32   [[MASK:%r[0-9]+]], [match.any.sync.i32_param_0];
-  ; CHECK: ld.param.u32   [[VALUE:%r[0-9]+]], [match.any.sync.i32_param_1];
+; CHECK-LABEL: .func{{.*}}match_any_sync_i32
+define i32 @match_any_sync_i32(i32 %mask, i32 %value) {
+  ; CHECK: ld.param.u32   [[MASK:%r[0-9]+]], [match_any_sync_i32_param_0];
+  ; CHECK: ld.param.u32   [[VALUE:%r[0-9]+]], [match_any_sync_i32_param_1];
 
   ; CHECK:  match.any.sync.b32  [[V0:%r[0-9]+]], [[VALUE]], [[MASK]];
   %v0 = call i32 @llvm.nvvm.match.any.sync.i32(i32 %mask, i32 %value)
   ; CHECK:  match.any.sync.b32  [[V1:%r[0-9]+]], [[VALUE]], 1;
   %v1 = call i32 @llvm.nvvm.match.any.sync.i32(i32 1, i32 %value)
   ; CHECK:  match.any.sync.b32  [[V2:%r[0-9]+]], 2, [[MASK]];
   %v2 = call i32 @llvm.nvvm.match.any.sync.i32(i32 %mask, i32 2)
   ; CHECK:  match.any.sync.b32  [[V3:%r[0-9]+]], 4, 3;
   %v3 = call i32 @llvm.nvvm.match.any.sync.i32(i32 3, i32 4)
   %sum1 = add i32 %v0, %v1
   %sum2 = add i32 %v2, %v3
   %sum3 = add i32 %sum1, %sum2
   ret i32 %sum3;
 }
 
-; CHECK-LABEL: .func{{.*}}match.any.sync.i64
-define i32 @match.any.sync.i64(i32 %mask, i64 %value) {
-  ; CHECK: ld.param.u32   [[MASK:%r[0-9]+]], [match.any.sync.i64_param_0];
-  ; CHECK: ld.param.u64   [[VALUE:%rd[0-9]+]], [match.any.sync.i64_param_1];
+; CHECK-LABEL: .func{{.*}}match_any_sync_i64
+define i32 @match_any_sync_i64(i32 %mask, i64 %value) {
+  ; CHECK: ld.param.u32   [[MASK:%r[0-9]+]], [match_any_sync_i64_param_0];
+  ; CHECK: ld.param.u64   [[VALUE:%rd[0-9]+]], [match_any_sync_i64_param_1];
 
   ; CHECK:  match.any.sync.b64  [[V0:%r[0-9]+]], [[VALUE]], [[MASK]];
   %v0 = call i32 @llvm.nvvm.match.any.sync.i64(i32 %mask, i64 %value)
   ; CHECK:  match.any.sync.b64  [[V1:%r[0-9]+]], [[VALUE]], 1;
   %v1 = call i32 @llvm.nvvm.match.any.sync.i64(i32 1, i64 %value)
   ; CHECK:  match.any.sync.b64  [[V2:%r[0-9]+]], 2, [[MASK]];
   %v2 = call i32 @llvm.nvvm.match.any.sync.i64(i32 %mask, i64 2)
   ; CHECK:  match.any.sync.b64  [[V3:%r[0-9]+]], 4, 3;
   %v3 = call i32 @llvm.nvvm.match.any.sync.i64(i32 3, i64 4)
   %sum1 = add i32 %v0, %v1
   %sum2 = add i32 %v2, %v3
   %sum3 = add i32 %sum1, %sum2
   ret i32 %sum3;
 }
 
 declare {i32, i1} @llvm.nvvm.match.all.sync.i32p(i32, i32)
 declare {i32, i1} @llvm.nvvm.match.all.sync.i64p(i32, i64)
 
-; CHECK-LABEL: .func{{.*}}match.all.sync.i32p(
-define {i32,i1} @match.all.sync.i32p(i32 %mask, i32 %value) {
-  ; CHECK: ld.param.u32   [[MASK:%r[0-9]+]], [match.all.sync.i32p_param_0];
-  ; CHECK: ld.param.u32   [[VALUE:%r[0-9]+]], [match.all.sync.i32p_param_1];
+; CHECK-LABEL: .func{{.*}}match_all_sync_i32p(
+define {i32,i1} @match_all_sync_i32p(i32 %mask, i32 %value) {
+  ; CHECK: ld.param.u32   [[MASK:%r[0-9]+]], [match_all_sync_i32p_param_0];
+  ; CHECK: ld.param.u32   [[VALUE:%r[0-9]+]], [match_all_sync_i32p_param_1];
 
   ; CHECK:  match.all.sync.b32 {{%r[0-9]+\|%p[0-9]+}}, [[VALUE]], [[MASK]];
   %r1 = call {i32, i1} @llvm.nvvm.match.all.sync.i32p(i32 %mask, i32 %value)
   %v1 = extractvalue {i32, i1} %r1, 0
   %p1 = extractvalue {i32, i1} %r1, 1
 
   ; CHECK:  match.all.sync.b32 {{%r[0-9]+\|%p[0-9]+}}, 1, [[MASK]];
   %r2 = call {i32, i1} @llvm.nvvm.match.all.sync.i32p(i32 %mask, i32 1)
define {i32,i1} @match_all_sync_i32p(i32 %mask, i32 %value) {
   %psum1 = add i1 %p1, %p2
   %psum2 = add i1 %p3, %p4
   %psum3 = add i1 %psum1, %psum2
   %ret0 = insertvalue {i32, i1} undef, i32 %vsum3, 0
   %ret1 = insertvalue {i32, i1} %ret0, i1 %psum3, 1
   ret {i32, i1} %ret1;
 }
 
-; CHECK-LABEL: .func{{.*}}match.all.sync.i64p(
-define {i32,i1} @match.all.sync.i64p(i32 %mask, i64 %value) {
-  ; CHECK: ld.param.u32   [[MASK:%r[0-9]+]], [match.all.sync.i64p_param_0];
-  ; CHECK: ld.param.u64   [[VALUE:%rd[0-9]+]], [match.all.sync.i64p_param_1];
+; CHECK-LABEL: .func{{.*}}match_all_sync_i64p(
+define {i32,i1} @match_all_sync_i64p(i32 %mask, i64 %value) {
+  ; CHECK: ld.param.u32   [[MASK:%r[0-9]+]], [match_all_sync_i64p_param_0];
+  ; CHECK: ld.param.u64   [[VALUE:%rd[0-9]+]], [match_all_sync_i64p_param_1];
 
   ; CHECK:  match.all.sync.b64 {{%r[0-9]+\|%p[0-9]+}}, [[VALUE]], [[MASK]];
   %r1 = call {i32, i1} @llvm.nvvm.match.all.sync.i64p(i32 %mask, i64 %value)
   %v1 = extractvalue {i32, i1} %r1, 0
   %p1 = extractvalue {i32, i1} %r1, 1
 
   ; CHECK:  match.all.sync.b64 {{%r[0-9]+\|%p[0-9]+}}, 1, [[MASK]];
   %r2 = call {i32, i1} @llvm.nvvm.match.all.sync.i64p(i32 %mask, i64 1)

llvm/test/CodeGen/NVPTX/shfl-p.ll

 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_30 | FileCheck %s
 
 declare {i32, i1} @llvm.nvvm.shfl.down.i32p(i32, i32, i32)
 declare {float, i1} @llvm.nvvm.shfl.down.f32p(float, i32, i32)
 declare {i32, i1} @llvm.nvvm.shfl.up.i32p(i32, i32, i32)
 declare {float, i1} @llvm.nvvm.shfl.up.f32p(float, i32, i32)
 declare {i32, i1} @llvm.nvvm.shfl.bfly.i32p(i32, i32, i32)
 declare {float, i1} @llvm.nvvm.shfl.bfly.f32p(float, i32, i32)
 declare {i32, i1} @llvm.nvvm.shfl.idx.i32p(i32, i32, i32)
 declare {float, i1} @llvm.nvvm.shfl.idx.f32p(float, i32, i32)
 
-; CHECK-LABEL: .func{{.*}}shfl.i32.rrr
-define {i32, i1} @shfl.i32.rrr(i32 %a, i32 %b, i32 %c) {
+; CHECK-LABEL: .func{{.*}}shfl_i32_rrr
+define {i32, i1} @shfl_i32_rrr(i32 %a, i32 %b, i32 %c) {
   ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
   ; CHECK: shfl.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], [[C]];
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {i32, i1} @llvm.nvvm.shfl.down.i32p(i32 %a, i32 %b, i32 %c)
   ret {i32, i1} %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.i32.irr
-define {i32, i1} @shfl.i32.irr(i32 %a, i32 %b, i32 %c) {
+; CHECK-LABEL: .func{{.*}}shfl_i32_irr
+define {i32, i1} @shfl_i32_irr(i32 %a, i32 %b, i32 %c) {
   ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
   ; CHECK: shfl.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], [[C]];
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {i32, i1} @llvm.nvvm.shfl.down.i32p(i32 %a, i32 %b, i32 %c)
   ret {i32, i1} %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.i32.rri
-define {i32, i1} @shfl.i32.rri(i32 %a, i32 %b) {
+; CHECK-LABEL: .func{{.*}}shfl_i32_rri
+define {i32, i1} @shfl_i32_rri(i32 %a, i32 %b) {
   ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
   ; CHECK: shfl.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], 1;
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {i32, i1} @llvm.nvvm.shfl.down.i32p(i32 %a, i32 %b, i32 1)
   ret {i32, i1} %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.i32.iri
-define {i32, i1} @shfl.i32.iri(i32 %a, i32 %b) {
+; CHECK-LABEL: .func{{.*}}shfl_i32_iri
+define {i32, i1} @shfl_i32_iri(i32 %a, i32 %b) {
   ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
   ; CHECK: shfl.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], 2;
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {i32, i1} @llvm.nvvm.shfl.down.i32p(i32 %a, i32 %b, i32 2)
   ret {i32, i1} %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.i32.rir
-define {i32, i1} @shfl.i32.rir(i32 %a, i32 %c) {
+; CHECK-LABEL: .func{{.*}}shfl_i32_rir
+define {i32, i1} @shfl_i32_rir(i32 %a, i32 %c) {
   ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
   ; CHECK: shfl.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 1, [[C]];
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {i32, i1} @llvm.nvvm.shfl.down.i32p(i32 %a, i32 1, i32 %c)
   ret {i32, i1} %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.i32.iir
-define {i32, i1} @shfl.i32.iir(i32 %a, i32 %c) {
+; CHECK-LABEL: .func{{.*}}shfl_i32_iir
+define {i32, i1} @shfl_i32_iir(i32 %a, i32 %c) {
   ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
   ; CHECK: shfl.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 2, [[C]];
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {i32, i1} @llvm.nvvm.shfl.down.i32p(i32 %a, i32 2, i32 %c)
   ret {i32, i1} %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.i32.rii
-define {i32, i1} @shfl.i32.rii(i32 %a) {
+; CHECK-LABEL: .func{{.*}}shfl_i32_rii
+define {i32, i1} @shfl_i32_rii(i32 %a) {
   ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
   ; CHECK: shfl.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 1, 2;
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {i32, i1} @llvm.nvvm.shfl.down.i32p(i32 %a, i32 1, i32 2)
   ret {i32, i1} %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.i32.iii
-define {i32, i1} @shfl.i32.iii(i32 %a, i32 %b) {
+; CHECK-LABEL: .func{{.*}}shfl_i32_iii
+define {i32, i1} @shfl_i32_iii(i32 %a, i32 %b) {
   ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
   ; CHECK: shfl.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 2, 3;
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {i32, i1} @llvm.nvvm.shfl.down.i32p(i32 %a, i32 2, i32 3)
   ret {i32, i1} %val
 }
 
 ;; Same intrinsics, but for float
 
-; CHECK-LABEL: .func{{.*}}shfl.f32.rrr
-define {float, i1} @shfl.f32.rrr(float %a, i32 %b, i32 %c) {
+; CHECK-LABEL: .func{{.*}}shfl_f32_rrr
+define {float, i1} @shfl_f32_rrr(float %a, i32 %b, i32 %c) {
   ; CHECK: ld.param.f32 [[A:%f[0-9]+]]
   ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
   ; CHECK: shfl.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], [[C]];
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {float, i1} @llvm.nvvm.shfl.down.f32p(float %a, i32 %b, i32 %c)
   ret {float, i1} %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.f32.irr
-define {float, i1} @shfl.f32.irr(float %a, i32 %b, i32 %c) {
+; CHECK-LABEL: .func{{.*}}shfl_f32_irr
+define {float, i1} @shfl_f32_irr(float %a, i32 %b, i32 %c) {
   ; CHECK: ld.param.f32 [[A:%f[0-9]+]]
   ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
   ; CHECK: shfl.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], [[C]];
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {float, i1} @llvm.nvvm.shfl.down.f32p(float %a, i32 %b, i32 %c)
   ret {float, i1} %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.f32.rri
-define {float, i1} @shfl.f32.rri(float %a, i32 %b) {
+; CHECK-LABEL: .func{{.*}}shfl_f32_rri
+define {float, i1} @shfl_f32_rri(float %a, i32 %b) {
   ; CHECK: ld.param.f32 [[A:%f[0-9]+]]
   ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
   ; CHECK: shfl.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], 1;
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {float, i1} @llvm.nvvm.shfl.down.f32p(float %a, i32 %b, i32 1)
   ret {float, i1} %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.f32.iri
-define {float, i1} @shfl.f32.iri(float %a, i32 %b) {
+; CHECK-LABEL: .func{{.*}}shfl_f32_iri
+define {float, i1} @shfl_f32_iri(float %a, i32 %b) {
   ; CHECK: ld.param.f32 [[A:%f[0-9]+]]
   ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
   ; CHECK: shfl.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], 2;
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {float, i1} @llvm.nvvm.shfl.down.f32p(float %a, i32 %b, i32 2)
   ret {float, i1} %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.f32.rir
-define {float, i1} @shfl.f32.rir(float %a, i32 %c) {
+; CHECK-LABEL: .func{{.*}}shfl_f32_rir
+define {float, i1} @shfl_f32_rir(float %a, i32 %c) {
   ; CHECK: ld.param.f32 [[A:%f[0-9]+]]
   ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
   ; CHECK: shfl.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 1, [[C]];
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {float, i1} @llvm.nvvm.shfl.down.f32p(float %a, i32 1, i32 %c)
   ret {float, i1} %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.f32.iir
-define {float, i1} @shfl.f32.iir(float %a, i32 %c) {
+; CHECK-LABEL: .func{{.*}}shfl_f32_iir
+define {float, i1} @shfl_f32_iir(float %a, i32 %c) {
   ; CHECK: ld.param.f32 [[A:%f[0-9]+]]
   ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
   ; CHECK: shfl.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 2, [[C]];
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {float, i1} @llvm.nvvm.shfl.down.f32p(float %a, i32 2, i32 %c)
   ret {float, i1} %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.f32.rii
-define {float, i1} @shfl.f32.rii(float %a) {
+; CHECK-LABEL: .func{{.*}}shfl_f32_rii
+define {float, i1} @shfl_f32_rii(float %a) {
   ; CHECK: ld.param.f32 [[A:%f[0-9]+]]
   ; CHECK: shfl.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 1, 2;
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {float, i1} @llvm.nvvm.shfl.down.f32p(float %a, i32 1, i32 2)
   ret {float, i1} %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.f32.iii
-define {float, i1} @shfl.f32.iii(float %a, i32 %b) {
+; CHECK-LABEL: .func{{.*}}shfl_f32_iii
+define {float, i1} @shfl_f32_iii(float %a, i32 %b) {
   ; CHECK: ld.param.f32 [[A:%f[0-9]+]]
   ; CHECK: shfl.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 2, 3;
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {float, i1} @llvm.nvvm.shfl.down.f32p(float %a, i32 2, i32 3)
   ret {float, i1} %val
 }

llvm/test/CodeGen/NVPTX/shfl-sync-p.ll

 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_30 -mattr=+ptx60 | FileCheck %s
 
 declare {i32, i1} @llvm.nvvm.shfl.sync.down.i32p(i32, i32, i32, i32)
 declare {float, i1} @llvm.nvvm.shfl.sync.down.f32p(i32, float, i32, i32)
 declare {i32, i1} @llvm.nvvm.shfl.sync.up.i32p(i32, i32, i32, i32)
 declare {float, i1} @llvm.nvvm.shfl.sync.up.f32p(i32, float, i32, i32)
 declare {i32, i1} @llvm.nvvm.shfl.sync.bfly.i32p(i32, i32, i32, i32)
 declare {float, i1} @llvm.nvvm.shfl.sync.bfly.f32p(i32, float, i32, i32)
 declare {i32, i1} @llvm.nvvm.shfl.sync.idx.i32p(i32, i32, i32, i32)
 declare {float, i1} @llvm.nvvm.shfl.sync.idx.f32p(i32, float, i32, i32)
 
-; CHECK-LABEL: .func{{.*}}shfl.sync.i32.rrr
-define {i32, i1} @shfl.sync.i32.rrr(i32 %mask, i32 %a, i32 %b, i32 %c) {
+; CHECK-LABEL: .func{{.*}}shfl_sync_i32_rrr
+define {i32, i1} @shfl_sync_i32_rrr(i32 %mask, i32 %a, i32 %b, i32 %c) {
   ; CHECK: ld.param.u32 [[MASK:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
   ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], [[C]], [[MASK]];
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {i32, i1} @llvm.nvvm.shfl.sync.down.i32p(i32 %mask, i32 %a, i32 %b, i32 %c)
   ret {i32, i1} %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.sync.i32.irr
-define {i32, i1} @shfl.sync.i32.irr(i32 %a, i32 %b, i32 %c) {
+; CHECK-LABEL: .func{{.*}}shfl_sync_i32_irr
+define {i32, i1} @shfl_sync_i32_irr(i32 %a, i32 %b, i32 %c) {
   ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
   ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], [[C]], 1;
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {i32, i1} @llvm.nvvm.shfl.sync.down.i32p(i32 1, i32 %a, i32 %b, i32 %c)
   ret {i32, i1} %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.sync.i32.rri
-define {i32, i1} @shfl.sync.i32.rri(i32 %mask, i32 %a, i32 %b) {
+; CHECK-LABEL: .func{{.*}}shfl_sync_i32_rri
+define {i32, i1} @shfl_sync_i32_rri(i32 %mask, i32 %a, i32 %b) {
   ; CHECK: ld.param.u32 [[MASK:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
   ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], 1, [[MASK]];
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {i32, i1} @llvm.nvvm.shfl.sync.down.i32p(i32 %mask, i32 %a, i32 %b, i32 1)
   ret {i32, i1} %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.sync.i32.iri
-define {i32, i1} @shfl.sync.i32.iri(i32 %a, i32 %b) {
+; CHECK-LABEL: .func{{.*}}shfl_sync_i32_iri
+define {i32, i1} @shfl_sync_i32_iri(i32 %a, i32 %b) {
   ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
   ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], 2, 1;
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {i32, i1} @llvm.nvvm.shfl.sync.down.i32p(i32 1, i32 %a, i32 %b, i32 2)
   ret {i32, i1} %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.sync.i32.rir
-define {i32, i1} @shfl.sync.i32.rir(i32 %mask, i32 %a, i32 %c) {
+; CHECK-LABEL: .func{{.*}}shfl_sync_i32_rir
+define {i32, i1} @shfl_sync_i32_rir(i32 %mask, i32 %a, i32 %c) {
   ; CHECK: ld.param.u32 [[MASK:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
   ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 1, [[C]], [[MASK]];
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {i32, i1} @llvm.nvvm.shfl.sync.down.i32p(i32 %mask, i32 %a, i32 1, i32 %c)
   ret {i32, i1} %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.sync.i32.iir
-define {i32, i1} @shfl.sync.i32.iir(i32 %a, i32 %c) {
+; CHECK-LABEL: .func{{.*}}shfl_sync_i32_iir
+define {i32, i1} @shfl_sync_i32_iir(i32 %a, i32 %c) {
   ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
   ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 2, [[C]], 1;
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {i32, i1} @llvm.nvvm.shfl.sync.down.i32p(i32 1, i32 %a, i32 2, i32 %c)
   ret {i32, i1} %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.sync.i32.rii
-define {i32, i1} @shfl.sync.i32.rii(i32 %mask, i32 %a) {
+; CHECK-LABEL: .func{{.*}}shfl_sync_i32_rii
+define {i32, i1} @shfl_sync_i32_rii(i32 %mask, i32 %a) {
   ; CHECK: ld.param.u32 [[MASK:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
   ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 1, 2, [[MASK]];
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {i32, i1} @llvm.nvvm.shfl.sync.down.i32p(i32 %mask, i32 %a, i32 1, i32 2)
   ret {i32, i1} %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.sync.i32.iii
-define {i32, i1} @shfl.sync.i32.iii(i32 %a, i32 %b) {
+; CHECK-LABEL: .func{{.*}}shfl_sync_i32_iii
+define {i32, i1} @shfl_sync_i32_iii(i32 %a, i32 %b) {
   ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
   ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 2, 3, 1;
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {i32, i1} @llvm.nvvm.shfl.sync.down.i32p(i32 1, i32 %a, i32 2, i32 3)
   ret {i32, i1} %val
 }
 
 ;; Same intrinsics, but for float
 
-; CHECK-LABEL: .func{{.*}}shfl.sync.f32.rrr
-define {float, i1} @shfl.sync.f32.rrr(i32 %mask, float %a, i32 %b, i32 %c) {
+; CHECK-LABEL: .func{{.*}}shfl_sync_f32_rrr
+define {float, i1} @shfl_sync_f32_rrr(i32 %mask, float %a, i32 %b, i32 %c) {
   ; CHECK: ld.param.u32 [[MASK:%r[0-9]+]]
   ; CHECK: ld.param.f32 [[A:%f[0-9]+]]
   ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
   ; CHECK: shfl.sync.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], [[C]], [[MASK]];
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {float, i1} @llvm.nvvm.shfl.sync.down.f32p(i32 %mask, float %a, i32 %b, i32 %c)
   ret {float, i1} %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.sync.f32.irr
-define {float, i1} @shfl.sync.f32.irr(float %a, i32 %b, i32 %c) {
+; CHECK-LABEL: .func{{.*}}shfl_sync_f32_irr
+define {float, i1} @shfl_sync_f32_irr(float %a, i32 %b, i32 %c) {
   ; CHECK: ld.param.f32 [[A:%f[0-9]+]]
   ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
   ; CHECK: shfl.sync.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], [[C]], 1;
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {float, i1} @llvm.nvvm.shfl.sync.down.f32p(i32 1, float %a, i32 %b, i32 %c)
   ret {float, i1} %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.sync.f32.rri
-define {float, i1} @shfl.sync.f32.rri(i32 %mask, float %a, i32 %b) {
+; CHECK-LABEL: .func{{.*}}shfl_sync_f32_rri
+define {float, i1} @shfl_sync_f32_rri(i32 %mask, float %a, i32 %b) {
   ; CHECK: ld.param.u32 [[MASK:%r[0-9]+]]
   ; CHECK: ld.param.f32 [[A:%f[0-9]+]]
   ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
   ; CHECK: shfl.sync.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], 1, [[MASK]];
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {float, i1} @llvm.nvvm.shfl.sync.down.f32p(i32 %mask, float %a, i32 %b, i32 1)
   ret {float, i1} %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.sync.f32.iri
-define {float, i1} @shfl.sync.f32.iri(float %a, i32 %b) {
+; CHECK-LABEL: .func{{.*}}shfl_sync_f32_iri
+define {float, i1} @shfl_sync_f32_iri(float %a, i32 %b) {
   ; CHECK: ld.param.f32 [[A:%f[0-9]+]]
   ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
   ; CHECK: shfl.sync.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], [[B]], 2, 1;
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {float, i1} @llvm.nvvm.shfl.sync.down.f32p(i32 1, float %a, i32 %b, i32 2)
   ret {float, i1} %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.sync.f32.rir
-define {float, i1} @shfl.sync.f32.rir(i32 %mask, float %a, i32 %c) {
+; CHECK-LABEL: .func{{.*}}shfl_sync_f32_rir
+define {float, i1} @shfl_sync_f32_rir(i32 %mask, float %a, i32 %c) {
   ; CHECK: ld.param.u32 [[MASK:%r[0-9]+]]
   ; CHECK: ld.param.f32 [[A:%f[0-9]+]]
   ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
   ; CHECK: shfl.sync.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 1, [[C]], [[MASK]];
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {float, i1} @llvm.nvvm.shfl.sync.down.f32p(i32 %mask, float %a, i32 1, i32 %c)
   ret {float, i1} %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.sync.f32.iir
-define {float, i1} @shfl.sync.f32.iir(float %a, i32 %c) {
+; CHECK-LABEL: .func{{.*}}shfl_sync_f32_iir
+define {float, i1} @shfl_sync_f32_iir(float %a, i32 %c) {
   ; CHECK: ld.param.f32 [[A:%f[0-9]+]]
   ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
   ; CHECK: shfl.sync.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 2, [[C]], 1;
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {float, i1} @llvm.nvvm.shfl.sync.down.f32p(i32 1, float %a, i32 2, i32 %c)
   ret {float, i1} %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.sync.f32.rii
-define {float, i1} @shfl.sync.f32.rii(i32 %mask, float %a) {
+; CHECK-LABEL: .func{{.*}}shfl_sync_f32_rii
+define {float, i1} @shfl_sync_f32_rii(i32 %mask, float %a) {
   ; CHECK: ld.param.u32 [[MASK:%r[0-9]+]]
   ; CHECK: ld.param.f32 [[A:%f[0-9]+]]
   ; CHECK: shfl.sync.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 1, 2, [[MASK]];
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {float, i1} @llvm.nvvm.shfl.sync.down.f32p(i32 %mask, float %a, i32 1, i32 2)
   ret {float, i1} %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.sync.f32.iii
-define {float, i1} @shfl.sync.f32.iii(float %a, i32 %b) {
+; CHECK-LABEL: .func{{.*}}shfl_sync_f32_iii
+define {float, i1} @shfl_sync_f32_iii(float %a, i32 %b) {
   ; CHECK: ld.param.f32 [[A:%f[0-9]+]]
   ; CHECK: shfl.sync.down.b32 [[OUT:%f[0-9]+]]|[[OUTP:%p[0-9]+]], [[A]], 2, 3, 1;
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call {float, i1} @llvm.nvvm.shfl.sync.down.f32p(i32 1, float %a, i32 2, i32 3)
   ret {float, i1} %val
 }

llvm/test/CodeGen/NVPTX/shfl-sync.ll

 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_30 -mattr=+ptx60 | FileCheck %s
 
 declare i32 @llvm.nvvm.shfl.sync.down.i32(i32, i32, i32, i32)
 declare float @llvm.nvvm.shfl.sync.down.f32(float, i32, i32, i32)
 declare i32 @llvm.nvvm.shfl.sync.up.i32(i32, i32, i32, i32)
 declare float @llvm.nvvm.shfl.sync.up.f32(float, i32, i32, i32)
 declare i32 @llvm.nvvm.shfl.sync.bfly.i32(i32, i32, i32, i32)
 declare float @llvm.nvvm.shfl.sync.bfly.f32(float, i32, i32, i32)
 declare i32 @llvm.nvvm.shfl.sync.idx.i32(i32, i32, i32, i32)
 declare float @llvm.nvvm.shfl.sync.idx.f32(float, i32, i32, i32)
 
-; CHECK-LABEL: .func{{.*}}shfl.sync.rrr
-define i32 @shfl.sync.rrr(i32 %mask, i32 %a, i32 %b, i32 %c) {
+; CHECK-LABEL: .func{{.*}}shfl_sync_rrr
+define i32 @shfl_sync_rrr(i32 %mask, i32 %a, i32 %b, i32 %c) {
   ; CHECK: ld.param.u32 [[MASK:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
   ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], [[C]], [[MASK]];
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 %mask, i32 %a, i32 %b, i32 %c)
   ret i32 %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.sync.irr
-define i32 @shfl.sync.irr(i32 %a, i32 %b, i32 %c) {
+; CHECK-LABEL: .func{{.*}}shfl_sync_irr
+define i32 @shfl_sync_irr(i32 %a, i32 %b, i32 %c) {
   ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
   ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], [[C]], 1;
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 1, i32 %a, i32 %b, i32 %c)
   ret i32 %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.sync.rri
-define i32 @shfl.sync.rri(i32 %mask, i32 %a, i32 %b) {
+; CHECK-LABEL: .func{{.*}}shfl_sync_rri
+define i32 @shfl_sync_rri(i32 %mask, i32 %a, i32 %b) {
   ; CHECK: ld.param.u32 [[MASK:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
   ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], 1, [[MASK]];
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 %mask, i32 %a, i32 %b, i32 1)
   ret i32 %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.sync.iri
-define i32 @shfl.sync.iri(i32 %a, i32 %b) {
+; CHECK-LABEL: .func{{.*}}shfl_sync_iri
+define i32 @shfl_sync_iri(i32 %a, i32 %b) {
   ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[B:%r[0-9]+]]
   ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], [[B]], 2, 1;
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 1, i32 %a, i32 %b, i32 2)
   ret i32 %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.sync.rir
-define i32 @shfl.sync.rir(i32 %mask, i32 %a, i32 %c) {
+; CHECK-LABEL: .func{{.*}}shfl_sync_rir
+define i32 @shfl_sync_rir(i32 %mask, i32 %a, i32 %c) {
   ; CHECK: ld.param.u32 [[MASK:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
   ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 1, [[C]], [[MASK]];
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 %mask, i32 %a, i32 1, i32 %c)
   ret i32 %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.sync.iir
-define i32 @shfl.sync.iir(i32 %a, i32 %c) {
+; CHECK-LABEL: .func{{.*}}shfl_sync_iir
+define i32 @shfl_sync_iir(i32 %a, i32 %c) {
   ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[C:%r[0-9]+]]
   ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 2, [[C]], 1;
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 1, i32 %a, i32 2, i32 %c)
   ret i32 %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.sync.rii
-define i32 @shfl.sync.rii(i32 %mask, i32 %a) {
+; CHECK-LABEL: .func{{.*}}shfl_sync_rii
+define i32 @shfl_sync_rii(i32 %mask, i32 %a) {
   ; CHECK: ld.param.u32 [[MASK:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
   ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 1, 2, [[MASK]];
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 %mask, i32 %a, i32 1, i32 2)
   ret i32 %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.sync.iii
-define i32 @shfl.sync.iii(i32 %a, i32 %b) {
+; CHECK-LABEL: .func{{.*}}shfl_sync_iii
+define i32 @shfl_sync_iii(i32 %a, i32 %b) {
   ; CHECK: ld.param.u32 [[A:%r[0-9]+]]
   ; CHECK: shfl.sync.down.b32 [[OUT:%r[0-9]+]], [[A]], 2, 3, 1;
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call i32 @llvm.nvvm.shfl.sync.down.i32(i32 1, i32 %a, i32 2, i32 3)
   ret i32 %val
 }

llvm/test/CodeGen/NVPTX/shfl.ll

 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_30 | FileCheck %s
 
 declare i32 @llvm.nvvm.shfl.down.i32(i32, i32, i32)
 declare float @llvm.nvvm.shfl.down.f32(float, i32, i32)
 declare i32 @llvm.nvvm.shfl.up.i32(i32, i32, i32)
 declare float @llvm.nvvm.shfl.up.f32(float, i32, i32)
 declare i32 @llvm.nvvm.shfl.bfly.i32(i32, i32, i32)
 declare float @llvm.nvvm.shfl.bfly.f32(float, i32, i32)
 declare i32 @llvm.nvvm.shfl.idx.i32(i32, i32, i32)
 declare float @llvm.nvvm.shfl.idx.f32(float, i32, i32)
 
 ; Try all four permutations of register and immediate parameters with
 ; shfl.down.
 
-; CHECK-LABEL: .func{{.*}}shfl.down1
-define i32 @shfl.down1(i32 %in) {
+; CHECK-LABEL: .func{{.*}}shfl_down1
+define i32 @shfl_down1(i32 %in) {
   ; CHECK: ld.param.u32 [[IN:%r[0-9]+]]
   ; CHECK: shfl.down.b32 [[OUT:%r[0-9]+]], [[IN]], 1, 2;
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %val = call i32 @llvm.nvvm.shfl.down.i32(i32 %in, i32 1, i32 2)
   ret i32 %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.down2
-define i32 @shfl.down2(i32 %in, i32 %width) {
+; CHECK-LABEL: .func{{.*}}shfl_down2
+define i32 @shfl_down2(i32 %in, i32 %width) {
   ; CHECK: ld.param.u32 [[IN1:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[IN2:%r[0-9]+]]
   ; CHECK: shfl.down.{{.}}32 %r{{[0-9]+}}, [[IN1]], [[IN2]], 3;
   %val = call i32 @llvm.nvvm.shfl.down.i32(i32 %in, i32 %width, i32 3)
   ret i32 %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.down3
-define i32 @shfl.down3(i32 %in, i32 %mask) {
+; CHECK-LABEL: .func{{.*}}shfl_down3
+define i32 @shfl_down3(i32 %in, i32 %mask) {
   ; CHECK: ld.param.u32 [[IN1:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[IN2:%r[0-9]+]]
   ; CHECK: shfl.down.{{.}}32 %r{{[0-9]+}}, [[IN1]], 4, [[IN2]];
   %val = call i32 @llvm.nvvm.shfl.down.i32(i32 %in, i32 4, i32 %mask)
   ret i32 %val
 }
 
-; CHECK-LABEL: .func{{.*}}shfl.down4
-define i32 @shfl.down4(i32 %in, i32 %width, i32 %mask) {
+; CHECK-LABEL: .func{{.*}}shfl_down4
+define i32 @shfl_down4(i32 %in, i32 %width, i32 %mask) {
   ; CHECK: ld.param.u32 [[IN1:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[IN2:%r[0-9]+]]
   ; CHECK: ld.param.u32 [[IN3:%r[0-9]+]]
   ; CHECK: shfl.down.{{.}}32 %r{{[0-9]+}}, [[IN1]], [[IN2]], [[IN3]];
   %val = call i32 @llvm.nvvm.shfl.down.i32(i32 %in, i32 %width, i32 %mask)
   ret i32 %val
 }
 
 ; Try shfl.down with floating-point params.
-; CHECK-LABEL: .func{{.*}}shfl.down.float
-define float @shfl.down.float(float %in) {
+; CHECK-LABEL: .func{{.*}}shfl_down_float
+define float @shfl_down_float(float %in) {
   ; CHECK: ld.param.f32 [[IN:%f[0-9]+]]
   ; CHECK: shfl.down.b32 [[OUT:%f[0-9]+]], [[IN]], 5, 6;
   ; CHECK: st.param.{{.}}32 {{.*}}, [[OUT]]
   %out = call float @llvm.nvvm.shfl.down.f32(float %in, i32 5, i32 6)
   ret float %out
 }
 
 ; Try the rest of the shfl modes.  Hopefully they're declared in such a way
 ; that if shfl.down works correctly, they also work correctly.
-define void @shfl.rest(i32 %in_i32, float %in_float, i32* %out_i32, float* %out_float) {
+define void @shfl_rest(i32 %in_i32, float %in_float, i32* %out_i32, float* %out_float) {
   ; CHECK: shfl.up.b32 %r{{[0-9]+}}, %r{{[0-9]+}}, 1, 2;
   %up_i32 = call i32 @llvm.nvvm.shfl.up.i32(i32 %in_i32, i32 1, i32 2)
   store i32 %up_i32, i32* %out_i32
 
   ; CHECK: shfl.up.b32 %f{{[0-9]+}}, %f{{[0-9]+}}, 3, 4;
   %up_float = call float @llvm.nvvm.shfl.up.f32(float %in_float, i32 3, i32 4)
   store float %up_float, float* %out_float
 

llvm/test/CodeGen/NVPTX/vote.ll

 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_30 -mattr=+ptx60 | FileCheck %s
 
 declare i1 @llvm.nvvm.vote.all(i1)
-; CHECK-LABEL: .func{{.*}}vote.all
-define i1 @vote.all(i1 %pred) {
+; CHECK-LABEL: .func{{.*}}vote_all
+define i1 @vote_all(i1 %pred) {
   ; CHECK: vote.all.pred
   %val = call i1 @llvm.nvvm.vote.all(i1 %pred)
   ret i1 %val
 }
 
 declare i1 @llvm.nvvm.vote.any(i1)
-; CHECK-LABEL: .func{{.*}}vote.any
-define i1 @vote.any(i1 %pred) {
+; CHECK-LABEL: .func{{.*}}vote_any
+define i1 @vote_any(i1 %pred) {
   ; CHECK: vote.any.pred
   %val = call i1 @llvm.nvvm.vote.any(i1 %pred)
   ret i1 %val
 }
 
 declare i1 @llvm.nvvm.vote.uni(i1)
-; CHECK-LABEL: .func{{.*}}vote.uni
-define i1 @vote.uni(i1 %pred) {
+; CHECK-LABEL: .func{{.*}}vote_uni
+define i1 @vote_uni(i1 %pred) {
   ; CHECK: vote.uni.pred
   %val = call i1 @llvm.nvvm.vote.uni(i1 %pred)
   ret i1 %val
 }
 
 declare i32 @llvm.nvvm.vote.ballot(i1)
-; CHECK-LABEL: .func{{.*}}vote.ballot
-define i32 @vote.ballot(i1 %pred) {
+; CHECK-LABEL: .func{{.*}}vote_ballot
+define i32 @vote_ballot(i1 %pred) {
   ; CHECK: vote.ballot.b32
   %val = call i32 @llvm.nvvm.vote.ballot(i1 %pred)
   ret i32 %val
 }
 
 declare i1 @llvm.nvvm.vote.all.sync(i32, i1)
-; CHECK-LABEL: .func{{.*}}vote.sync.all
-define i1 @vote.sync.all(i32 %mask, i1 %pred) {
+; CHECK-LABEL: .func{{.*}}vote_sync_all
+define i1 @vote_sync_all(i32 %mask, i1 %pred) {
   ; CHECK: vote.sync.all.pred
   %val = call i1 @llvm.nvvm.vote.all.sync(i32 %mask, i1 %pred)
   ret i1 %val
 }
 
 declare i1 @llvm.nvvm.vote.any.sync(i32, i1)
-; CHECK-LABEL: .func{{.*}}vote.sync.any
-define i1 @vote.sync.any(i32 %mask, i1 %pred) {
+; CHECK-LABEL: .func{{.*}}vote_sync_any
+define i1 @vote_sync_any(i32 %mask, i1 %pred) {
   ; CHECK: vote.sync.any.pred
   %val = call i1 @llvm.nvvm.vote.any.sync(i32 %mask, i1 %pred)
   ret i1 %val
 }
 
 declare i1 @llvm.nvvm.vote.uni.sync(i32, i1)
-; CHECK-LABEL: .func{{.*}}vote.sync.uni
-define i1 @vote.sync.uni(i32 %mask, i1 %pred) {
+; CHECK-LABEL: .func{{.*}}vote_sync_uni
+define i1 @vote_sync_uni(i32 %mask, i1 %pred) {
   ; CHECK: vote.sync.uni.pred
   %val = call i1 @llvm.nvvm.vote.uni.sync(i32 %mask, i1 %pred)
   ret i1 %val
 }
 
 declare i32 @llvm.nvvm.vote.ballot.sync(i32, i1)
-; CHECK-LABEL: .func{{.*}}vote.sync.ballot
-define i32 @vote.sync.ballot(i32 %mask, i1 %pred) {
+; CHECK-LABEL: .func{{.*}}vote_sync_ballot
+define i32 @vote_sync_ballot(i32 %mask, i1 %pred) {
   ; CHECK: vote.sync.ballot.b32
   %val = call i32 @llvm.nvvm.vote.ballot.sync(i32 %mask, i1 %pred)
   ret i32 %val
 }