diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td index 0b183fc30068e..ea58985cbebda 100644 --- a/llvm/include/llvm/IR/IntrinsicsNVVM.td +++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td @@ -5274,7 +5274,7 @@ foreach shape = ["16x64b", "16x128b", "16x256b", "32x32b", "16x32bx2"] in { // def int_nvvm_st_bulk: DefaultAttrsIntrinsic<[], - [llvm_global_ptr_ty, llvm_i64_ty, llvm_i64_ty], + [llvm_ptr_ty, llvm_i64_ty, llvm_i64_ty], [IntrArgMemOnly, IntrWriteMem, WriteOnly>, NoCapture>, ImmArg>]>; diff --git a/llvm/test/CodeGen/NVPTX/st_bulk.ll b/llvm/test/CodeGen/NVPTX/st_bulk.ll index 085df7f1d8d3f..785f78a6f9519 100644 --- a/llvm/test/CodeGen/NVPTX/st_bulk.ll +++ b/llvm/test/CodeGen/NVPTX/st_bulk.ll @@ -4,8 +4,8 @@ ; RUN: %if ptxas-12.8 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx86 | %ptxas-verify -arch=sm_100 %} ; RUN: %if ptxas-12.8 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx86 --nvptx-short-ptr | %ptxas-verify -arch=sm_100 %} -declare void @llvm.nvvm.st.bulk(ptr addrspace(1), i64, i64) -define void @st_bulk(ptr addrspace(1) %dest_addr, i64 %size) { +declare void @llvm.nvvm.st.bulk(ptr, i64, i64) +define void @st_bulk(ptr %dest_addr, i64 %size) { ; CHECK-LABEL: st_bulk( ; CHECK: { ; CHECK-NEXT: .reg .b64 %rd<3>; @@ -15,7 +15,7 @@ define void @st_bulk(ptr addrspace(1) %dest_addr, i64 %size) { ; CHECK-NEXT: ld.param.u64 %rd2, [st_bulk_param_1]; ; CHECK-NEXT: st.bulk [%rd1], %rd2, 0; ; CHECK-NEXT: ret; - call void @llvm.nvvm.st.bulk(ptr addrspace(1) %dest_addr, i64 %size, i64 0) + call void @llvm.nvvm.st.bulk(ptr %dest_addr, i64 %size, i64 0) ret void }