This is an archive of the discontinued LLVM Phabricator instance.

Differential D100394

[Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX cp.async instructions
ClosedPublic

Authored by nyalloc on Apr 13 2021, 10:03 AM.

Details

Reviewers
tra
jholewinski
jdoerfert
Commits
rG02c2468864bb: [Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX cp.async…
Summary

Adds NVPTX builtins and intrinsics for the CUDA PTX cp.async instructions for sm_80 architecture or newer.

PTX ISA description of cp.async:
https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-asynchronous-copy
https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-cp-async-mbarrier-arrive

Authored-by: Stuart Adams <stuart.adams@codeplay.com>
Co-Authored-by: Alexander Johnston <alexander@codeplay.com>

Diff Detail

Event Timeline

nyalloc created this revision.Apr 13 2021, 10:03 AM
Herald added subscribers: hiraditya, yaxunl. · View Herald TranscriptApr 13 2021, 10:03 AM
nyalloc requested review of this revision.Apr 13 2021, 10:03 AM
Herald added a subscriber: llvm-commits. · View Herald TranscriptApr 13 2021, 10:03 AM
steffenlarsen added a subscriber: steffenlarsen.Apr 13 2021, 10:04 AM
nyalloc edited the summary of this revision. (Show Details)Apr 13 2021, 10:05 AM
nyalloc edited the summary of this revision. (Show Details)
tra added inline comments.Apr 13 2021, 12:22 PM
clang/include/clang/Basic/BuiltinsNVPTX.def
756–759

For cp.async.mbarrier instructions to work we do need to have mbarrier.initto init the barrier object and other mbarrier ops to use them inpractice. Perhaps these should be added if/when all mbarrier instructions are added.

Also, mbarrier object has additional requirements for the pointer (aligned by 8, in shared space): https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier-size-alignment

So, technically they all should use v*3, but I don't think it will work well in practice -- a lot of things assume that we start with all the pointers bein in generic AS. Nevertheless, we do want to have some sort of safeguards for these builtins.

Perhaps it would make sense to add a custom type checker and only allow references to __shared__ variables.

767

I think this should be "vIi" as the instruction only accepts an integer constant as an argument.

llvm/include/llvm/IR/IntrinsicsNVVM.td
1074–1089

These all should probably have IntrArgMemOnly attribute and, possibly WriteOnly<0>,ReadOnly<1> and, maybe NoAlias on both arguments, too, because src/dest are in different nonoverlapping address spaces.

Also, the PTX spec is not clear on whether cp.async expects to see the pointer arguments in generic AS, or do they need to be converted to shared/global ones first. Normally, the instructions with .shared or .global in the name expect specific address space. If that's the case here, then we may need to use qualified pointer types here, too.

1095–1097

This should have ImmArg as the argument must be an immediate value.

llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
353–356

This does not look right. If I read the PTX spec correctly, the argument can't be a register.

https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-wait-group

Operand N is an integer constant.

llvm/test/CodeGen/NVPTX/async-copy.ll
4

No tests for cp.async.wait* and cp.async.commit_group

Harbormaster completed remote builds in B98511: Diff 337192.Apr 13 2021, 4:07 PM
Alexander-Johnston added a subscriber: Alexander-Johnston.Apr 14 2021, 2:41 AM
vchuravy added a subscriber: vchuravy.Apr 14 2021, 6:03 AM
nyalloc updated this revision to Diff 338059.EditedApr 16 2021, 4:05 AM

Updated diff to address review comments.

  • mbarrier intrinsics, builtins and tests are now included
  • alignment / address space updated for cp.async.mbarrier instructions
  • cp_async_wait_all intrinsics and builtins updated to use immediate values
  • copy intrinsics updated with appropriate attributes
  • copy intrinsics / builtins arguments are updated to use the appropriate address spaces
  • missing tests added
Harbormaster completed remote builds in B99138: Diff 338059.Apr 16 2021, 5:03 AM
tra added a subscriber: • wash.Apr 16 2021, 10:55 AM

Overall the patch looks good. We may still need to tweak intrinsic properties later, but this is a good starting point.

I'm not familiar enough with the new instructions, so my suggestions are based on just reading the PTX spec and there's a good change I didn't get it all right.
If someone from NVIDIA is watching, now would be a good chance to chime in.
@wash, @jholewinski - any comments on how these instructions should be handled?

clang/include/clang/Basic/BuiltinsNVPTX.def
470

I think _b64 is redundant for the mbarrier instructions -- that's the only type they accept.

clang/test/CodeGen/builtins-nvptx.c
682

I'd add CHECK-LABEL: <function_name> here and in other functions.

718

Extra //

llvm/include/llvm/IR/IntrinsicsNVVM.td
34–39

llvm_globali8ptr_ty -> llvm_global_i8ptr_ty would make it a bit easier to read.

1112–1122

These are probably safe to mark as IntrWriteMem, IntrArgMemOnly, WriteOnly<ArgIndex<0>>, NoCapture<ArgIndex<0>>

1156–1158

This one can be IntrNoMem.

llvm/test/CodeGen/NVPTX/async-copy.ll
6–8

I'd recommend adding a common check label (le'ts say ALL and running the tests with --check-prefixes=ALL,CHECK_PTX64.

This would allow you to use ALL for things that do not change. In this case ; ALL: cp.async.wait_group 8;
It would also be great to add ALL-LABEL: <function_name> for each function to limit the range FileCheck operates for.

nyalloc updated this revision to Diff 338848.Apr 20 2021, 6:38 AM

Addressed @tra's review comments.

  • _b64 postfix is removed from mbarrier intrinsics and builtins.
  • CHECK-LABEL is introduced in builtins-nvptx.c
  • Code style updated in IntrinsicsNVVM.td: llvm_globali8ptr_ty -> llvm_global_i8ptr_ty etc
  • int_nvvm_mbarrier_inval && int_nvvm_mbarrier_inval_shared marked as IntrWriteMem, IntrArgMemOnly, WriteOnly<ArgIndex<0>>, NoCapture<ArgIndex<0>>
  • int_nvvm_mbarrier_pending_count marked as IntrNoMem
  • ALL check prefix added to async-copy.ll, ALL-LABEL used appropriately
Harbormaster completed remote builds in B99706: Diff 338848.Apr 20 2021, 7:45 AM
tra accepted this revision.Apr 20 2021, 10:02 AM

LGTM overall, modulo few test and naming nits.

llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
415

Does ptxas accept all-lower-case nocomplete?

The no*C*omplete stands out as a sore thumb. Capital letters are used in few LLVM intrinsics, so it's not a showstopper, but I think lower case everywhere makes more sense. WDYT?

llvm/test/CodeGen/NVPTX/async-copy.ll
41

All functions in the file should use -LABEL checks.

llvm/test/CodeGen/NVPTX/mbarrier.ll
8

Same here. Please add -LABEL checks for all functions.

This revision is now accepted and ready to land.Apr 20 2021, 10:02 AM
nyalloc marked 13 inline comments as done.Apr 20 2021, 10:10 AM
nyalloc added inline comments.
clang/include/clang/Basic/BuiltinsNVPTX.def
756–759

Well spotted, I'll add the mbarrier intrinsics and builtins to this patch. They were originally going to be put up separately but I'll add them in to this seen as they are related.

llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
415

I 100% agree. For some reason the PTX decide to use camel case here. They also throw some snake case in other names. We can change the LLVM intrinsics to use a more consistent naming scheme, but it will come at the cost that it's no longer a clean mapping of names to the PTX.

llvm/test/CodeGen/NVPTX/mbarrier.ll
8

Will do!

nyalloc updated this revision to Diff 339145.Apr 21 2021, 2:07 AM
nyalloc marked an inline comment as done.

All functions in the new tests now use -LABEL checks.

Harbormaster completed remote builds in B99917: Diff 339145.Apr 21 2021, 2:59 AM
steffenlarsen mentioned this in D100124: [Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX redux.sync instructions.May 13 2021, 10:07 AM
tra added a comment.May 17 2021, 9:40 AM

I'll land this patch along with D100124

Closed by commit rG02c2468864bb: [Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX cp.async… (authored by nyalloc, committed by tra). · Explain WhyMay 17 2021, 9:47 AM
This revision was automatically updated to reflect the committed changes.
tra added a commit: rG02c2468864bb: [Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX cp.async….
Herald added a subscriber: cfe-commits. · View Herald TranscriptMay 17 2021, 9:47 AM
nirvedhmeshram added a subscriber: nirvedhmeshram.Apr 21 2022, 5:54 PM

Hello, I was interested in using llvm.nvvm.cp.async.cg.shared.global.8 and llvm.nvvm.cp.async.cg.shared.global.4 and was wondering if there is some fundamental reason they were not added here. I only see the ca variants for these.

Herald added a project: Restricted Project. · View Herald TranscriptApr 21 2022, 5:54 PM
Herald added subscribers: mattd, gchakrabarti, asavonic. · View Herald Transcript
steffenlarsen added a comment.Apr 22 2022, 12:36 AM
In D100394#3466316, @nirvedhmeshram wrote:

Hello, I was interested in using llvm.nvvm.cp.async.cg.shared.global.8 and llvm.nvvm.cp.async.cg.shared.global.4 and was wondering if there is some fundamental reason they were not added here. I only see the ca variants for these.

Hi @nirvedhmeshram! According to the PTX ISA there is only a 16 variant of cp.async.cg.shared.global. That said, they have an example further down using 8 with it, so it seems there's either a problem in the Syntax subsection or the examples. Either way, that is the explanation as to why it was not added with this.

tra added a comment.EditedMay 17 2023, 2:05 PM

Hi. It looks like CUDA-11+ headers need a variant of cm.async intrinsics which provides the optional src_size argument.

I'm planning to add it to the existing intrinsics in NVPTX. It's just a heads-up in case you may have existing uses of them that may need to be updated.

See https://reviews.llvm.org/D150820 for the proposed changes.

Revision Contents

PathSize
clang/
include/
clang/
Basic/
38 lines
test/
CodeGen/
81 lines
llvm/
include/
llvm/
IR/
105 lines
lib/
Target/
NVPTX/
2 lines
205 lines
test/
CodeGen/
NVPTX/
101 lines
123 lines

Diff 338059

clang/include/clang/Basic/BuiltinsNVPTX.def

Show First 20 Lines • Show All 459 Lines • ▼ Show 20 Lines
TARGET_BUILTIN(__nvvm_match_all_sync_i64p, "WiUiWii*", "", PTX60) TARGET_BUILTIN(__nvvm_match_all_sync_i64p, "WiUiWii*", "", PTX60)
// Membar // Membar
BUILTIN(__nvvm_membar_cta, "v", "") BUILTIN(__nvvm_membar_cta, "v", "")
BUILTIN(__nvvm_membar_gl, "v", "") BUILTIN(__nvvm_membar_gl, "v", "")
BUILTIN(__nvvm_membar_sys, "v", "") BUILTIN(__nvvm_membar_sys, "v", "")
// mbarrier
TARGET_BUILTIN(__nvvm_mbarrier_init_b64, "vWi*i", "", AND(SM_80,PTX70))
traUnsubmitted
Done
ReplyInline Actions

I think _b64 is redundant for the mbarrier instructions -- that's the only type they accept.

tra: I think _b64 is redundant for the `mbarrier` instructions -- that's the only type they accept.
TARGET_BUILTIN(__nvvm_mbarrier_init_shared_b64, "vWi*3i", "", AND(SM_80,PTX70))
TARGET_BUILTIN(__nvvm_mbarrier_inval_b64, "vWi*", "", AND(SM_80,PTX70))
TARGET_BUILTIN(__nvvm_mbarrier_inval_shared_b64, "vWi*3", "", AND(SM_80,PTX70))
TARGET_BUILTIN(__nvvm_mbarrier_arrive_b64, "WiWi*", "", AND(SM_80,PTX70))
TARGET_BUILTIN(__nvvm_mbarrier_arrive_shared_b64, "WiWi*3", "", AND(SM_80,PTX70))
TARGET_BUILTIN(__nvvm_mbarrier_arrive_noComplete_b64, "WiWi*i", "", AND(SM_80,PTX70))
TARGET_BUILTIN(__nvvm_mbarrier_arrive_noComplete_shared_b64, "WiWi*3i", "", AND(SM_80,PTX70))
TARGET_BUILTIN(__nvvm_mbarrier_arrive_drop_b64, "WiWi*", "", AND(SM_80,PTX70))
TARGET_BUILTIN(__nvvm_mbarrier_arrive_drop_shared_b64, "WiWi*3", "", AND(SM_80,PTX70))
TARGET_BUILTIN(__nvvm_mbarrier_arrive_drop_noComplete_b64, "WiWi*i", "", AND(SM_80,PTX70))
TARGET_BUILTIN(__nvvm_mbarrier_arrive_drop_noComplete_shared_b64, "WiWi*3i", "", AND(SM_80,PTX70))
TARGET_BUILTIN(__nvvm_mbarrier_test_wait_b64, "bWi*Wi", "", AND(SM_80,PTX70))
TARGET_BUILTIN(__nvvm_mbarrier_test_wait_shared_b64, "bWi*3Wi", "", AND(SM_80,PTX70))
TARGET_BUILTIN(__nvvm_mbarrier_pending_count_b64, "iWi", "", AND(SM_80,PTX70))
// Memcpy, Memset // Memcpy, Memset
BUILTIN(__nvvm_memcpy, "vUc*Uc*zi","") BUILTIN(__nvvm_memcpy, "vUc*Uc*zi","")
BUILTIN(__nvvm_memset, "vUc*Uczi","") BUILTIN(__nvvm_memset, "vUc*Uczi","")
// Image // Image
BUILTIN(__builtin_ptx_read_image2Dfi_, "V4fiiii", "") BUILTIN(__builtin_ptx_read_image2Dfi_, "V4fiiii", "")
▲ Show 20 Lines • Show All 248 Lines • ▼ Show 20 Lines
TARGET_BUILTIN(__imma_m8n8k32_ld_a_u4, "vi*iC*UiIi", "", AND(SM_75,PTX63)) TARGET_BUILTIN(__imma_m8n8k32_ld_a_u4, "vi*iC*UiIi", "", AND(SM_75,PTX63))
TARGET_BUILTIN(__imma_m8n8k32_ld_b_s4, "vi*iC*UiIi", "", AND(SM_75,PTX63)) TARGET_BUILTIN(__imma_m8n8k32_ld_b_s4, "vi*iC*UiIi", "", AND(SM_75,PTX63))
TARGET_BUILTIN(__imma_m8n8k32_ld_b_u4, "vi*iC*UiIi", "", AND(SM_75,PTX63)) TARGET_BUILTIN(__imma_m8n8k32_ld_b_u4, "vi*iC*UiIi", "", AND(SM_75,PTX63))
TARGET_BUILTIN(__imma_m8n8k32_ld_c, "vi*iC*UiIi", "", AND(SM_75,PTX63)) TARGET_BUILTIN(__imma_m8n8k32_ld_c, "vi*iC*UiIi", "", AND(SM_75,PTX63))
TARGET_BUILTIN(__imma_m8n8k32_mma_s4, "vi*iC*iC*iC*IiIi", "", AND(SM_75,PTX63)) TARGET_BUILTIN(__imma_m8n8k32_mma_s4, "vi*iC*iC*iC*IiIi", "", AND(SM_75,PTX63))
TARGET_BUILTIN(__imma_m8n8k32_mma_u4, "vi*iC*iC*iC*IiIi", "", AND(SM_75,PTX63)) TARGET_BUILTIN(__imma_m8n8k32_mma_u4, "vi*iC*iC*iC*IiIi", "", AND(SM_75,PTX63))
TARGET_BUILTIN(__imma_m8n8k32_st_c_i32, "vi*iC*UiIi", "", AND(SM_75,PTX63)) TARGET_BUILTIN(__imma_m8n8k32_st_c_i32, "vi*iC*UiIi", "", AND(SM_75,PTX63))
// Async Copy
TARGET_BUILTIN(__nvvm_cp_async_mbarrier_arrive_b64, "vWi*", "", AND(SM_80,PTX70))
TARGET_BUILTIN(__nvvm_cp_async_mbarrier_arrive_shared_b64, "vWi*3", "", AND(SM_80,PTX70))
TARGET_BUILTIN(__nvvm_cp_async_mbarrier_arrive_noinc_b64, "vWi*", "", AND(SM_80,PTX70))
TARGET_BUILTIN(__nvvm_cp_async_mbarrier_arrive_noinc_shared_b64, "vWi*3", "", AND(SM_80,PTX70))
traUnsubmitted
Done
ReplyInline Actions

For cp.async.mbarrier instructions to work we do need to have mbarrier.initto init the barrier object and other mbarrier ops to use them inpractice. Perhaps these should be added if/when all mbarrier instructions are added.

Also, mbarrier object has additional requirements for the pointer (aligned by 8, in shared space): https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier-size-alignment

So, technically they all should use v*3, but I don't think it will work well in practice -- a lot of things assume that we start with all the pointers bein in generic AS. Nevertheless, we do want to have some sort of safeguards for these builtins.

Perhaps it would make sense to add a custom type checker and only allow references to __shared__ variables.

tra: For `cp.async.mbarrier` instructions to work we do need to have `mbarrier.init`to init the…
nyallocAuthorUnsubmitted
Done
ReplyInline Actions

Well spotted, I'll add the mbarrier intrinsics and builtins to this patch. They were originally going to be put up separately but I'll add them in to this seen as they are related.

nyalloc: Well spotted, I'll add the mbarrier intrinsics and builtins to this patch. They were originally…
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))
traUnsubmitted
Done
ReplyInline Actions

I think this should be "vIi" as the instruction only accepts an integer constant as an argument.

tra: I think this should be "vIi" as the instruction only accepts an integer constant as an argument.
TARGET_BUILTIN(__nvvm_cp_async_wait_all, "v", "", AND(SM_80,PTX70))
#undef BUILTIN #undef BUILTIN
#undef TARGET_BUILTIN #undef TARGET_BUILTIN
#pragma pop_macro("AND") #pragma pop_macro("AND")
#pragma pop_macro("SM_60") #pragma pop_macro("SM_60")
#pragma pop_macro("SM_70") #pragma pop_macro("SM_70")
#pragma pop_macro("SM_72") #pragma pop_macro("SM_72")
#pragma pop_macro("SM_75") #pragma pop_macro("SM_75")
#pragma pop_macro("SM_80") #pragma pop_macro("SM_80")
Show All 9 Lines

clang/test/CodeGen/builtins-nvptx.c

// REQUIRES: nvptx-registered-target // REQUIRES: nvptx-registered-target
// RUN: %clang_cc1 -ffp-contract=off -triple nvptx-unknown-unknown -target-cpu sm_80 -target-feature +ptx70 \
// RUN: -fcuda-is-device -S -emit-llvm -o - -x cuda %s \
// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX70_SM80 -check-prefix=LP32 %s
// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_80 -target-feature +ptx70 \
// RUN: -fcuda-is-device -S -emit-llvm -o - -x cuda %s \
// RUN: | FileCheck -check-prefix=CHECK -check-prefix=CHECK_PTX70_SM80 -check-prefix=LP64 %s
// RUN: %clang_cc1 -ffp-contract=off -triple nvptx-unknown-unknown -target-cpu sm_60 \ // RUN: %clang_cc1 -ffp-contract=off -triple nvptx-unknown-unknown -target-cpu sm_60 \
// RUN: -fcuda-is-device -S -emit-llvm -o - -x cuda %s \ // RUN: -fcuda-is-device -S -emit-llvm -o - -x cuda %s \
// RUN: | FileCheck -check-prefix=CHECK -check-prefix=LP32 %s // RUN: | FileCheck -check-prefix=CHECK -check-prefix=LP32 %s
// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_60 \ // RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_60 \
// RUN: -fcuda-is-device -S -emit-llvm -o - -x cuda %s \ // RUN: -fcuda-is-device -S -emit-llvm -o - -x cuda %s \
// RUN: | FileCheck -check-prefix=CHECK -check-prefix=LP64 %s // RUN: | FileCheck -check-prefix=CHECK -check-prefix=LP64 %s
// RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_61 \ // RUN: %clang_cc1 -ffp-contract=off -triple nvptx64-unknown-unknown -target-cpu sm_61 \
// RUN: -fcuda-is-device -S -emit-llvm -o - -x cuda %s \ // RUN: -fcuda-is-device -S -emit-llvm -o - -x cuda %s \
▲ Show 20 Lines • Show All 657 Lines • ▼ Show 20 Lines__device__ void nvvm_vote(int pred) {
// CHECK: call i1 @llvm.nvvm.vote.any(i1 // CHECK: call i1 @llvm.nvvm.vote.any(i1
__nvvm_vote_any(pred); __nvvm_vote_any(pred);
// CHECK: call i1 @llvm.nvvm.vote.uni(i1 // CHECK: call i1 @llvm.nvvm.vote.uni(i1
__nvvm_vote_uni(pred); __nvvm_vote_uni(pred);
// CHECK: call i32 @llvm.nvvm.vote.ballot(i1 // CHECK: call i32 @llvm.nvvm.vote.ballot(i1
__nvvm_vote_ballot(pred); __nvvm_vote_ballot(pred);
// CHECK: ret void // CHECK: ret void
} }
__device__ void nvvm_mbarrier(long long* addr, __attribute__((address_space(3))) long long* sharedAddr, int count, long long state) {
traUnsubmitted
Done
ReplyInline Actions

I'd add CHECK-LABEL: <function_name> here and in other functions.

tra: I'd add `CHECK-LABEL: <function_name>` here and in other functions.
#if __CUDA_ARCH__ >= 800
__nvvm_mbarrier_init_b64(addr, count);
// CHECK_PTX70_SM80: call void @llvm.nvvm.mbarrier.init.b64
__nvvm_mbarrier_init_shared_b64(sharedAddr, count);
// CHECK_PTX70_SM80: call void @llvm.nvvm.mbarrier.init.shared.b64
__nvvm_mbarrier_inval_b64(addr);
// CHECK_PTX70_SM80: call void @llvm.nvvm.mbarrier.inval.b64
__nvvm_mbarrier_inval_shared_b64(sharedAddr);
// CHECK_PTX70_SM80: call void @llvm.nvvm.mbarrier.inval.shared.b64
__nvvm_mbarrier_arrive_b64(addr);
// CHECK_PTX70_SM80: call i64 @llvm.nvvm.mbarrier.arrive.b64
__nvvm_mbarrier_arrive_shared_b64(sharedAddr);
// CHECK_PTX70_SM80: call i64 @llvm.nvvm.mbarrier.arrive.shared.b64
__nvvm_mbarrier_arrive_noComplete_b64(addr, count);
// CHECK_PTX70_SM80: call i64 @llvm.nvvm.mbarrier.arrive.noComplete.b64
__nvvm_mbarrier_arrive_noComplete_shared_b64(sharedAddr, count);
// CHECK_PTX70_SM80: call i64 @llvm.nvvm.mbarrier.arrive.noComplete.shared.b64
__nvvm_mbarrier_arrive_drop_b64(addr);
// CHECK_PTX70_SM80: call i64 @llvm.nvvm.mbarrier.arrive.drop.b64
__nvvm_mbarrier_arrive_drop_shared_b64(sharedAddr);
// CHECK_PTX70_SM80: call i64 @llvm.nvvm.mbarrier.arrive.drop.shared.b64
__nvvm_mbarrier_arrive_drop_noComplete_b64(addr, count);
// CHECK_PTX70_SM80: call i64 @llvm.nvvm.mbarrier.arrive.drop.noComplete.b64
__nvvm_mbarrier_arrive_drop_noComplete_shared_b64(sharedAddr, count);
// CHECK_PTX70_SM80: call i64 @llvm.nvvm.mbarrier.arrive.drop.noComplete.shared.b64
__nvvm_mbarrier_test_wait_b64(addr, state);
// CHECK_PTX70_SM80: call i1 @llvm.nvvm.mbarrier.test.wait.b64
__nvvm_mbarrier_test_wait_shared_b64(sharedAddr, state);
// CHECK_PTX70_SM80: call i1 @llvm.nvvm.mbarrier.test.wait.shared.b64
__nvvm_mbarrier_pending_count_b64(state);
// // CHECK_PTX70_SM80: call i32 @llvm.nvvm.mbarrier.pending.count.b64
traUnsubmitted
Done
ReplyInline Actions

Extra //

tra: Extra `//`
#endif
// CHECK: ret void
}
__device__ void nvvm_async_copy(__attribute__((address_space(3))) void* dst, __attribute__((address_space(1))) const void* src, long long* addr, __attribute__((address_space(3))) long long* sharedAddr) {
#if __CUDA_ARCH__ >= 800
// CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.mbarrier.arrive.b64
__nvvm_cp_async_mbarrier_arrive_b64(addr);
// CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.mbarrier.arrive.shared.b64
__nvvm_cp_async_mbarrier_arrive_shared_b64(sharedAddr);
// CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.mbarrier.arrive.noinc.b64
__nvvm_cp_async_mbarrier_arrive_noinc_b64(addr);
// CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.mbarrier.arrive.noinc.shared.b64
__nvvm_cp_async_mbarrier_arrive_noinc_shared_b64(sharedAddr);
// CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.ca.shared.global.4
__nvvm_cp_async_ca_shared_global_4(dst, src);
// CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.ca.shared.global.8
__nvvm_cp_async_ca_shared_global_8(dst, src);
// CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.ca.shared.global.16
__nvvm_cp_async_ca_shared_global_16(dst, src);
// CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.cg.shared.global.16
__nvvm_cp_async_cg_shared_global_16(dst, src);
// 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)
__nvvm_cp_async_wait_group(0);
// CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.wait.group(i32 8)
__nvvm_cp_async_wait_group(8);
// CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.wait.group(i32 16)
__nvvm_cp_async_wait_group(16);
// CHECK_PTX70_SM80: call void @llvm.nvvm.cp.async.wait.all
__nvvm_cp_async_wait_all();
#endif
// CHECK: ret void
}
No newline at end of file

llvm/include/llvm/IR/IntrinsicsNVVM.td

Show All 25 Lines
// * llvm.nvvm.max.ui --> select(x uge y, x, y) // * llvm.nvvm.max.ui --> select(x uge y, x, y)
// * llvm.nvvm.max.ull --> ibid. // * llvm.nvvm.max.ull --> ibid.
// * llvm.nvvm.max.i --> select(x sle y, x, y) // * llvm.nvvm.max.i --> select(x sle y, x, y)
// * llvm.nvvm.max.ll --> ibid. // * llvm.nvvm.max.ll --> ibid.
// * llvm.nvvm.max.ui --> select(x ule y, x, y) // * llvm.nvvm.max.ui --> select(x ule y, x, y)
// * llvm.nvvm.max.ull --> ibid. // * llvm.nvvm.max.ull --> ibid.
// * llvm.nvvm.h2f --> llvm.convert.to.fp16.f32 // * llvm.nvvm.h2f --> llvm.convert.to.fp16.f32
def llvm_i8ptr_ty : LLVMPointerType<llvm_i8_ty>; // i8*
def llvm_globali8ptr_ty : LLVMQualPointerType<llvm_i8_ty, 1>; // (global)i8*
def llvm_sharedi8ptr_ty : LLVMQualPointerType<llvm_i8_ty, 3>; // (shared)i8*
def llvm_i64ptr_ty : LLVMPointerType<llvm_i64_ty>; // i64*
def llvm_anyi64ptr_ty : LLVMAnyPointerType<llvm_i64_ty>; // (space)i64* def llvm_anyi64ptr_ty : LLVMAnyPointerType<llvm_i64_ty>; // (space)i64*
def llvm_sharedi64ptr_ty : LLVMQualPointerType<llvm_i64_ty, 3>; // (shared)i64*
traUnsubmitted
Done
ReplyInline Actions

llvm_globali8ptr_ty -> llvm_global_i8ptr_ty would make it a bit easier to read.

tra: `llvm_globali8ptr_ty` -> `llvm_global_i8ptr_ty` would make it a bit easier to read.
// //
// MISC // MISC
// //
// Helper class that represents a 'fragment' of an NVPTX *MMA instruction. // Helper class that represents a 'fragment' of an NVPTX *MMA instruction.
// Geom: m<M>n<N>k<K>. E.g. m8n32k16 // Geom: m<M>n<N>k<K>. E.g. m8n32k16
// Frag: [abcd] // Frag: [abcd]
▲ Show 20 Lines • Show All 1,004 Lines • ▼ Show 20 Lines// Bar.Sync
// Membar // Membar
def int_nvvm_membar_cta : GCCBuiltin<"__nvvm_membar_cta">, def int_nvvm_membar_cta : GCCBuiltin<"__nvvm_membar_cta">,
Intrinsic<[], [], []>; Intrinsic<[], [], []>;
def int_nvvm_membar_gl : GCCBuiltin<"__nvvm_membar_gl">, def int_nvvm_membar_gl : GCCBuiltin<"__nvvm_membar_gl">,
Intrinsic<[], [], []>; Intrinsic<[], [], []>;
def int_nvvm_membar_sys : GCCBuiltin<"__nvvm_membar_sys">, def int_nvvm_membar_sys : GCCBuiltin<"__nvvm_membar_sys">,
Intrinsic<[], [], []>; Intrinsic<[], [], []>;
// Async Copy
def int_nvvm_cp_async_mbarrier_arrive_b64 :
GCCBuiltin<"__nvvm_cp_async_mbarrier_arrive_b64">,
Intrinsic<[],[llvm_i64ptr_ty],[IntrConvergent]>;
def int_nvvm_cp_async_mbarrier_arrive_shared_b64 :
GCCBuiltin<"__nvvm_cp_async_mbarrier_arrive_shared_b64">,
Intrinsic<[],[llvm_sharedi64ptr_ty],[IntrConvergent]>;
def int_nvvm_cp_async_mbarrier_arrive_noinc_b64 :
GCCBuiltin<"__nvvm_cp_async_mbarrier_arrive_noinc_b64">,
Intrinsic<[],[llvm_i64ptr_ty],[IntrConvergent]>;
def int_nvvm_cp_async_mbarrier_arrive_noinc_shared_b64 :
GCCBuiltin<"__nvvm_cp_async_mbarrier_arrive_noinc_shared_b64">,
Intrinsic<[],[llvm_sharedi64ptr_ty],[IntrConvergent]>;
def int_nvvm_cp_async_ca_shared_global_4 :
GCCBuiltin<"__nvvm_cp_async_ca_shared_global_4">,
Intrinsic<[],[llvm_sharedi8ptr_ty, llvm_globali8ptr_ty],
[IntrArgMemOnly, 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 :
GCCBuiltin<"__nvvm_cp_async_ca_shared_global_8">,
Intrinsic<[],[llvm_sharedi8ptr_ty, llvm_globali8ptr_ty],
[IntrArgMemOnly, 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 :
GCCBuiltin<"__nvvm_cp_async_ca_shared_global_16">,
Intrinsic<[],[llvm_sharedi8ptr_ty, llvm_globali8ptr_ty],
[IntrArgMemOnly, NoAlias<ArgIndex<0>>, NoAlias<ArgIndex<1>>,
traUnsubmitted
Done
ReplyInline Actions

These all should probably have IntrArgMemOnly attribute and, possibly WriteOnly<0>,ReadOnly<1> and, maybe NoAlias on both arguments, too, because src/dest are in different nonoverlapping address spaces.

Also, the PTX spec is not clear on whether cp.async expects to see the pointer arguments in generic AS, or do they need to be converted to shared/global ones first. Normally, the instructions with .shared or .global in the name expect specific address space. If that's the case here, then we may need to use qualified pointer types here, too.

tra: These all should probably have `IntrArgMemOnly` attribute and, possibly `WriteOnly<0>…
WriteOnly<ArgIndex<0>>, ReadOnly<ArgIndex<1>>],
"llvm.nvvm.cp.async.ca.shared.global.16">;
def int_nvvm_cp_async_cg_shared_global_16 :
GCCBuiltin<"__nvvm_cp_async_cg_shared_global_16">,
Intrinsic<[],[llvm_sharedi8ptr_ty, llvm_globali8ptr_ty],
[IntrArgMemOnly, NoAlias<ArgIndex<0>>, NoAlias<ArgIndex<1>>,
WriteOnly<ArgIndex<0>>, ReadOnly<ArgIndex<1>>],
"llvm.nvvm.cp.async.cg.shared.global.16">;
traUnsubmitted
Done
ReplyInline Actions

This should have ImmArg as the argument must be an immediate value.

tra: This should have `ImmArg` as the argument must be an immediate value.
def int_nvvm_cp_async_commit_group :
GCCBuiltin<"__nvvm_cp_async_commit_group">,
Intrinsic<[],[],[]>;
def int_nvvm_cp_async_wait_group :
GCCBuiltin<"__nvvm_cp_async_wait_group">,
Intrinsic<[],[llvm_i32_ty],[ImmArg<ArgIndex<0>>]>;
def int_nvvm_cp_async_wait_all :
GCCBuiltin<"__nvvm_cp_async_wait_all">,
Intrinsic<[],[],[]>;
// mbarrier
def int_nvvm_mbarrier_init_b64 : GCCBuiltin<"__nvvm_mbarrier_init_b64">,
Intrinsic<[],[llvm_i64ptr_ty, llvm_i32_ty],[IntrConvergent]>;
def int_nvvm_mbarrier_init_shared_b64 :
GCCBuiltin<"__nvvm_mbarrier_init_shared_b64">,
Intrinsic<[],[llvm_sharedi64ptr_ty, llvm_i32_ty],[IntrConvergent]>;
def int_nvvm_mbarrier_inval_b64 : GCCBuiltin<"__nvvm_mbarrier_inval_b64">,
Intrinsic<[],[llvm_i64ptr_ty],[IntrConvergent]>;
def int_nvvm_mbarrier_inval_shared_b64 :
GCCBuiltin<"__nvvm_mbarrier_inval_shared_b64">,
Intrinsic<[],[llvm_sharedi64ptr_ty],[IntrConvergent]>;
traUnsubmitted
Done
ReplyInline Actions

These are probably safe to mark as IntrWriteMem, IntrArgMemOnly, WriteOnly<ArgIndex<0>>, NoCapture<ArgIndex<0>>

tra: These are probably safe to mark as `IntrWriteMem, IntrArgMemOnly, WriteOnly<ArgIndex<0>>…
def int_nvvm_mbarrier_arrive_b64 : GCCBuiltin<"__nvvm_mbarrier_arrive_b64">,
Intrinsic<[llvm_i64_ty],[llvm_i64ptr_ty],[IntrConvergent]>;
def int_nvvm_mbarrier_arrive_shared_b64 :
GCCBuiltin<"__nvvm_mbarrier_arrive_shared_b64">,
Intrinsic<[llvm_i64_ty],[llvm_sharedi64ptr_ty],[IntrConvergent]>;
def int_nvvm_mbarrier_arrive_noComplete_b64 :
GCCBuiltin<"__nvvm_mbarrier_arrive_noComplete_b64">,
Intrinsic<[llvm_i64_ty],[llvm_i64ptr_ty, llvm_i32_ty],[IntrConvergent]>;
def int_nvvm_mbarrier_arrive_noComplete_shared_b64 :
GCCBuiltin<"__nvvm_mbarrier_arrive_noComplete_shared_b64">,
Intrinsic<[llvm_i64_ty],[llvm_sharedi64ptr_ty, llvm_i32_ty],[IntrConvergent]>;
def int_nvvm_mbarrier_arrive_drop_b64 :
GCCBuiltin<"__nvvm_mbarrier_arrive_drop_b64">,
Intrinsic<[llvm_i64_ty],[llvm_i64ptr_ty],[IntrConvergent]>;
def int_nvvm_mbarrier_arrive_drop_shared_b64 :
GCCBuiltin<"__nvvm_mbarrier_arrive_drop_shared_b64">,
Intrinsic<[llvm_i64_ty],[llvm_sharedi64ptr_ty],[IntrConvergent]>;
def int_nvvm_mbarrier_arrive_drop_noComplete_b64 :
GCCBuiltin<"__nvvm_mbarrier_arrive_drop_noComplete_b64">,
Intrinsic<[llvm_i64_ty],[llvm_i64ptr_ty, llvm_i32_ty],[IntrConvergent]>;
def int_nvvm_mbarrier_arrive_drop_noComplete_shared_b64 :
GCCBuiltin<"__nvvm_mbarrier_arrive_drop_noComplete_shared_b64">,
Intrinsic<[llvm_i64_ty],[llvm_sharedi64ptr_ty, llvm_i32_ty],[IntrConvergent]>;
def int_nvvm_mbarrier_test_wait_b64 :
GCCBuiltin<"__nvvm_mbarrier_test_wait_b64">,
Intrinsic<[llvm_i1_ty],[llvm_i64ptr_ty, llvm_i64_ty],[IntrConvergent]>;
def int_nvvm_mbarrier_test_wait_shared_b64 :
GCCBuiltin<"__nvvm_mbarrier_test_wait_shared_b64">,
Intrinsic<[llvm_i1_ty],[llvm_sharedi64ptr_ty, llvm_i64_ty],[IntrConvergent]>;
def int_nvvm_mbarrier_pending_count_b64 :
GCCBuiltin<"__nvvm_mbarrier_pending_count_b64">,
Intrinsic<[llvm_i32_ty],[llvm_i64_ty],[IntrConvergent]>;
traUnsubmitted
Done
ReplyInline Actions

This one can be IntrNoMem.

tra: This one can be `IntrNoMem`.
// Generated within nvvm. Use for ldu on sm_20 or later. Second arg is the // Generated within nvvm. Use for ldu on sm_20 or later. Second arg is the
// pointer's alignment. // pointer's alignment.
def int_nvvm_ldu_global_i : Intrinsic<[llvm_anyint_ty], def int_nvvm_ldu_global_i : Intrinsic<[llvm_anyint_ty],
[LLVMAnyPointerType<LLVMMatchType<0>>, llvm_i32_ty], [LLVMAnyPointerType<LLVMMatchType<0>>, llvm_i32_ty],
[IntrReadMem, IntrArgMemOnly, NoCapture<ArgIndex<0>>], [IntrReadMem, IntrArgMemOnly, NoCapture<ArgIndex<0>>],
"llvm.nvvm.ldu.global.i">; "llvm.nvvm.ldu.global.i">;
def int_nvvm_ldu_global_f : Intrinsic<[llvm_anyfloat_ty], def int_nvvm_ldu_global_f : Intrinsic<[llvm_anyfloat_ty],
[LLVMAnyPointerType<LLVMMatchType<0>>, llvm_i32_ty], [LLVMAnyPointerType<LLVMMatchType<0>>, llvm_i32_ty],
▲ Show 20 Lines • Show All 3,101 Lines • Show Last 20 Lines

llvm/lib/Target/NVPTX/NVPTXInstrInfo.td

Show First 20 Lines • Show All 138 Lines • ▼ Show 20 Lines
def True : Predicate<"true">; def True : Predicate<"true">;
def hasPTX31 : Predicate<"Subtarget->getPTXVersion() >= 31">; def hasPTX31 : Predicate<"Subtarget->getPTXVersion() >= 31">;
def hasPTX60 : Predicate<"Subtarget->getPTXVersion() >= 60">; def hasPTX60 : Predicate<"Subtarget->getPTXVersion() >= 60">;
def hasPTX61 : Predicate<"Subtarget->getPTXVersion() >= 61">; def hasPTX61 : Predicate<"Subtarget->getPTXVersion() >= 61">;
def hasPTX63 : Predicate<"Subtarget->getPTXVersion() >= 63">; def hasPTX63 : Predicate<"Subtarget->getPTXVersion() >= 63">;
def hasPTX64 : Predicate<"Subtarget->getPTXVersion() >= 64">; def hasPTX64 : Predicate<"Subtarget->getPTXVersion() >= 64">;
def hasPTX70 : Predicate<"Subtarget->getPTXVersion() >= 70">;
def hasSM30 : Predicate<"Subtarget->getSmVersion() >= 30">; def hasSM30 : Predicate<"Subtarget->getSmVersion() >= 30">;
def hasSM70 : Predicate<"Subtarget->getSmVersion() >= 70">; def hasSM70 : Predicate<"Subtarget->getSmVersion() >= 70">;
def hasSM72 : Predicate<"Subtarget->getSmVersion() >= 72">; def hasSM72 : Predicate<"Subtarget->getSmVersion() >= 72">;
def hasSM75 : Predicate<"Subtarget->getSmVersion() >= 75">; def hasSM75 : Predicate<"Subtarget->getSmVersion() >= 75">;
def hasSM80 : Predicate<"Subtarget->getSmVersion() >= 80">;
// non-sync shfl instructions are not available on sm_70+ in PTX6.4+ // non-sync shfl instructions are not available on sm_70+ in PTX6.4+
def hasSHFL : Predicate<"!(Subtarget->getSmVersion() >= 70" def hasSHFL : Predicate<"!(Subtarget->getSmVersion() >= 70"
"&& Subtarget->getPTXVersion() >= 64)">; "&& Subtarget->getPTXVersion() >= 64)">;
def useShortPtr : Predicate<"useShortPointers()">; def useShortPtr : Predicate<"useShortPointers()">;
def useFP16Math: Predicate<"Subtarget->allowFP16Math()">; def useFP16Math: Predicate<"Subtarget->allowFP16Math()">;
▲ Show 20 Lines • Show All 2,984 Lines • Show Last 20 Lines

llvm/lib/Target/NVPTX/NVPTXIntrinsics.td

  • This file is larger than 256 KB, so syntax highlighting is disabled by default.
Show First 20 Lines • Show All 283 Lines • ▼ Show 20 Linesclass MEMBAR<string StrOp, Intrinsic IntOP> :
StrOp, [(IntOP)]>; StrOp, [(IntOP)]>;
def INT_MEMBAR_CTA : MEMBAR<"membar.cta;", int_nvvm_membar_cta>; def INT_MEMBAR_CTA : MEMBAR<"membar.cta;", int_nvvm_membar_cta>;
def INT_MEMBAR_GL : MEMBAR<"membar.gl;", int_nvvm_membar_gl>; def INT_MEMBAR_GL : MEMBAR<"membar.gl;", int_nvvm_membar_gl>;
def INT_MEMBAR_SYS : MEMBAR<"membar.sys;", int_nvvm_membar_sys>; def INT_MEMBAR_SYS : MEMBAR<"membar.sys;", int_nvvm_membar_sys>;
//----------------------------------- //-----------------------------------
// Async Copy Functions
//-----------------------------------
multiclass CP_ASYNC_MBARRIER_ARRIVE<string NoInc, string AddrSpace, Intrinsic Intrin> {
def _32 : NVPTXInst<(outs), (ins Int32Regs:$addr),
!strconcat("cp.async.mbarrier.arrive", NoInc, AddrSpace, ".b64 [$addr];"),
[(Intrin Int32Regs:$addr)]>,
Requires<[hasPTX70, hasSM80]>;
def _64 : NVPTXInst<(outs), (ins Int64Regs:$addr),
!strconcat("cp.async.mbarrier.arrive", NoInc, AddrSpace, ".b64 [$addr];"),
[(Intrin Int64Regs:$addr)]>,
Requires<[hasPTX70, hasSM80]>;
}
defm CP_ASYNC_MBARRIER_ARRIVE :
CP_ASYNC_MBARRIER_ARRIVE<"", "", int_nvvm_cp_async_mbarrier_arrive_b64>;
defm CP_ASYNC_MBARRIER_ARRIVE_SHARED :
CP_ASYNC_MBARRIER_ARRIVE<"", ".shared", int_nvvm_cp_async_mbarrier_arrive_shared_b64>;
defm CP_ASYNC_MBARRIER_ARRIVE_NOINC :
CP_ASYNC_MBARRIER_ARRIVE<".noinc", "", int_nvvm_cp_async_mbarrier_arrive_noinc_b64>;
defm CP_ASYNC_MBARRIER_ARRIVE_NOINC_SHARED :
CP_ASYNC_MBARRIER_ARRIVE<".noinc", ".shared", int_nvvm_cp_async_mbarrier_arrive_noinc_shared_b64>;
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)]>,
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)]>,
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>;
defm CP_ASYNC_CA_SHARED_GLOBAL_8 :
CP_ASYNC_CA_SHARED_GLOBAL_I<"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]>;
}
defm CP_ASYNC_CG_SHARED_GLOBAL_16 :
CP_ASYNC_CG_SHARED_GLOBAL<"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)]>,
Requires<[hasPTX70, hasSM80]>;
def CP_ASYNC_WAIT_GROUP :
NVPTXInst<(outs), (ins i32imm:$n), "cp.async.wait_group $n;",
[(int_nvvm_cp_async_wait_group (i32 timm:$n))]>,
Requires<[hasPTX70, hasSM80]>;
traUnsubmitted
Done
ReplyInline Actions

This does not look right. If I read the PTX spec correctly, the argument can't be a register.

https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-wait-group

Operand N is an integer constant.

tra: This does not look right. If I read the PTX spec correctly, the argument can't be a register.
def CP_ASYNC_WAIT_ALL :
NVPTXInst<(outs), (ins), "cp.async.wait_all;",
[(int_nvvm_cp_async_wait_all)]>,
Requires<[hasPTX70, hasSM80]>;
//-----------------------------------
// MBarrier Functions
//-----------------------------------
multiclass MBARRIER_INIT<string AddrSpace, Intrinsic Intrin> {
def _32 : NVPTXInst<(outs), (ins Int32Regs:$addr, Int32Regs:$count),
!strconcat("mbarrier.init", AddrSpace, ".b64 [$addr], $count;"),
[(Intrin Int32Regs:$addr, Int32Regs:$count)]>,
Requires<[hasPTX70, hasSM80]>;
def _64 : NVPTXInst<(outs), (ins Int64Regs:$addr, Int32Regs:$count),
!strconcat("mbarrier.init", AddrSpace, ".b64 [$addr], $count;"),
[(Intrin Int64Regs:$addr, Int32Regs:$count)]>,
Requires<[hasPTX70, hasSM80]>;
}
defm MBARRIER_INIT : MBARRIER_INIT<"", int_nvvm_mbarrier_init_b64>;
defm MBARRIER_INIT_SHARED : MBARRIER_INIT<".shared",
int_nvvm_mbarrier_init_shared_b64>;
multiclass MBARRIER_INVAL<string AddrSpace, Intrinsic Intrin> {
def _32 : NVPTXInst<(outs), (ins Int32Regs:$addr),
!strconcat("mbarrier.inval", AddrSpace, ".b64 [$addr];"),
[(Intrin Int32Regs:$addr)]>,
Requires<[hasPTX70, hasSM80]>;
def _64 : NVPTXInst<(outs), (ins Int64Regs:$addr),
!strconcat("mbarrier.inval", AddrSpace, ".b64 [$addr];"),
[(Intrin Int64Regs:$addr)]>,
Requires<[hasPTX70, hasSM80]>;
}
defm MBARRIER_INVAL : MBARRIER_INVAL<"", int_nvvm_mbarrier_inval_b64>;
defm MBARRIER_INVAL_SHARED : MBARRIER_INVAL<".shared",
int_nvvm_mbarrier_inval_shared_b64>;
multiclass MBARRIER_ARRIVE<string AddrSpace, Intrinsic Intrin> {
def _32 : NVPTXInst<(outs Int64Regs:$state), (ins Int32Regs:$addr),
!strconcat("mbarrier.arrive", AddrSpace, ".b64 $state, [$addr];"),
[(set Int64Regs:$state, (Intrin Int32Regs:$addr))]>,
Requires<[hasPTX70, hasSM80]>;
def _64 : NVPTXInst<(outs Int64Regs:$state), (ins Int64Regs:$addr),
!strconcat("mbarrier.arrive", AddrSpace, ".b64 $state, [$addr];"),
[(set Int64Regs:$state, (Intrin Int64Regs:$addr))]>,
Requires<[hasPTX70, hasSM80]>;
}
defm MBARRIER_ARRIVE : MBARRIER_ARRIVE<"", int_nvvm_mbarrier_arrive_b64>;
defm MBARRIER_ARRIVE_SHARED :
MBARRIER_ARRIVE<".shared", int_nvvm_mbarrier_arrive_shared_b64>;
multiclass MBARRIER_ARRIVE_NOCOMPLETE<string AddrSpace, Intrinsic Intrin> {
def _32 : NVPTXInst<(outs Int64Regs:$state),
(ins Int32Regs:$addr, Int32Regs:$count),
!strconcat("mbarrier.arrive.noComplete", AddrSpace,
traUnsubmitted
Not Done
ReplyInline Actions

Does ptxas accept all-lower-case nocomplete?

The no*C*omplete stands out as a sore thumb. Capital letters are used in few LLVM intrinsics, so it's not a showstopper, but I think lower case everywhere makes more sense. WDYT?

tra: Does ptxas accept all-lower-case `nocomplete`? The `no*C*omplete` stands out as a sore thumb.
nyallocAuthorUnsubmitted
Done
ReplyInline Actions

I 100% agree. For some reason the PTX decide to use camel case here. They also throw some snake case in other names. We can change the LLVM intrinsics to use a more consistent naming scheme, but it will come at the cost that it's no longer a clean mapping of names to the PTX.

nyalloc: I 100% agree. For some reason the PTX decide to use camel case here. They also throw some snake…
".b64 $state, [$addr], $count;"),
[(set Int64Regs:$state, (Intrin Int32Regs:$addr, Int32Regs:$count))]>,
Requires<[hasPTX70, hasSM80]>;
def _64 : NVPTXInst<(outs Int64Regs:$state),
(ins Int64Regs:$addr, Int32Regs:$count),
!strconcat("mbarrier.arrive.noComplete", AddrSpace,
".b64 $state, [$addr], $count;"),
[(set Int64Regs:$state, (Intrin Int64Regs:$addr, Int32Regs:$count))]>,
Requires<[hasPTX70, hasSM80]>;
}
defm MBARRIER_ARRIVE_NOCOMPLETE :
MBARRIER_ARRIVE_NOCOMPLETE<"", int_nvvm_mbarrier_arrive_noComplete_b64>;
defm MBARRIER_ARRIVE_NOCOMPLETE_SHARED :
MBARRIER_ARRIVE_NOCOMPLETE<".shared", int_nvvm_mbarrier_arrive_noComplete_shared_b64>;
multiclass MBARRIER_ARRIVE_DROP<string AddrSpace, Intrinsic Intrin> {
def _32 : NVPTXInst<(outs Int64Regs:$state), (ins Int32Regs:$addr),
!strconcat("mbarrier.arrive_drop", AddrSpace,
".b64 $state, [$addr];"),
[(set Int64Regs:$state, (Intrin Int32Regs:$addr))]>,
Requires<[hasPTX70, hasSM80]>;
def _64 : NVPTXInst<(outs Int64Regs:$state), (ins Int64Regs:$addr),
!strconcat("mbarrier.arrive_drop", AddrSpace,
".b64 $state, [$addr];"),
[(set Int64Regs:$state, (Intrin Int64Regs:$addr))]>,
Requires<[hasPTX70, hasSM80]>;
}
defm MBARRIER_ARRIVE_DROP :
MBARRIER_ARRIVE_DROP<"", int_nvvm_mbarrier_arrive_drop_b64>;
defm MBARRIER_ARRIVE_DROP_SHARED :
MBARRIER_ARRIVE_DROP<".shared", int_nvvm_mbarrier_arrive_drop_shared_b64>;
multiclass MBARRIER_ARRIVE_DROP_NOCOMPLETE<string AddrSpace, Intrinsic Intrin> {
def _32 : NVPTXInst<(outs Int64Regs:$state),
(ins Int32Regs:$addr, Int32Regs:$count),
!strconcat("mbarrier.arrive_drop.noComplete", AddrSpace,
".b64 $state, [$addr], $count;"),
[(set Int64Regs:$state, (Intrin Int32Regs:$addr, Int32Regs:$count))]>,
Requires<[hasPTX70, hasSM80]>;
def _64 : NVPTXInst<(outs Int64Regs:$state),
(ins Int64Regs:$addr, Int32Regs:$count),
!strconcat("mbarrier.arrive_drop.noComplete", AddrSpace,
".b64 $state, [$addr], $count;"),
[(set Int64Regs:$state, (Intrin Int64Regs:$addr, Int32Regs:$count))]>,
Requires<[hasPTX70, hasSM80]>;
}
defm MBARRIER_ARRIVE_DROP_NOCOMPLETE :
MBARRIER_ARRIVE_DROP_NOCOMPLETE<"", int_nvvm_mbarrier_arrive_drop_noComplete_b64>;
defm MBARRIER_ARRIVE_DROP_NOCOMPLETE_SHARED :
MBARRIER_ARRIVE_DROP_NOCOMPLETE<".shared",
int_nvvm_mbarrier_arrive_drop_noComplete_shared_b64>;
multiclass MBARRIER_TEST_WAIT<string AddrSpace, Intrinsic Intrin> {
def _32 : NVPTXInst<(outs Int1Regs:$res), (ins Int32Regs:$addr, Int64Regs:$state),
!strconcat("mbarrier.test_wait", AddrSpace, ".b64 $res, [$addr], $state;"),
[(set Int1Regs:$res, (Intrin Int32Regs:$addr, Int64Regs:$state))]>,
Requires<[hasPTX70, hasSM80]>;
def _64 : NVPTXInst<(outs Int1Regs:$res), (ins Int64Regs:$addr, Int64Regs:$state),
!strconcat("mbarrier.test_wait", AddrSpace, ".b64 $res, [$addr], $state;"),
[(set Int1Regs:$res, (Intrin Int64Regs:$addr, Int64Regs:$state))]>,
Requires<[hasPTX70, hasSM80]>;
}
defm MBARRIER_TEST_WAIT :
MBARRIER_TEST_WAIT<"", int_nvvm_mbarrier_test_wait_b64>;
defm MBARRIER_TEST_WAIT_SHARED :
MBARRIER_TEST_WAIT<".shared", int_nvvm_mbarrier_test_wait_shared_b64>;
class MBARRIER_PENDING_COUNT<Intrinsic Intrin> :
NVPTXInst<(outs Int32Regs:$res), (ins Int64Regs:$state),
"mbarrier.pending_count.b64 $res, $state;",
[(set Int32Regs:$res, (Intrin Int64Regs:$state))]>,
Requires<[hasPTX70, hasSM80]>;
def MBARRIER_PENDING_COUNT :
MBARRIER_PENDING_COUNT<int_nvvm_mbarrier_pending_count_b64>;
//-----------------------------------
// Math Functions // Math Functions
//----------------------------------- //-----------------------------------
// Map min(1.0, max(0.0, x)) to sat(x) // Map min(1.0, max(0.0, x)) to sat(x)
// Note that max(0.0, min(x, 1.0)) cannot be mapped to sat(x) because when x is // Note that max(0.0, min(x, 1.0)) cannot be mapped to sat(x) because when x is
// NaN // NaN
// max(0.0, min(x, 1.0)) is 1.0 while sat(x) is 0. // max(0.0, min(x, 1.0)) is 1.0 while sat(x) is 0.
// Same story for fmax, fmin. // Same story for fmax, fmin.
▲ Show 20 Lines • Show All 7,313 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/async-copy.ll

  • This file was added.
; RUN: llc < %s -march=nvptx -mcpu=sm_80 -mattr=+ptx70 | FileCheck %s -check-prefix=CHECK_PTX32
; RUN: llc < %s -march=nvptx64 -mcpu=sm_80 -mattr=+ptx70 | FileCheck %s -check-prefix=CHECK_PTX64
declare void @llvm.nvvm.cp.async.wait.group(i32)
traUnsubmitted
Done
ReplyInline Actions

No tests for cp.async.wait* and cp.async.commit_group

tra: No tests for `cp.async.wait*` and `cp.async.commit_group`
define void @asyncwaitgroup() {
; CHECK_PTX32: cp.async.wait_group 8;
; CHECK_PTX64: cp.async.wait_group 8;
traUnsubmitted
Done
ReplyInline Actions

I'd recommend adding a common check label (le'ts say ALL and running the tests with --check-prefixes=ALL,CHECK_PTX64.

This would allow you to use ALL for things that do not change. In this case ; ALL: cp.async.wait_group 8;
It would also be great to add ALL-LABEL: <function_name> for each function to limit the range FileCheck operates for.

tra: I'd recommend adding a common check label (le'ts say `ALL` and running the tests with `--check…
tail call void @llvm.nvvm.cp.async.wait.group(i32 8)
; CHECK_PTX32: cp.async.wait_group 0;
; CHECK_PTX64: cp.async.wait_group 0;
tail call void @llvm.nvvm.cp.async.wait.group(i32 0)
; CHECK_PTX32: cp.async.wait_group 16;
; CHECK_PTX64: 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()
define void @asyncwaitall() {
; CHECK_PTX32: cp.async.wait_all
; CHECK_PTX64: cp.async.wait_all
tail call void @llvm.nvvm.cp.async.wait.all()
ret void
}
declare void @llvm.nvvm.cp.async.commit.group()
define void @asynccommitgroup() {
; CHECK_PTX32: cp.async.commit_group
; CHECK_PTX64: cp.async.commit_group
tail call void @llvm.nvvm.cp.async.commit.group()
ret void
}
declare void @llvm.nvvm.cp.async.mbarrier.arrive.b64(i64* %a)
declare void @llvm.nvvm.cp.async.mbarrier.arrive.shared.b64(i64 addrspace(3)* %a)
declare void @llvm.nvvm.cp.async.mbarrier.arrive.noinc.b64(i64* %a)
declare void @llvm.nvvm.cp.async.mbarrier.arrive.noinc.shared.b64(i64 addrspace(3)* %a)
traUnsubmitted
Not Done
ReplyInline Actions

All functions in the file should use -LABEL checks.

tra: All functions in the file should use `-LABEL` checks.
define void @asyncmbarrier(i64* %a) {
; CHECK_PTX32: cp.async.mbarrier.arrive.b64 [%r{{[0-9]+}}];
; CHECK_PTX64: cp.async.mbarrier.arrive.b64 [%rd{{[0-9]+}}];
tail call void @llvm.nvvm.cp.async.mbarrier.arrive.b64(i64* %a)
ret void
}
define void @asyncmbarriershared(i64 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]+}}];
tail call void @llvm.nvvm.cp.async.mbarrier.arrive.shared.b64(i64 addrspace(3)* %a)
ret void
}
define void @asyncmbarriernoinc(i64* %a) {
; CHECK_PTX32: cp.async.mbarrier.arrive.noinc.b64 [%r{{[0-9]+}}];
; CHECK_PTX64: cp.async.mbarrier.arrive.noinc.b64 [%rd{{[0-9]+}}];
tail call void @llvm.nvvm.cp.async.mbarrier.arrive.noinc.b64(i64* %a)
ret void
}
define void @asyncmbarriernoincshared(i64 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]+}}];
tail call void @llvm.nvvm.cp.async.mbarrier.arrive.noinc.shared.b64(i64 addrspace(3)* %a)
ret void
}
declare void @llvm.nvvm.cp.async.ca.shared.global.4(i8 addrspace(3)* %a, i8 addrspace(1)* %b)
define void @asynccasharedglobal4i8(i8 addrspace(3)* %a, i8 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(i8 addrspace(3)* %a, i8 addrspace(1)* %b)
ret void
}
declare void @llvm.nvvm.cp.async.ca.shared.global.8(i8 addrspace(3)* %a, i8 addrspace(1)* %b)
define void @asynccasharedglobal8i8(i8 addrspace(3)* %a, i8 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(i8 addrspace(3)* %a, i8 addrspace(1)* %b)
ret void
}
declare void @llvm.nvvm.cp.async.ca.shared.global.16(i8 addrspace(3)* %a, i8 addrspace(1)* %b)
define void @asynccasharedglobal16i8(i8 addrspace(3)* %a, i8 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(i8 addrspace(3)* %a, i8 addrspace(1)* %b)
ret void
}
declare void @llvm.nvvm.cp.async.cg.shared.global.16(i8 addrspace(3)* %a, i8 addrspace(1)* %b)
define void @asynccgsharedglobal16i8(i8 addrspace(3)* %a, i8 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(i8 addrspace(3)* %a, i8 addrspace(1)* %b)
ret void
}

llvm/test/CodeGen/NVPTX/mbarrier.ll

  • This file was added.
; RUN: llc < %s -march=nvptx -mcpu=sm_80 | FileCheck %s -check-prefix=x32
; RUN: llc < %s -march=nvptx64 -mcpu=sm_80 | FileCheck %s -check-prefix=x64
declare void @llvm.nvvm.mbarrier.init.b64(i64* %a, i32 %b)
declare void @llvm.nvvm.mbarrier.init.shared.b64(i64 addrspace(3)* %a, i32 %b)
define void @barrierinit(i64* %a, i32 %b) {
; x32: mbarrier.init.b64 [%r{{[0-9]+}}], %r{{[0-9]+}};
traUnsubmitted
Not Done
ReplyInline Actions

Same here. Please add -LABEL checks for all functions.

tra: Same here. Please add `-LABEL` checks for all functions.
nyallocAuthorUnsubmitted
Done
ReplyInline Actions

Will do!

nyalloc: Will do!
; x64: mbarrier.init.b64 [%rd{{[0-9]+}}], %r{{[0-9]+}};
tail call void @llvm.nvvm.mbarrier.init.b64(i64* %a, i32 %b)
ret void
}
define void @barrierinitshared(i64 addrspace(3)* %a, i32 %b) {
; x32: mbarrier.init.shared.b64 [%r{{[0-9]+}}], %r{{[0-9]+}};
; x64: mbarrier.init.shared.b64 [%rd{{[0-9]+}}], %r{{[0-9]+}};
tail call void @llvm.nvvm.mbarrier.init.shared.b64(i64 addrspace(3)* %a, i32 %b)
ret void
}
declare void @llvm.nvvm.mbarrier.inval.b64(i64* %a)
declare void @llvm.nvvm.mbarrier.inval.shared.b64(i64 addrspace(3)* %a)
define void @barrierinval(i64* %a) {
; x32: mbarrier.inval.b64 [%r{{[0-1]+}}];
; x64: mbarrier.inval.b64 [%rd{{[0-1]+}}];
tail call void @llvm.nvvm.mbarrier.inval.b64(i64* %a)
ret void
}
define void @barrierinvalshared(i64 addrspace(3)* %a) {
; x32: mbarrier.inval.shared.b64 [%r{{[0-1]+}}];
; x64: mbarrier.inval.shared.b64 [%rd{{[0-1]+}}];
tail call void @llvm.nvvm.mbarrier.inval.shared.b64(i64 addrspace(3)* %a)
ret void
}
declare i64 @llvm.nvvm.mbarrier.arrive.b64(i64* %a)
declare i64 @llvm.nvvm.mbarrier.arrive.shared.b64(i64 addrspace(3)* %a)
define void @barrierarrive(i64* %a) {
; x32: mbarrier.arrive.b64 %rd{{[0-9]+}}, [%r{{[0-9]+}}];
; x64: mbarrier.arrive.b64 %rd{{[0-9]+}}, [%rd{{[0-9]+}}];
%ret = tail call i64 @llvm.nvvm.mbarrier.arrive.b64(i64* %a)
ret void
}
define void @barrierarriveshared(i64 addrspace(3)* %a) {
; x32: mbarrier.arrive.shared.b64 %rd{{[0-9]+}}, [%r{{[0-9]+}}];
; x64: mbarrier.arrive.shared.b64 %rd{{[0-9]+}}, [%rd{{[0-9]+}}];
%ret = tail call i64 @llvm.nvvm.mbarrier.arrive.shared.b64(i64 addrspace(3)* %a)
ret void
}
declare i64 @llvm.nvvm.mbarrier.arrive.noComplete.b64(i64* %a, i32 %b)
declare i64 @llvm.nvvm.mbarrier.arrive.noComplete.shared.b64(i64 addrspace(3)* %a, i32 %b)
define void @barrierarrivenoComplete(i64* %a, i32 %b) {
; x32: mbarrier.arrive.noComplete.b64 %rd{{[0-9]+}}, [%r{{[0-9]+}}], %r{{[0-9]+}};
; x64: mbarrier.arrive.noComplete.b64 %rd{{[0-9]+}}, [%rd{{[0-9]+}}], %r{{[0-9]+}};
%ret = tail call i64 @llvm.nvvm.mbarrier.arrive.noComplete.b64(i64* %a, i32 %b)
ret void
}
define void @barrierarrivenoCompleteshared(i64 addrspace(3)* %a, i32 %b) {
; x32: mbarrier.arrive.noComplete.shared.b64 %rd{{[0-9]+}}, [%r{{[0-9]+}}], %r{{[0-9]+}};
; x64: mbarrier.arrive.noComplete.shared.b64 %rd{{[0-9]+}}, [%rd{{[0-9]+}}], %r{{[0-9]+}};
%ret = tail call i64 @llvm.nvvm.mbarrier.arrive.noComplete.shared.b64(i64 addrspace(3)* %a, i32 %b)
ret void
}
declare i64 @llvm.nvvm.mbarrier.arrive.drop.b64(i64* %a)
declare i64 @llvm.nvvm.mbarrier.arrive.drop.shared.b64(i64 addrspace(3)* %a)
define void @barrierarrivedrop(i64* %a) {
; x32: mbarrier.arrive_drop.b64 %rd{{[0-9]+}}, [%r{{[0-9]+}}];
; x64: mbarrier.arrive_drop.b64 %rd{{[0-9]+}}, [%rd{{[0-9]+}}];
%ret = tail call i64 @llvm.nvvm.mbarrier.arrive.drop.b64(i64* %a)
ret void
}
define void @barrierarrivedropshared(i64 addrspace(3)* %a) {
; x32: mbarrier.arrive_drop.shared.b64 %rd{{[0-9]+}}, [%r{{[0-9]+}}];
; x64: mbarrier.arrive_drop.shared.b64 %rd{{[0-9]+}}, [%rd{{[0-9]+}}];
%ret = tail call i64 @llvm.nvvm.mbarrier.arrive.drop.shared.b64(i64 addrspace(3)* %a)
ret void
}
declare i64 @llvm.nvvm.mbarrier.arrive.drop.noComplete.b64(i64* %a, i32 %b)
declare i64 @llvm.nvvm.mbarrier.arrive.drop.noComplete.shared.b64(i64 addrspace(3)* %a, i32 %b)
define void @barrierarrivedropnoComplete(i64* %a, i32 %b) {
; x32: mbarrier.arrive_drop.noComplete.b64 %rd{{[0-9]+}}, [%r{{[0-9]+}}], %r{{[0-9]+}};
; x64: mbarrier.arrive_drop.noComplete.b64 %rd{{[0-9]+}}, [%rd{{[0-9]+}}], %r{{[0-9]+}};
%ret = tail call i64 @llvm.nvvm.mbarrier.arrive.drop.noComplete.b64(i64* %a, i32 %b)
ret void
}
define void @barrierarrivedropnoCompleteshared(i64 addrspace(3)* %a, i32 %b) {
; x32: mbarrier.arrive_drop.noComplete.shared.b64 %rd{{[0-9]+}}, [%r{{[0-9]+}}], %r{{[0-9]+}};
; x64: mbarrier.arrive_drop.noComplete.shared.b64 %rd{{[0-9]+}}, [%rd{{[0-9]+}}], %r{{[0-9]+}};
%ret = tail call i64 @llvm.nvvm.mbarrier.arrive.drop.noComplete.shared.b64(i64 addrspace(3)* %a, i32 %b)
ret void
}
declare i1 @llvm.nvvm.mbarrier.test.wait.b64(i64* %a, i64 %b)
declare i1 @llvm.nvvm.mbarrier.test.wait.shared.b64(i64 addrspace(3)* %a, i64 %b)
define void @barriertestwait(i64* %a, i64 %b) {
; x32: mbarrier.test_wait.b64 %p{{[0-9]+}}, [%r{{[0-9]+}}], %rd{{[0-9]+}};
; x64: mbarrier.test_wait.b64 %p{{[0-9]+}}, [%rd{{[0-9]+}}], %rd{{[0-9]+}};
%ret = tail call i1 @llvm.nvvm.mbarrier.test.wait.b64(i64* %a, i64 %b)
ret void
}
define void @barriertestwaitshared(i64 addrspace(3)* %a, i64 %b) {
; x32: mbarrier.test_wait.shared.b64 %p{{[0-9]+}}, [%r{{[0-9]+}}], %rd{{[0-9]+}};
; x64: mbarrier.test_wait.shared.b64 %p{{[0-9]+}}, [%rd{{[0-9]+}}], %rd{{[0-9]+}};
%ret = tail call i1 @llvm.nvvm.mbarrier.test.wait.shared.b64(i64 addrspace(3)* %a, i64 %b)
ret void
}
declare i32 @llvm.nvvm.mbarrier.pending.count.b64(i64 %b)
define void @barrierpendingcount(i64* %a, i64 %b) {
; x32: mbarrier.pending_count.b64 %r{{[0-9]+}}, %rd{{[0-9]+}};
; x64: mbarrier.pending_count.b64 %r{{[0-9]+}}, %rd{{[0-9]+}};
%ret = tail call i32 @llvm.nvvm.mbarrier.pending.count.b64(i64 %b)
ret void
}