tra created this revision.
Herald added subscribers: mattd, gchakrabarti, asavonic, bixia, hiraditya.
Herald added a project: All.
tra updated this revision to Diff 523216.
tra added a comment.
tra retitled this revision from "[NVPTX] added src_size argument to
__nvvm_cp_async* intrinsics." to "[NVPTX, CUDA] added optional src_size
argument to __nvvm_cp_async*".
tra edited the summary of this revision.
Herald added a subscriber: yaxunl.
tra published this revision for review.
tra added reviewers: jlebar, nyalloc.
Herald added subscribers: llvm-commits, cfe-commits, jdoerfert, jholewinski.
Herald added projects: clang, LLVM.
Updated clang side.
The optional argument is needed for CUDA-11+ headers when we're compiling for
sm_80+ GPUs.
For the intrinsics, the src_size argument is required now. Old calls w/o the
src_size argument can be upgraded by using src_size=transfer size of the
intrinsic.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D150820
Files:
clang/include/clang/Basic/BuiltinsNVPTX.def
clang/include/clang/Sema/Sema.h
clang/lib/CodeGen/CGBuiltin.cpp
clang/lib/Sema/SemaChecking.cpp
clang/test/CodeGen/builtins-nvptx.c
llvm/include/llvm/IR/IntrinsicsNVVM.td
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
llvm/test/CodeGen/NVPTX/async-copy.ll
Index: llvm/test/CodeGen/NVPTX/async-copy.ll
===================================================================
--- llvm/test/CodeGen/NVPTX/async-copy.ll
+++ llvm/test/CodeGen/NVPTX/async-copy.ll
@@ -1,35 +1,35 @@
-; RUN: llc < %s -march=nvptx -mcpu=sm_80 -mattr=+ptx70 | FileCheck -check-prefixes=ALL,CHECK_PTX32 %s
-; RUN: llc < %s -march=nvptx64 -mcpu=sm_80 -mattr=+ptx70 | FileCheck -check-prefixes=ALL,CHECK_PTX64 %s
+; RUN: llc < %s -march=nvptx -mcpu=sm_80 -mattr=+ptx70 | FileCheck -check-prefixes=CHECK,CHECK_PTX32 %s
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_80 -mattr=+ptx70 | FileCheck -check-prefixes=CHECK,CHECK_PTX64 %s
; RUN: %if ptxas-11.0 %{ llc < %s -march=nvptx -mcpu=sm_80 -mattr=+ptx70 | %ptxas-verify -arch=sm_80 %}
; RUN: %if ptxas-11.0 %{ llc < %s -march=nvptx64 -mcpu=sm_80 -mattr=+ptx70 | %ptxas-verify -arch=sm_80 %}
declare void @llvm.nvvm.cp.async.wait.group(i32)
-; ALL-LABEL: asyncwaitgroup
+; CHECK-LABEL: asyncwaitgroup
define void @asyncwaitgroup() {
- ; ALL: cp.async.wait_group 8;
+ ; CHECK: cp.async.wait_group 8;
tail call void @llvm.nvvm.cp.async.wait.group(i32 8)
- ; ALL: cp.async.wait_group 0;
+ ; CHECK: cp.async.wait_group 0;
tail call void @llvm.nvvm.cp.async.wait.group(i32 0)
- ; ALL: cp.async.wait_group 16;
+ ; CHECK: cp.async.wait_group 16;
tail call void @llvm.nvvm.cp.async.wait.group(i32 16)
ret void
}
declare void @llvm.nvvm.cp.async.wait.all()
-; ALL-LABEL: asyncwaitall
+; CHECK-LABEL: asyncwaitall
define void @asyncwaitall() {
-; ALL: cp.async.wait_all
+; CHECK: cp.async.wait_all
tail call void @llvm.nvvm.cp.async.wait.all()
ret void
}
declare void @llvm.nvvm.cp.async.commit.group()
-; ALL-LABEL: asynccommitgroup
+; CHECK-LABEL: asynccommitgroup
define void @asynccommitgroup() {
-; ALL: cp.async.commit_group
+; CHECK: cp.async.commit_group
tail call void @llvm.nvvm.cp.async.commit.group()
ret void
}
@@ -41,72 +41,75 @@
; CHECK-LABEL: asyncmbarrier
define void @asyncmbarrier(ptr %a) {
-; CHECK_PTX32: cp.async.mbarrier.arrive.b64 [%r{{[0-9]+}}];
-; CHECK_PTX64: cp.async.mbarrier.arrive.b64 [%rd{{[0-9]+}}];
+; The distinction between PTX32/PTX64 here is only to capture pointer register type
+; in R to be used in subsequent tests.
+; CHECK_PTX32: cp.async.mbarrier.arrive.b64 [%[[R:r]]{{[0-9]+}}];
+; CHECK_PTX64: cp.async.mbarrier.arrive.b64 [%[[R:rd]]{{[0-9]+}}];
tail call void @llvm.nvvm.cp.async.mbarrier.arrive(ptr %a)
ret void
}
; CHECK-LABEL: asyncmbarriershared
define void @asyncmbarriershared(ptr addrspace(3) %a) {
-; CHECK_PTX32: cp.async.mbarrier.arrive.shared.b64 [%r{{[0-9]+}}];
-; CHECK_PTX64: cp.async.mbarrier.arrive.shared.b64 [%rd{{[0-9]+}}];
+; CHECK: cp.async.mbarrier.arrive.shared.b64 [%[[R]]{{[0-9]+}}];
tail call void @llvm.nvvm.cp.async.mbarrier.arrive.shared(ptr addrspace(3) %a)
ret void
}
; CHECK-LABEL: asyncmbarriernoinc
define void @asyncmbarriernoinc(ptr %a) {
-; CHECK_PTX32: cp.async.mbarrier.arrive.noinc.b64 [%r{{[0-9]+}}];
-; CHECK_PTX64: cp.async.mbarrier.arrive.noinc.b64 [%rd{{[0-9]+}}];
+; CHECK_PTX64: cp.async.mbarrier.arrive.noinc.b64 [%[[R]]{{[0-9]+}}];
tail call void @llvm.nvvm.cp.async.mbarrier.arrive.noinc(ptr %a)
ret void
}
; CHECK-LABEL: asyncmbarriernoincshared
define void @asyncmbarriernoincshared(ptr addrspace(3) %a) {
-; CHECK_PTX32: cp.async.mbarrier.arrive.noinc.shared.b64 [%r{{[0-9]+}}];
-; CHECK_PTX64: cp.async.mbarrier.arrive.noinc.shared.b64 [%rd{{[0-9]+}}];
+; CHECK: cp.async.mbarrier.arrive.noinc.shared.b64 [%[[R]]{{[0-9]+}}];
tail call void @llvm.nvvm.cp.async.mbarrier.arrive.noinc.shared(ptr addrspace(3) %a)
ret void
}
-declare void @llvm.nvvm.cp.async.ca.shared.global.4(ptr addrspace(3) %a, ptr addrspace(1) %b)
+declare void @llvm.nvvm.cp.async.ca.shared.global.4(ptr addrspace(3) %a, ptr addrspace(1) %b, i32 %c)
; CHECK-LABEL: asynccasharedglobal4i8
-define void @asynccasharedglobal4i8(ptr addrspace(3) %a, ptr addrspace(1) %b) {
-; CHECK_PTX32: cp.async.ca.shared.global [%r{{[0-9]+}}], [%r{{[0-9]+}}], 4;
-; CHECK_PTX64: cp.async.ca.shared.global [%rd{{[0-9]+}}], [%rd{{[0-9]+}}], 4;
- tail call void @llvm.nvvm.cp.async.ca.shared.global.4(ptr addrspace(3) %a, ptr addrspace(1) %b)
+define void @asynccasharedglobal4i8(ptr addrspace(3) %a, ptr addrspace(1) %b, i32 %c) {
+; CHECK: cp.async.ca.shared.global [%[[R]]{{[0-9]+}}], [%[[R]]{{[0-9]+}}], 4, %r{{[0-9]+}};
+; CHECK: cp.async.ca.shared.global [%[[R]]{{[0-9]+}}], [%[[R]]{{[0-9]+}}], 4, 1;
+ tail call void @llvm.nvvm.cp.async.ca.shared.global.4(ptr addrspace(3) %a, ptr addrspace(1) %b, i32 %c)
+ tail call void @llvm.nvvm.cp.async.ca.shared.global.4(ptr addrspace(3) %a, ptr addrspace(1) %b, i32 1)
ret void
}
-declare void @llvm.nvvm.cp.async.ca.shared.global.8(ptr addrspace(3) %a, ptr addrspace(1) %b)
+declare void @llvm.nvvm.cp.async.ca.shared.global.8(ptr addrspace(3) %a, ptr addrspace(1) %b, i32 %c)
; CHECK-LABEL: asynccasharedglobal8i8
-define void @asynccasharedglobal8i8(ptr addrspace(3) %a, ptr addrspace(1) %b) {
-; CHECK_PTX32: cp.async.ca.shared.global [%r{{[0-9]+}}], [%r{{[0-9]+}}], 8;
-; CHECK_PTX64: cp.async.ca.shared.global [%rd{{[0-9]+}}], [%rd{{[0-9]+}}], 8;
- tail call void @llvm.nvvm.cp.async.ca.shared.global.8(ptr addrspace(3) %a, ptr addrspace(1) %b)
+define void @asynccasharedglobal8i8(ptr addrspace(3) %a, ptr addrspace(1) %b, i32 %c) {
+; CHECK: cp.async.ca.shared.global [%[[R]]{{[0-9]+}}], [%[[R]]{{[0-9]+}}], 8, %r{{[0-9]+}};
+; CHECK: cp.async.ca.shared.global [%[[R]]{{[0-9]+}}], [%[[R]]{{[0-9]+}}], 8, 1;
+ tail call void @llvm.nvvm.cp.async.ca.shared.global.8(ptr addrspace(3) %a, ptr addrspace(1) %b, i32 %c)
+ tail call void @llvm.nvvm.cp.async.ca.shared.global.8(ptr addrspace(3) %a, ptr addrspace(1) %b, i32 1)
ret void
}
-declare void @llvm.nvvm.cp.async.ca.shared.global.16(ptr addrspace(3) %a, ptr addrspace(1) %b)
+declare void @llvm.nvvm.cp.async.ca.shared.global.16(ptr addrspace(3) %a, ptr addrspace(1) %b, i32 %c)
; CHECK-LABEL: asynccasharedglobal16i8
-define void @asynccasharedglobal16i8(ptr addrspace(3) %a, ptr addrspace(1) %b) {
-; CHECK_PTX32: cp.async.ca.shared.global [%r{{[0-9]+}}], [%r{{[0-9]+}}], 16;
-; CHECK_PTX64: cp.async.ca.shared.global [%rd{{[0-9]+}}], [%rd{{[0-9]+}}], 16;
- tail call void @llvm.nvvm.cp.async.ca.shared.global.16(ptr addrspace(3) %a, ptr addrspace(1) %b)
+define void @asynccasharedglobal16i8(ptr addrspace(3) %a, ptr addrspace(1) %b, i32 %c) {
+; CHECK: cp.async.ca.shared.global [%[[R]]{{[0-9]+}}], [%[[R]]{{[0-9]+}}], 16, %r{{[0-9]+}};
+; CHECK: cp.async.ca.shared.global [%[[R]]{{[0-9]+}}], [%[[R]]{{[0-9]+}}], 16, 1;
+ tail call void @llvm.nvvm.cp.async.ca.shared.global.16(ptr addrspace(3) %a, ptr addrspace(1) %b, i32 %c)
+ tail call void @llvm.nvvm.cp.async.ca.shared.global.16(ptr addrspace(3) %a, ptr addrspace(1) %b, i32 1)
ret void
}
-declare void @llvm.nvvm.cp.async.cg.shared.global.16(ptr addrspace(3) %a, ptr addrspace(1) %b)
+declare void @llvm.nvvm.cp.async.cg.shared.global.16(ptr addrspace(3) %a, ptr addrspace(1) %b, i32 %c)
; CHECK-LABEL: asynccgsharedglobal16i8
-define void @asynccgsharedglobal16i8(ptr addrspace(3) %a, ptr addrspace(1) %b) {
-; CHECK_PTX32: cp.async.cg.shared.global [%r{{[0-9]+}}], [%r{{[0-9]+}}], 16;
-; CHECK_PTX64: cp.async.cg.shared.global [%rd{{[0-9]+}}], [%rd{{[0-9]+}}], 16;
- tail call void @llvm.nvvm.cp.async.cg.shared.global.16(ptr addrspace(3) %a, ptr addrspace(1) %b)
+define void @asynccgsharedglobal16i8(ptr addrspace(3) %a, ptr addrspace(1) %b, i32 %c) {
+; CHECK: cp.async.cg.shared.global [%[[R]]{{[0-9]+}}], [%[[R]]{{[0-9]+}}], 16, %r{{[0-9]+}};
+; CHECK: cp.async.cg.shared.global [%[[R]]{{[0-9]+}}], [%[[R]]{{[0-9]+}}], 16, 1;
+ tail call void @llvm.nvvm.cp.async.cg.shared.global.16(ptr addrspace(3) %a, ptr addrspace(1) %b, i32 %c)
+ tail call void @llvm.nvvm.cp.async.cg.shared.global.16(ptr addrspace(3) %a, ptr addrspace(1) %b, i32 1)
ret void
}
Index: llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
===================================================================
--- llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -328,39 +328,36 @@
defm CP_ASYNC_MBARRIER_ARRIVE_NOINC_SHARED :
CP_ASYNC_MBARRIER_ARRIVE<".noinc", ".shared", int_nvvm_cp_async_mbarrier_arrive_noinc_shared>;
-multiclass CP_ASYNC_CA_SHARED_GLOBAL_I<string cpsize, Intrinsic Intrin> {
- def _32 : NVPTXInst<(outs), (ins Int32Regs:$dst, Int32Regs:$src),
- !strconcat("cp.async.ca.shared.global [$dst], [$src], ", cpsize, ";"),
- [(Intrin Int32Regs:$dst, Int32Regs:$src)]>,
+multiclass CP_ASYNC_SHARED_GLOBAL_I<string cc, string cpsize, Intrinsic Intrin> {
+ def _32 : NVPTXInst<(outs), (ins Int32Regs:$dst, Int32Regs:$src, Int32Regs:$src_size),
+ !strconcat("cp.async.", cc, ".shared.global [$dst], [$src], ", cpsize, ", $src_size;"),
+ [(Intrin Int32Regs:$dst, Int32Regs:$src, Int32Regs:$src_size)]>,
Requires<[hasPTX70, hasSM80]>;
- def _64 : NVPTXInst<(outs), (ins Int64Regs:$dst, Int64Regs:$src),
- !strconcat("cp.async.ca.shared.global [$dst], [$src], ", cpsize, ";"),
- [(Intrin Int64Regs:$dst, Int64Regs:$src)]>,
+ def _32i: NVPTXInst<(outs), (ins Int32Regs:$dst, Int32Regs:$src, i32imm:$src_size),
+ !strconcat("cp.async.", cc, ".shared.global [$dst], [$src], ", cpsize, ", $src_size;"),
+ [(Intrin Int32Regs:$dst, Int32Regs:$src, imm:$src_size)]>,
+ Requires<[hasPTX70, hasSM80]>;
+ def _64 : NVPTXInst<(outs), (ins Int64Regs:$dst, Int64Regs:$src, Int32Regs:$src_size),
+ !strconcat("cp.async.", cc, ".shared.global [$dst], [$src], ", cpsize, ", $src_size;"),
+ [(Intrin Int64Regs:$dst, Int64Regs:$src, Int32Regs:$src_size)]>,
+ Requires<[hasPTX70, hasSM80]>;
+ def _64i: NVPTXInst<(outs), (ins Int64Regs:$dst, Int64Regs:$src, i32imm:$src_size),
+ !strconcat("cp.async.", cc, ".shared.global [$dst], [$src], ", cpsize, ", $src_size;"),
+ [(Intrin Int64Regs:$dst, Int64Regs:$src, imm:$src_size)]>,
Requires<[hasPTX70, hasSM80]>;
}
defm CP_ASYNC_CA_SHARED_GLOBAL_4 :
- CP_ASYNC_CA_SHARED_GLOBAL_I<"4", int_nvvm_cp_async_ca_shared_global_4>;
+ CP_ASYNC_SHARED_GLOBAL_I<"ca", "4", int_nvvm_cp_async_ca_shared_global_4>;
defm CP_ASYNC_CA_SHARED_GLOBAL_8 :
- CP_ASYNC_CA_SHARED_GLOBAL_I<"8", int_nvvm_cp_async_ca_shared_global_8>;
+ CP_ASYNC_SHARED_GLOBAL_I<"ca", "8", int_nvvm_cp_async_ca_shared_global_8>;
defm CP_ASYNC_CA_SHARED_GLOBAL_16 :
- CP_ASYNC_CA_SHARED_GLOBAL_I<"16", int_nvvm_cp_async_ca_shared_global_16>;
-
-multiclass CP_ASYNC_CG_SHARED_GLOBAL<string cpsize, Intrinsic Intrin> {
- def _32 : NVPTXInst<(outs), (ins Int32Regs:$dst, Int32Regs:$src),
- !strconcat("cp.async.cg.shared.global [$dst], [$src], ", cpsize, ";"),
- [(Intrin Int32Regs:$dst, Int32Regs:$src)]>,
- Requires<[hasPTX70, hasSM80]>;
- def _64 : NVPTXInst<(outs), (ins Int64Regs:$dst, Int64Regs:$src),
- !strconcat("cp.async.cg.shared.global [$dst], [$src], ", cpsize, ";"),
- [(Intrin Int64Regs:$dst, Int64Regs:$src)]>,
- Requires<[hasPTX70, hasSM80]>;
-}
+ CP_ASYNC_SHARED_GLOBAL_I<"ca", "16", int_nvvm_cp_async_ca_shared_global_16>;
defm CP_ASYNC_CG_SHARED_GLOBAL_16 :
- CP_ASYNC_CG_SHARED_GLOBAL<"16", int_nvvm_cp_async_cg_shared_global_16>;
+ CP_ASYNC_SHARED_GLOBAL_I<"cg", "16", int_nvvm_cp_async_cg_shared_global_16>;
def CP_ASYNC_COMMIT_GROUP :
NVPTXInst<(outs), (ins), "cp.async.commit_group;", [(int_nvvm_cp_async_commit_group)]>,
Index: llvm/include/llvm/IR/IntrinsicsNVVM.td
===================================================================
--- llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -1381,26 +1381,22 @@
Intrinsic<[],[llvm_shared_i64ptr_ty],[IntrConvergent, IntrNoCallback]>;
def int_nvvm_cp_async_ca_shared_global_4 :
- ClangBuiltin<"__nvvm_cp_async_ca_shared_global_4">,
- Intrinsic<[],[llvm_shared_i8ptr_ty, llvm_global_i8ptr_ty],
+ Intrinsic<[],[llvm_shared_i8ptr_ty, llvm_global_i8ptr_ty, llvm_i32_ty],
[IntrArgMemOnly, IntrNoCallback, NoAlias<ArgIndex<0>>, NoAlias<ArgIndex<1>>,
WriteOnly<ArgIndex<0>>, ReadOnly<ArgIndex<1>>],
"llvm.nvvm.cp.async.ca.shared.global.4">;
def int_nvvm_cp_async_ca_shared_global_8 :
- ClangBuiltin<"__nvvm_cp_async_ca_shared_global_8">,
- Intrinsic<[],[llvm_shared_i8ptr_ty, llvm_global_i8ptr_ty],
+ Intrinsic<[],[llvm_shared_i8ptr_ty, llvm_global_i8ptr_ty, llvm_i32_ty],
[IntrArgMemOnly, IntrNoCallback, NoAlias<ArgIndex<0>>, NoAlias<ArgIndex<1>>,
WriteOnly<ArgIndex<0>>, ReadOnly<ArgIndex<1>>],
"llvm.nvvm.cp.async.ca.shared.global.8">;
def int_nvvm_cp_async_ca_shared_global_16 :
- ClangBuiltin<"__nvvm_cp_async_ca_shared_global_16">,
- Intrinsic<[],[llvm_shared_i8ptr_ty, llvm_global_i8ptr_ty],
+ Intrinsic<[],[llvm_shared_i8ptr_ty, llvm_global_i8ptr_ty, llvm_i32_ty],
[IntrArgMemOnly, IntrNoCallback, NoAlias<ArgIndex<0>>, NoAlias<ArgIndex<1>>,
WriteOnly<ArgIndex<0>>, ReadOnly<ArgIndex<1>>],
"llvm.nvvm.cp.async.ca.shared.global.16">;
def int_nvvm_cp_async_cg_shared_global_16 :
- ClangBuiltin<"__nvvm_cp_async_cg_shared_global_16">,
- Intrinsic<[],[llvm_shared_i8ptr_ty, llvm_global_i8ptr_ty],
+ Intrinsic<[],[llvm_shared_i8ptr_ty, llvm_global_i8ptr_ty, llvm_i32_ty],
[IntrArgMemOnly, IntrNoCallback, NoAlias<ArgIndex<0>>, NoAlias<ArgIndex<1>>,
WriteOnly<ArgIndex<0>>, ReadOnly<ArgIndex<1>>],
"llvm.nvvm.cp.async.cg.shared.global.16">;
Index: clang/test/CodeGen/builtins-nvptx.c
===================================================================
--- clang/test/CodeGen/builtins-nvptx.c
+++ clang/test/CodeGen/builtins-nvptx.c
@@ -830,15 +830,24 @@
// CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.mbarrier.arrive.noinc.shared
__nvvm_cp_async_mbarrier_arrive_noinc_shared(sharedAddr);
- // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.ca.shared.global.4
+ // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.ca.shared.global.4({{.*}}, i32 4)
__nvvm_cp_async_ca_shared_global_4(dst, src);
- // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.ca.shared.global.8
+ // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.ca.shared.global.8({{.*}}, i32 8)
__nvvm_cp_async_ca_shared_global_8(dst, src);
- // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.ca.shared.global.16
+ // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.ca.shared.global.16({{.*}}, i32 16)
__nvvm_cp_async_ca_shared_global_16(dst, src);
- // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.cg.shared.global.16
+ // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.cg.shared.global.16({{.*}}, i32 16)
__nvvm_cp_async_cg_shared_global_16(dst, src);
+ // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.ca.shared.global.4({{.*}}, i32 2)
+ __nvvm_cp_async_ca_shared_global_4(dst, src, 2);
+ // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.ca.shared.global.8({{.*}}, i32 2)
+ __nvvm_cp_async_ca_shared_global_8(dst, src, 2);
+ // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.ca.shared.global.16({{.*}}, i32 2)
+ __nvvm_cp_async_ca_shared_global_16(dst, src, 2);
+ // CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.cg.shared.global.16({{.*}}, i32 2)
+ __nvvm_cp_async_cg_shared_global_16(dst, src, 2);
+
// CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.commit.group
__nvvm_cp_async_commit_group();
// CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.wait.group(i32 0)
Index: clang/lib/Sema/SemaChecking.cpp
===================================================================
--- clang/lib/Sema/SemaChecking.cpp
+++ clang/lib/Sema/SemaChecking.cpp
@@ -4815,6 +4815,20 @@
return false;
}
+bool Sema::CheckNVPTXBuiltinFunctionCall(const TargetInfo &TI,
+ unsigned BuiltinID,
+ CallExpr *TheCall) {
+ switch (BuiltinID) {
+ case NVPTX::BI__nvvm_cp_async_ca_shared_global_4:
+ case NVPTX::BI__nvvm_cp_async_ca_shared_global_8:
+ case NVPTX::BI__nvvm_cp_async_ca_shared_global_16:
+ case NVPTX::BI__nvvm_cp_async_cg_shared_global_16:
+ return checkArgCountAtMost(*this, TheCall, 3);
+ }
+
+ return false;
+}
+
/// SemaBuiltinCpuSupports - Handle __builtin_cpu_supports(char *).
/// This checks that the target supports __builtin_cpu_supports and
/// that the string argument is constant and valid.
Index: clang/lib/CodeGen/CGBuiltin.cpp
===================================================================
--- clang/lib/CodeGen/CGBuiltin.cpp
+++ clang/lib/CodeGen/CGBuiltin.cpp
@@ -18177,6 +18177,15 @@
{Ptr, CGF.EmitScalarExpr(E->getArg(1))});
}
+static Value *MakeCpAsync(unsigned IntrinsicID, CodeGenFunction &CGF,
+ const CallExpr *E, int SrcSize) {
+ Value *SrcSizeArg = E->getNumArgs() == 3 ? CGF.EmitScalarExpr(E->getArg(2))
+ : CGF.Builder.getInt32(SrcSize);
+ return CGF.Builder.CreateCall(CGF.CGM.getIntrinsic(IntrinsicID),
+ {CGF.EmitScalarExpr(E->getArg(0)),
+ CGF.EmitScalarExpr(E->getArg(1)), SrcSizeArg});
+}
+
static Value *MakeHalfType(unsigned IntrinsicID, unsigned BuiltinID,
const CallExpr *E, CodeGenFunction &CGF) {
auto &C = CGF.CGM.getContext();
@@ -18840,6 +18849,18 @@
case NVPTX::BI__nvvm_ldu_h2: {
return MakeHalfType(Intrinsic::nvvm_ldu_global_f, BuiltinID, E, *this);
}
+ case NVPTX::BI__nvvm_cp_async_ca_shared_global_4:
+ return MakeCpAsync(Intrinsic::nvvm_cp_async_ca_shared_global_4, *this, E,
+ 4);
+ case NVPTX::BI__nvvm_cp_async_ca_shared_global_8:
+ return MakeCpAsync(Intrinsic::nvvm_cp_async_ca_shared_global_8, *this, E,
+ 8);
+ case NVPTX::BI__nvvm_cp_async_ca_shared_global_16:
+ return MakeCpAsync(Intrinsic::nvvm_cp_async_ca_shared_global_16, *this, E,
+ 16);
+ case NVPTX::BI__nvvm_cp_async_cg_shared_global_16:
+ return MakeCpAsync(Intrinsic::nvvm_cp_async_cg_shared_global_16, *this, E,
+ 16);
default:
return nullptr;
}
Index: clang/include/clang/Sema/Sema.h
===================================================================
--- clang/include/clang/Sema/Sema.h
+++ clang/include/clang/Sema/Sema.h
@@ -13558,6 +13558,8 @@
bool CheckWebAssemblyBuiltinFunctionCall(const TargetInfo &TI,
unsigned BuiltinID,
CallExpr *TheCall);
+ bool CheckNVPTXBuiltinFunctionCall(const TargetInfo &TI, unsigned BuiltinID,
+ CallExpr *TheCall);
bool SemaBuiltinVAStart(unsigned BuiltinID, CallExpr *TheCall);
bool SemaBuiltinVAStartARMMicrosoft(CallExpr *Call);
Index: clang/include/clang/Basic/BuiltinsNVPTX.def
===================================================================
--- clang/include/clang/Basic/BuiltinsNVPTX.def
+++ clang/include/clang/Basic/BuiltinsNVPTX.def
@@ -968,10 +968,10 @@
TARGET_BUILTIN(__nvvm_cp_async_mbarrier_arrive_noinc, "vWi*", "", AND(SM_80,PTX70))
TARGET_BUILTIN(__nvvm_cp_async_mbarrier_arrive_noinc_shared, "vWi*3", "", AND(SM_80,PTX70))
-TARGET_BUILTIN(__nvvm_cp_async_ca_shared_global_4, "vv*3vC*1", "", AND(SM_80,PTX70))
-TARGET_BUILTIN(__nvvm_cp_async_ca_shared_global_8, "vv*3vC*1", "", AND(SM_80,PTX70))
-TARGET_BUILTIN(__nvvm_cp_async_ca_shared_global_16, "vv*3vC*1", "", AND(SM_80,PTX70))
-TARGET_BUILTIN(__nvvm_cp_async_cg_shared_global_16, "vv*3vC*1", "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_cp_async_ca_shared_global_4, "vv*3vC*1.", "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_cp_async_ca_shared_global_8, "vv*3vC*1.", "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_cp_async_ca_shared_global_16, "vv*3vC*1.", "", AND(SM_80,PTX70))
+TARGET_BUILTIN(__nvvm_cp_async_cg_shared_global_16, "vv*3vC*1.", "", AND(SM_80,PTX70))
TARGET_BUILTIN(__nvvm_cp_async_commit_group, "v", "", AND(SM_80,PTX70))
TARGET_BUILTIN(__nvvm_cp_async_wait_group, "vIi", "", AND(SM_80,PTX70))
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
