[NVPTX] Support short entry param pointers (#206512)
This commit is contained in:
Alex MacLean
GitHub
parent
3782c00acd
commit
aef1e4f249
5 changed files with 518 additions and 244 deletions
|
|
@ -144,7 +144,7 @@
|
|||
|
||||
// RUN: %clang_cc1 -triple nvptx-unknown -o - -emit-llvm %s | \
|
||||
// RUN: FileCheck %s -check-prefix=NVPTX
|
||||
// NVPTX: target datalayout = "e-p:32:32-p6:32:32-p7:32:32-i64:64-i128:128-i256:256-v16:16-v32:32-n16:32:64"
|
||||
// NVPTX: target datalayout = "e-p:32:32-i64:64-i128:128-i256:256-v16:16-v32:32-n16:32:64"
|
||||
|
||||
// RUN: %clang_cc1 -triple nvptx64-unknown -o - -emit-llvm %s | \
|
||||
// RUN: FileCheck %s -check-prefix=NVPTX64
|
||||
|
|
|
|||
|
|
@ -4071,11 +4071,13 @@ SDValue NVPTXTargetLowering::LowerFormalArguments(
|
|||
SelectionDAG &DAG, SmallVectorImpl<SDValue> &InVals) const {
|
||||
const DataLayout &DL = DAG.getDataLayout();
|
||||
LLVMContext &Ctx = *DAG.getContext();
|
||||
auto PtrVT = getPointerTy(DAG.getDataLayout());
|
||||
|
||||
const Function &F = DAG.getMachineFunction().getFunction();
|
||||
const bool IsKernel = isKernelFunction(F);
|
||||
|
||||
const MVT PtrVT = getPointerTy(DL, IsKernel ? ADDRESS_SPACE_ENTRY_PARAM
|
||||
: ADDRESS_SPACE_LOCAL);
|
||||
|
||||
SDValue Root = DAG.getRoot();
|
||||
SmallVector<SDValue, 16> OutChains;
|
||||
|
||||
|
|
@ -4129,18 +4131,17 @@ SDValue NVPTXTargetLowering::LowerFormalArguments(
|
|||
const auto &ByvalIn = ArgIns[0];
|
||||
assert(getValueType(DL, Ty) == ByvalIn.VT &&
|
||||
"Ins type did not match function type");
|
||||
assert(ByvalIn.VT == PtrVT && "ByVal argument must be a pointer");
|
||||
|
||||
SDValue P;
|
||||
if (IsKernel) {
|
||||
assert(Arg.getType()->getPointerAddressSpace() ==
|
||||
ADDRESS_SPACE_ENTRY_PARAM &&
|
||||
assert(Ty->getPointerAddressSpace() == ADDRESS_SPACE_ENTRY_PARAM &&
|
||||
"Kernel ByVal argument must be lowered to the param address "
|
||||
"space by NVPTXLowerArgs");
|
||||
P = ArgSymbol;
|
||||
P.getNode()->setIROrder(Arg.getArgNo() + 1);
|
||||
} else {
|
||||
P = DAG.getNode(NVPTXISD::MoveParam, dl, ByvalIn.VT, ArgSymbol);
|
||||
P = DAG.getNode(NVPTXISD::MoveParam, dl, ArgSymbol.getValueType(),
|
||||
ArgSymbol);
|
||||
P.getNode()->setIROrder(Arg.getArgNo() + 1);
|
||||
P = DAG.getAddrSpaceCast(dl, ByvalIn.VT, P, ADDRESS_SPACE_LOCAL,
|
||||
ADDRESS_SPACE_GENERIC);
|
||||
|
|
|
|||
|
|
@ -456,19 +456,31 @@ static std::string computeX86DataLayout(const Triple &TT) {
|
|||
}
|
||||
|
||||
static std::string computeNVPTXDataLayout(const Triple &T, StringRef ABIName) {
|
||||
bool Is64Bit = T.getArch() == Triple::nvptx64;
|
||||
const bool Is32Bit = T.getArch() == Triple::nvptx;
|
||||
const bool IsShortPtr = ABIName == "shortptr";
|
||||
std::string Ret = "e";
|
||||
|
||||
// Tensor Memory (addrspace:6) is always 32-bits.
|
||||
// Distributed Shared Memory (addrspace:7) follows shared memory
|
||||
// (addrspace:3).
|
||||
if (!Is64Bit)
|
||||
Ret += "-p:32:32-p6:32:32-p7:32:32";
|
||||
else if (ABIName == "shortptr")
|
||||
Ret += "-p3:32:32-p4:32:32-p5:32:32-p6:32:32-p7:32:32";
|
||||
else
|
||||
if (Is32Bit) {
|
||||
Ret += "-p:32:32";
|
||||
} else {
|
||||
// Keep the pointer specifications sorted by address space.
|
||||
//
|
||||
// In shortptr mode, specify the following address spaces as 32-bits:
|
||||
// - shared (addrspace:3)
|
||||
// - constant (addrspace:4)
|
||||
// - local (addrspace:5)
|
||||
// - shared cluster (addrspace:7)
|
||||
// - entry parameter (addrspace:101)
|
||||
if (IsShortPtr)
|
||||
Ret += "-p3:32:32-p4:32:32-p5:32:32";
|
||||
|
||||
// Tensor Memory (addrspace:6) is always 32-bits.
|
||||
Ret += "-p6:32:32";
|
||||
|
||||
if (IsShortPtr)
|
||||
Ret += "-p7:32:32-p101:32:32";
|
||||
}
|
||||
|
||||
Ret += "-i64:64-i128:128-i256:256-v16:16-v32:32-n16:32:64";
|
||||
|
||||
return Ret;
|
||||
|
|
|
|||
|
|
@ -1,7 +1,9 @@
|
|||
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
|
||||
; RUN: opt < %s -S -nvptx-lower-args --mtriple nvptx64-nvidia-cuda -mcpu=sm_70 -mattr=+ptx77 | FileCheck %s --check-prefixes OPT
|
||||
; RUN: llc < %s --mtriple nvptx64-nvidia-cuda -mcpu=sm_70 -mattr=+ptx77 -O1 | FileCheck %s --check-prefixes PTX
|
||||
; RUN: opt < %s -S -nvptx-lower-args -mcpu=sm_70 -mattr=+ptx77 | FileCheck %s --check-prefixes OPT
|
||||
; RUN: llc < %s -mcpu=sm_70 -mattr=+ptx77 -O1 | FileCheck %s --check-prefixes PTX,PTX-DEFAULT
|
||||
; RUN: llc < %s -mcpu=sm_70 -mattr=+ptx77 -O1 --nvptx-short-ptr | FileCheck %s --check-prefixes PTX,PTX-SHORT-PTR
|
||||
|
||||
target triple = "nvptx64-nvidia-cuda"
|
||||
%struct.uint4 = type { i32, i32, i32, i32 }
|
||||
|
||||
@gi = dso_local addrspace(1) externally_initialized global %struct.uint4 { i32 50462976, i32 117835012, i32 185207048, i32 252579084 }, align 16
|
||||
|
|
@ -18,36 +20,68 @@ define dso_local noundef i32 @non_kernel_function(ptr nocapture noundef readonly
|
|||
; OPT-NEXT: [[TMP0:%.*]] = load i32, ptr [[ADD_PTR]], align 1
|
||||
; OPT-NEXT: ret i32 [[TMP0]]
|
||||
;
|
||||
; PTX-LABEL: non_kernel_function(
|
||||
; PTX: {
|
||||
; PTX-NEXT: .reg .pred %p<2>;
|
||||
; PTX-NEXT: .reg .b16 %rs<3>;
|
||||
; PTX-NEXT: .reg .b32 %r<11>;
|
||||
; PTX-NEXT: .reg .b64 %rd<8>;
|
||||
; PTX-EMPTY:
|
||||
; PTX-NEXT: // %bb.0: // %entry
|
||||
; PTX-NEXT: mov.b64 %rd1, non_kernel_function_param_0;
|
||||
; PTX-NEXT: cvta.local.u64 %rd2, %rd1;
|
||||
; PTX-NEXT: ld.param.b8 %rs1, [non_kernel_function_param_1];
|
||||
; PTX-NEXT: and.b16 %rs2, %rs1, 1;
|
||||
; PTX-NEXT: setp.ne.b16 %p1, %rs2, 0;
|
||||
; PTX-NEXT: mov.b64 %rd3, gi;
|
||||
; PTX-NEXT: cvta.global.u64 %rd4, %rd3;
|
||||
; PTX-NEXT: selp.b64 %rd5, %rd2, %rd4, %p1;
|
||||
; PTX-NEXT: ld.param.s32 %rd6, [non_kernel_function_param_2];
|
||||
; PTX-NEXT: add.s64 %rd7, %rd5, %rd6;
|
||||
; PTX-NEXT: ld.b8 %r1, [%rd7];
|
||||
; PTX-NEXT: ld.b8 %r2, [%rd7+1];
|
||||
; PTX-NEXT: shl.b32 %r3, %r2, 8;
|
||||
; PTX-NEXT: or.b32 %r4, %r3, %r1;
|
||||
; PTX-NEXT: ld.b8 %r5, [%rd7+2];
|
||||
; PTX-NEXT: shl.b32 %r6, %r5, 16;
|
||||
; PTX-NEXT: ld.b8 %r7, [%rd7+3];
|
||||
; PTX-NEXT: shl.b32 %r8, %r7, 24;
|
||||
; PTX-NEXT: or.b32 %r9, %r8, %r6;
|
||||
; PTX-NEXT: or.b32 %r10, %r9, %r4;
|
||||
; PTX-NEXT: st.param.b32 [func_retval0], %r10;
|
||||
; PTX-NEXT: ret;
|
||||
; PTX-DEFAULT-LABEL: non_kernel_function(
|
||||
; PTX-DEFAULT: {
|
||||
; PTX-DEFAULT-NEXT: .reg .pred %p<2>;
|
||||
; PTX-DEFAULT-NEXT: .reg .b16 %rs<3>;
|
||||
; PTX-DEFAULT-NEXT: .reg .b32 %r<11>;
|
||||
; PTX-DEFAULT-NEXT: .reg .b64 %rd<8>;
|
||||
; PTX-DEFAULT-EMPTY:
|
||||
; PTX-DEFAULT-NEXT: // %bb.0: // %entry
|
||||
; PTX-DEFAULT-NEXT: mov.b64 %rd1, non_kernel_function_param_0;
|
||||
; PTX-DEFAULT-NEXT: cvta.local.u64 %rd2, %rd1;
|
||||
; PTX-DEFAULT-NEXT: ld.param.b8 %rs1, [non_kernel_function_param_1];
|
||||
; PTX-DEFAULT-NEXT: and.b16 %rs2, %rs1, 1;
|
||||
; PTX-DEFAULT-NEXT: setp.ne.b16 %p1, %rs2, 0;
|
||||
; PTX-DEFAULT-NEXT: mov.b64 %rd3, gi;
|
||||
; PTX-DEFAULT-NEXT: cvta.global.u64 %rd4, %rd3;
|
||||
; PTX-DEFAULT-NEXT: selp.b64 %rd5, %rd2, %rd4, %p1;
|
||||
; PTX-DEFAULT-NEXT: ld.param.s32 %rd6, [non_kernel_function_param_2];
|
||||
; PTX-DEFAULT-NEXT: add.s64 %rd7, %rd5, %rd6;
|
||||
; PTX-DEFAULT-NEXT: ld.b8 %r1, [%rd7];
|
||||
; PTX-DEFAULT-NEXT: ld.b8 %r2, [%rd7+1];
|
||||
; PTX-DEFAULT-NEXT: shl.b32 %r3, %r2, 8;
|
||||
; PTX-DEFAULT-NEXT: or.b32 %r4, %r3, %r1;
|
||||
; PTX-DEFAULT-NEXT: ld.b8 %r5, [%rd7+2];
|
||||
; PTX-DEFAULT-NEXT: shl.b32 %r6, %r5, 16;
|
||||
; PTX-DEFAULT-NEXT: ld.b8 %r7, [%rd7+3];
|
||||
; PTX-DEFAULT-NEXT: shl.b32 %r8, %r7, 24;
|
||||
; PTX-DEFAULT-NEXT: or.b32 %r9, %r8, %r6;
|
||||
; PTX-DEFAULT-NEXT: or.b32 %r10, %r9, %r4;
|
||||
; PTX-DEFAULT-NEXT: st.param.b32 [func_retval0], %r10;
|
||||
; PTX-DEFAULT-NEXT: ret;
|
||||
;
|
||||
; PTX-SHORT-PTR-LABEL: non_kernel_function(
|
||||
; PTX-SHORT-PTR: {
|
||||
; PTX-SHORT-PTR-NEXT: .reg .pred %p<2>;
|
||||
; PTX-SHORT-PTR-NEXT: .reg .b16 %rs<3>;
|
||||
; PTX-SHORT-PTR-NEXT: .reg .b32 %r<12>;
|
||||
; PTX-SHORT-PTR-NEXT: .reg .b64 %rd<8>;
|
||||
; PTX-SHORT-PTR-EMPTY:
|
||||
; PTX-SHORT-PTR-NEXT: // %bb.0: // %entry
|
||||
; PTX-SHORT-PTR-NEXT: ld.param.b8 %rs1, [non_kernel_function_param_1];
|
||||
; PTX-SHORT-PTR-NEXT: and.b16 %rs2, %rs1, 1;
|
||||
; PTX-SHORT-PTR-NEXT: setp.ne.b16 %p1, %rs2, 0;
|
||||
; PTX-SHORT-PTR-NEXT: mov.b32 %r1, non_kernel_function_param_0;
|
||||
; PTX-SHORT-PTR-NEXT: cvt.u64.u32 %rd1, %r1;
|
||||
; PTX-SHORT-PTR-NEXT: cvta.local.u64 %rd2, %rd1;
|
||||
; PTX-SHORT-PTR-NEXT: mov.b64 %rd3, gi;
|
||||
; PTX-SHORT-PTR-NEXT: cvta.global.u64 %rd4, %rd3;
|
||||
; PTX-SHORT-PTR-NEXT: selp.b64 %rd5, %rd2, %rd4, %p1;
|
||||
; PTX-SHORT-PTR-NEXT: ld.param.s32 %rd6, [non_kernel_function_param_2];
|
||||
; PTX-SHORT-PTR-NEXT: add.s64 %rd7, %rd5, %rd6;
|
||||
; PTX-SHORT-PTR-NEXT: ld.b8 %r2, [%rd7];
|
||||
; PTX-SHORT-PTR-NEXT: ld.b8 %r3, [%rd7+1];
|
||||
; PTX-SHORT-PTR-NEXT: shl.b32 %r4, %r3, 8;
|
||||
; PTX-SHORT-PTR-NEXT: or.b32 %r5, %r4, %r2;
|
||||
; PTX-SHORT-PTR-NEXT: ld.b8 %r6, [%rd7+2];
|
||||
; PTX-SHORT-PTR-NEXT: shl.b32 %r7, %r6, 16;
|
||||
; PTX-SHORT-PTR-NEXT: ld.b8 %r8, [%rd7+3];
|
||||
; PTX-SHORT-PTR-NEXT: shl.b32 %r9, %r8, 24;
|
||||
; PTX-SHORT-PTR-NEXT: or.b32 %r10, %r9, %r7;
|
||||
; PTX-SHORT-PTR-NEXT: or.b32 %r11, %r10, %r5;
|
||||
; PTX-SHORT-PTR-NEXT: st.param.b32 [func_retval0], %r11;
|
||||
; PTX-SHORT-PTR-NEXT: ret;
|
||||
entry:
|
||||
%a. = select i1 %b, ptr %a, ptr addrspacecast (ptr addrspace(1) @gi to ptr)
|
||||
%idx.ext = sext i32 %c to i64
|
||||
|
|
@ -76,7 +110,6 @@ define ptx_kernel void @grid_const_int(ptr byval(i32) align 4 "nvvm.grid_constan
|
|||
; OPT-NEXT: [[ADD:%.*]] = add i32 [[TMP]], [[INPUT2]]
|
||||
; OPT-NEXT: store i32 [[ADD]], ptr [[OUT]], align 4
|
||||
; OPT-NEXT: ret void
|
||||
;
|
||||
%tmp = load i32, ptr %input1, align 4
|
||||
%add = add i32 %tmp, %input2
|
||||
store i32 %add, ptr %out
|
||||
|
|
@ -108,7 +141,6 @@ define ptx_kernel void @grid_const_struct(ptr byval(%struct.s) align 4 "nvvm.gri
|
|||
; OPT-NEXT: [[ADD:%.*]] = add i32 [[TMP1]], [[TMP2]]
|
||||
; OPT-NEXT: store i32 [[ADD]], ptr [[OUT]], align 4
|
||||
; OPT-NEXT: ret void
|
||||
;
|
||||
%gep1 = getelementptr inbounds %struct.s, ptr %input, i32 0, i32 0
|
||||
%gep2 = getelementptr inbounds %struct.s, ptr %input, i32 0, i32 1
|
||||
%int1 = load i32, ptr %gep1
|
||||
|
|
@ -119,63 +151,115 @@ define ptx_kernel void @grid_const_struct(ptr byval(%struct.s) align 4 "nvvm.gri
|
|||
}
|
||||
|
||||
define ptx_kernel void @grid_const_escape(ptr byval(%struct.s) align 4 "nvvm.grid_constant" %input) {
|
||||
; PTX-LABEL: grid_const_escape(
|
||||
; PTX: {
|
||||
; PTX-NEXT: .reg .b64 %rd<4>;
|
||||
; PTX-NEXT: prototype_0 : .callprototype (.param .b32 _) _ (.param .b64 _);
|
||||
; PTX-NEXT: // %bb.0:
|
||||
; PTX-NEXT: mov.b64 %rd1, grid_const_escape_param_0;
|
||||
; PTX-NEXT: cvta.param.u64 %rd2, %rd1;
|
||||
; PTX-NEXT: { // callseq 0, 0
|
||||
; PTX-NEXT: .param .b64 param0;
|
||||
; PTX-NEXT: .param .b32 retval0;
|
||||
; PTX-NEXT: st.param.b64 [param0], %rd2;
|
||||
; PTX-NEXT: mov.b64 %rd3, escape;
|
||||
; PTX-NEXT: call (retval0), %rd3, (param0), prototype_0;
|
||||
; PTX-NEXT: } // callseq 0
|
||||
; PTX-NEXT: ret;
|
||||
; PTX-DEFAULT-LABEL: grid_const_escape(
|
||||
; PTX-DEFAULT: {
|
||||
; PTX-DEFAULT-NEXT: .reg .b64 %rd<4>;
|
||||
; PTX-DEFAULT-NEXT: prototype_0 : .callprototype (.param .b32 _) _ (.param .b64 _);
|
||||
; PTX-DEFAULT-NEXT: // %bb.0:
|
||||
; PTX-DEFAULT-NEXT: mov.b64 %rd1, grid_const_escape_param_0;
|
||||
; PTX-DEFAULT-NEXT: cvta.param.u64 %rd2, %rd1;
|
||||
; PTX-DEFAULT-NEXT: { // callseq 0, 0
|
||||
; PTX-DEFAULT-NEXT: .param .b64 param0;
|
||||
; PTX-DEFAULT-NEXT: .param .b32 retval0;
|
||||
; PTX-DEFAULT-NEXT: st.param.b64 [param0], %rd2;
|
||||
; PTX-DEFAULT-NEXT: mov.b64 %rd3, escape;
|
||||
; PTX-DEFAULT-NEXT: call (retval0), %rd3, (param0), prototype_0;
|
||||
; PTX-DEFAULT-NEXT: } // callseq 0
|
||||
; PTX-DEFAULT-NEXT: ret;
|
||||
;
|
||||
; PTX-SHORT-PTR-LABEL: grid_const_escape(
|
||||
; PTX-SHORT-PTR: {
|
||||
; PTX-SHORT-PTR-NEXT: .reg .b32 %r<2>;
|
||||
; PTX-SHORT-PTR-NEXT: .reg .b64 %rd<4>;
|
||||
; PTX-SHORT-PTR-NEXT: prototype_0 : .callprototype (.param .b32 _) _ (.param .b64 _);
|
||||
; PTX-SHORT-PTR-NEXT: // %bb.0:
|
||||
; PTX-SHORT-PTR-NEXT: mov.b32 %r1, grid_const_escape_param_0;
|
||||
; PTX-SHORT-PTR-NEXT: cvt.u64.u32 %rd1, %r1;
|
||||
; PTX-SHORT-PTR-NEXT: cvta.param.u64 %rd2, %rd1;
|
||||
; PTX-SHORT-PTR-NEXT: { // callseq 0, 0
|
||||
; PTX-SHORT-PTR-NEXT: .param .b64 param0;
|
||||
; PTX-SHORT-PTR-NEXT: .param .b32 retval0;
|
||||
; PTX-SHORT-PTR-NEXT: st.param.b64 [param0], %rd2;
|
||||
; PTX-SHORT-PTR-NEXT: mov.b64 %rd3, escape;
|
||||
; PTX-SHORT-PTR-NEXT: call (retval0), %rd3, (param0), prototype_0;
|
||||
; PTX-SHORT-PTR-NEXT: } // callseq 0
|
||||
; PTX-SHORT-PTR-NEXT: ret;
|
||||
; OPT-LABEL: define ptx_kernel void @grid_const_escape(
|
||||
; OPT-SAME: ptr addrspace(101) readonly byval([[STRUCT_S:%.*]]) align 4 "nvvm.grid_constant" [[INPUT:%.*]]) #[[ATTR0]] {
|
||||
; OPT-NEXT: [[INPUT_PARAM_GEN:%.*]] = addrspacecast ptr addrspace(101) [[INPUT]] to ptr
|
||||
; OPT-NEXT: [[CALL:%.*]] = call i32 @escape(ptr [[INPUT_PARAM_GEN]])
|
||||
; OPT-NEXT: ret void
|
||||
;
|
||||
%call = call i32 @escape(ptr %input)
|
||||
ret void
|
||||
}
|
||||
|
||||
define ptx_kernel void @multiple_grid_const_escape(ptr byval(%struct.s) align 4 "nvvm.grid_constant" %input, i32 %a, ptr byval(i32) align 4 "nvvm.grid_constant" %b) {
|
||||
; PTX-LABEL: multiple_grid_const_escape(
|
||||
; PTX: {
|
||||
; PTX-NEXT: .local .align 4 .b8 __local_depot4[4];
|
||||
; PTX-NEXT: .reg .b64 %SP;
|
||||
; PTX-NEXT: .reg .b64 %SPL;
|
||||
; PTX-NEXT: .reg .b32 %r<2>;
|
||||
; PTX-NEXT: .reg .b64 %rd<8>;
|
||||
; PTX-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .b64 _, .param .b64 _, .param .b64 _);
|
||||
; PTX-NEXT: // %bb.0:
|
||||
; PTX-NEXT: mov.b64 %SPL, __local_depot4;
|
||||
; PTX-NEXT: cvta.local.u64 %SP, %SPL;
|
||||
; PTX-NEXT: mov.b64 %rd1, multiple_grid_const_escape_param_0;
|
||||
; PTX-NEXT: mov.b64 %rd2, multiple_grid_const_escape_param_2;
|
||||
; PTX-NEXT: cvta.param.u64 %rd3, %rd2;
|
||||
; PTX-NEXT: ld.param.b32 %r1, [multiple_grid_const_escape_param_1];
|
||||
; PTX-NEXT: cvta.param.u64 %rd4, %rd1;
|
||||
; PTX-NEXT: add.u64 %rd5, %SP, 0;
|
||||
; PTX-NEXT: add.u64 %rd6, %SPL, 0;
|
||||
; PTX-NEXT: st.local.b32 [%rd6], %r1;
|
||||
; PTX-NEXT: { // callseq 1, 0
|
||||
; PTX-NEXT: .param .b64 param0;
|
||||
; PTX-NEXT: .param .b64 param1;
|
||||
; PTX-NEXT: .param .b64 param2;
|
||||
; PTX-NEXT: .param .b32 retval0;
|
||||
; PTX-NEXT: st.param.b64 [param2], %rd3;
|
||||
; PTX-NEXT: st.param.b64 [param1], %rd5;
|
||||
; PTX-NEXT: st.param.b64 [param0], %rd4;
|
||||
; PTX-NEXT: mov.b64 %rd7, escape3;
|
||||
; PTX-NEXT: call (retval0), %rd7, (param0, param1, param2), prototype_1;
|
||||
; PTX-NEXT: } // callseq 1
|
||||
; PTX-NEXT: ret;
|
||||
; PTX-DEFAULT-LABEL: multiple_grid_const_escape(
|
||||
; PTX-DEFAULT: {
|
||||
; PTX-DEFAULT-NEXT: .local .align 4 .b8 __local_depot4[4];
|
||||
; PTX-DEFAULT-NEXT: .reg .b64 %SP;
|
||||
; PTX-DEFAULT-NEXT: .reg .b64 %SPL;
|
||||
; PTX-DEFAULT-NEXT: .reg .b32 %r<2>;
|
||||
; PTX-DEFAULT-NEXT: .reg .b64 %rd<8>;
|
||||
; PTX-DEFAULT-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .b64 _, .param .b64 _, .param .b64 _);
|
||||
; PTX-DEFAULT-NEXT: // %bb.0:
|
||||
; PTX-DEFAULT-NEXT: mov.b64 %SPL, __local_depot4;
|
||||
; PTX-DEFAULT-NEXT: cvta.local.u64 %SP, %SPL;
|
||||
; PTX-DEFAULT-NEXT: mov.b64 %rd1, multiple_grid_const_escape_param_0;
|
||||
; PTX-DEFAULT-NEXT: mov.b64 %rd2, multiple_grid_const_escape_param_2;
|
||||
; PTX-DEFAULT-NEXT: cvta.param.u64 %rd3, %rd2;
|
||||
; PTX-DEFAULT-NEXT: ld.param.b32 %r1, [multiple_grid_const_escape_param_1];
|
||||
; PTX-DEFAULT-NEXT: cvta.param.u64 %rd4, %rd1;
|
||||
; PTX-DEFAULT-NEXT: add.u64 %rd5, %SP, 0;
|
||||
; PTX-DEFAULT-NEXT: add.u64 %rd6, %SPL, 0;
|
||||
; PTX-DEFAULT-NEXT: st.local.b32 [%rd6], %r1;
|
||||
; PTX-DEFAULT-NEXT: { // callseq 1, 0
|
||||
; PTX-DEFAULT-NEXT: .param .b64 param0;
|
||||
; PTX-DEFAULT-NEXT: .param .b64 param1;
|
||||
; PTX-DEFAULT-NEXT: .param .b64 param2;
|
||||
; PTX-DEFAULT-NEXT: .param .b32 retval0;
|
||||
; PTX-DEFAULT-NEXT: st.param.b64 [param2], %rd3;
|
||||
; PTX-DEFAULT-NEXT: st.param.b64 [param1], %rd5;
|
||||
; PTX-DEFAULT-NEXT: st.param.b64 [param0], %rd4;
|
||||
; PTX-DEFAULT-NEXT: mov.b64 %rd7, escape3;
|
||||
; PTX-DEFAULT-NEXT: call (retval0), %rd7, (param0, param1, param2), prototype_1;
|
||||
; PTX-DEFAULT-NEXT: } // callseq 1
|
||||
; PTX-DEFAULT-NEXT: ret;
|
||||
;
|
||||
; PTX-SHORT-PTR-LABEL: multiple_grid_const_escape(
|
||||
; PTX-SHORT-PTR: {
|
||||
; PTX-SHORT-PTR-NEXT: .local .align 4 .b8 __local_depot4[4];
|
||||
; PTX-SHORT-PTR-NEXT: .reg .b64 %SP;
|
||||
; PTX-SHORT-PTR-NEXT: .reg .b64 %SPL;
|
||||
; PTX-SHORT-PTR-NEXT: .reg .b32 %r<5>;
|
||||
; PTX-SHORT-PTR-NEXT: .reg .b64 %rd<8>;
|
||||
; PTX-SHORT-PTR-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .b64 _, .param .b64 _, .param .b64 _);
|
||||
; PTX-SHORT-PTR-NEXT: // %bb.0:
|
||||
; PTX-SHORT-PTR-NEXT: mov.b64 %SPL, __local_depot4;
|
||||
; PTX-SHORT-PTR-NEXT: cvta.local.u64 %SP, %SPL;
|
||||
; PTX-SHORT-PTR-NEXT: mov.b32 %r1, multiple_grid_const_escape_param_0;
|
||||
; PTX-SHORT-PTR-NEXT: mov.b32 %r2, multiple_grid_const_escape_param_2;
|
||||
; PTX-SHORT-PTR-NEXT: cvt.u64.u32 %rd1, %r2;
|
||||
; PTX-SHORT-PTR-NEXT: cvta.param.u64 %rd2, %rd1;
|
||||
; PTX-SHORT-PTR-NEXT: ld.param.b32 %r3, [multiple_grid_const_escape_param_1];
|
||||
; PTX-SHORT-PTR-NEXT: cvt.u64.u32 %rd3, %r1;
|
||||
; PTX-SHORT-PTR-NEXT: cvta.param.u64 %rd4, %rd3;
|
||||
; PTX-SHORT-PTR-NEXT: add.u64 %rd5, %SP, 0;
|
||||
; PTX-SHORT-PTR-NEXT: add.u64 %rd6, %SPL, 0;
|
||||
; PTX-SHORT-PTR-NEXT: cvt.u32.u64 %r4, %rd6;
|
||||
; PTX-SHORT-PTR-NEXT: st.local.b32 [%r4], %r3;
|
||||
; PTX-SHORT-PTR-NEXT: { // callseq 1, 0
|
||||
; PTX-SHORT-PTR-NEXT: .param .b64 param0;
|
||||
; PTX-SHORT-PTR-NEXT: .param .b64 param1;
|
||||
; PTX-SHORT-PTR-NEXT: .param .b64 param2;
|
||||
; PTX-SHORT-PTR-NEXT: .param .b32 retval0;
|
||||
; PTX-SHORT-PTR-NEXT: st.param.b64 [param2], %rd2;
|
||||
; PTX-SHORT-PTR-NEXT: st.param.b64 [param1], %rd5;
|
||||
; PTX-SHORT-PTR-NEXT: st.param.b64 [param0], %rd4;
|
||||
; PTX-SHORT-PTR-NEXT: mov.b64 %rd7, escape3;
|
||||
; PTX-SHORT-PTR-NEXT: call (retval0), %rd7, (param0, param1, param2), prototype_1;
|
||||
; PTX-SHORT-PTR-NEXT: } // callseq 1
|
||||
; PTX-SHORT-PTR-NEXT: ret;
|
||||
; OPT-LABEL: define ptx_kernel void @multiple_grid_const_escape(
|
||||
; OPT-SAME: ptr addrspace(101) readonly byval([[STRUCT_S:%.*]]) align 4 "nvvm.grid_constant" [[INPUT:%.*]], i32 [[A:%.*]], ptr addrspace(101) readonly byval(i32) align 4 "nvvm.grid_constant" [[B:%.*]]) #[[ATTR0]] {
|
||||
; OPT-NEXT: [[B_PARAM_GEN:%.*]] = addrspacecast ptr addrspace(101) [[B]] to ptr
|
||||
|
|
@ -184,7 +268,6 @@ define ptx_kernel void @multiple_grid_const_escape(ptr byval(%struct.s) align 4
|
|||
; OPT-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 4
|
||||
; OPT-NEXT: [[CALL:%.*]] = call i32 @escape3(ptr [[INPUT_PARAM_GEN]], ptr [[A_ADDR]], ptr [[B_PARAM_GEN]])
|
||||
; OPT-NEXT: ret void
|
||||
;
|
||||
%a.addr = alloca i32, align 4
|
||||
store i32 %a, ptr %a.addr, align 4
|
||||
%call = call i32 @escape3(ptr %input, ptr %a.addr, ptr %b)
|
||||
|
|
@ -192,43 +275,74 @@ define ptx_kernel void @multiple_grid_const_escape(ptr byval(%struct.s) align 4
|
|||
}
|
||||
|
||||
define ptx_kernel void @grid_const_memory_escape(ptr byval(%struct.s) align 4 "nvvm.grid_constant" %input, ptr %addr) {
|
||||
; PTX-LABEL: grid_const_memory_escape(
|
||||
; PTX: {
|
||||
; PTX-NEXT: .reg .b64 %rd<5>;
|
||||
; PTX-EMPTY:
|
||||
; PTX-NEXT: // %bb.0:
|
||||
; PTX-NEXT: mov.b64 %rd1, grid_const_memory_escape_param_0;
|
||||
; PTX-NEXT: ld.param.b64 %rd2, [grid_const_memory_escape_param_1];
|
||||
; PTX-NEXT: cvta.to.global.u64 %rd3, %rd2;
|
||||
; PTX-NEXT: cvta.param.u64 %rd4, %rd1;
|
||||
; PTX-NEXT: st.global.b64 [%rd3], %rd4;
|
||||
; PTX-NEXT: ret;
|
||||
; PTX-DEFAULT-LABEL: grid_const_memory_escape(
|
||||
; PTX-DEFAULT: {
|
||||
; PTX-DEFAULT-NEXT: .reg .b64 %rd<5>;
|
||||
; PTX-DEFAULT-EMPTY:
|
||||
; PTX-DEFAULT-NEXT: // %bb.0:
|
||||
; PTX-DEFAULT-NEXT: mov.b64 %rd1, grid_const_memory_escape_param_0;
|
||||
; PTX-DEFAULT-NEXT: ld.param.b64 %rd2, [grid_const_memory_escape_param_1];
|
||||
; PTX-DEFAULT-NEXT: cvta.to.global.u64 %rd3, %rd2;
|
||||
; PTX-DEFAULT-NEXT: cvta.param.u64 %rd4, %rd1;
|
||||
; PTX-DEFAULT-NEXT: st.global.b64 [%rd3], %rd4;
|
||||
; PTX-DEFAULT-NEXT: ret;
|
||||
;
|
||||
; PTX-SHORT-PTR-LABEL: grid_const_memory_escape(
|
||||
; PTX-SHORT-PTR: {
|
||||
; PTX-SHORT-PTR-NEXT: .reg .b32 %r<2>;
|
||||
; PTX-SHORT-PTR-NEXT: .reg .b64 %rd<5>;
|
||||
; PTX-SHORT-PTR-EMPTY:
|
||||
; PTX-SHORT-PTR-NEXT: // %bb.0:
|
||||
; PTX-SHORT-PTR-NEXT: mov.b32 %r1, grid_const_memory_escape_param_0;
|
||||
; PTX-SHORT-PTR-NEXT: ld.param.b64 %rd1, [grid_const_memory_escape_param_1];
|
||||
; PTX-SHORT-PTR-NEXT: cvta.to.global.u64 %rd2, %rd1;
|
||||
; PTX-SHORT-PTR-NEXT: cvt.u64.u32 %rd3, %r1;
|
||||
; PTX-SHORT-PTR-NEXT: cvta.param.u64 %rd4, %rd3;
|
||||
; PTX-SHORT-PTR-NEXT: st.global.b64 [%rd2], %rd4;
|
||||
; PTX-SHORT-PTR-NEXT: ret;
|
||||
; OPT-LABEL: define ptx_kernel void @grid_const_memory_escape(
|
||||
; OPT-SAME: ptr addrspace(101) readonly byval([[STRUCT_S:%.*]]) align 4 "nvvm.grid_constant" [[INPUT:%.*]], ptr [[ADDR:%.*]]) #[[ATTR0]] {
|
||||
; OPT-NEXT: [[INPUT1:%.*]] = addrspacecast ptr addrspace(101) [[INPUT]] to ptr
|
||||
; OPT-NEXT: store ptr [[INPUT1]], ptr [[ADDR]], align 8
|
||||
; OPT-NEXT: ret void
|
||||
;
|
||||
store ptr %input, ptr %addr, align 8
|
||||
ret void
|
||||
}
|
||||
|
||||
define ptx_kernel void @grid_const_inlineasm_escape(ptr byval(%struct.s) align 4 "nvvm.grid_constant" %input, ptr %result) {
|
||||
; PTX-LABEL: grid_const_inlineasm_escape(
|
||||
; PTX: {
|
||||
; PTX-NEXT: .reg .b64 %rd<7>;
|
||||
; PTX-EMPTY:
|
||||
; PTX-NEXT: // %bb.0:
|
||||
; PTX-NEXT: mov.b64 %rd4, grid_const_inlineasm_escape_param_0;
|
||||
; PTX-NEXT: ld.param.b64 %rd5, [grid_const_inlineasm_escape_param_1];
|
||||
; PTX-NEXT: cvta.to.global.u64 %rd6, %rd5;
|
||||
; PTX-NEXT: cvta.param.u64 %rd2, %rd4;
|
||||
; PTX-NEXT: add.s64 %rd3, %rd2, 4;
|
||||
; PTX-NEXT: // begin inline asm
|
||||
; PTX-NEXT: add.s64 %rd1, %rd2, %rd3;
|
||||
; PTX-NEXT: // end inline asm
|
||||
; PTX-NEXT: st.global.b64 [%rd6], %rd1;
|
||||
; PTX-NEXT: ret;
|
||||
; PTX-DEFAULT-LABEL: grid_const_inlineasm_escape(
|
||||
; PTX-DEFAULT: {
|
||||
; PTX-DEFAULT-NEXT: .reg .b64 %rd<7>;
|
||||
; PTX-DEFAULT-EMPTY:
|
||||
; PTX-DEFAULT-NEXT: // %bb.0:
|
||||
; PTX-DEFAULT-NEXT: mov.b64 %rd4, grid_const_inlineasm_escape_param_0;
|
||||
; PTX-DEFAULT-NEXT: ld.param.b64 %rd5, [grid_const_inlineasm_escape_param_1];
|
||||
; PTX-DEFAULT-NEXT: cvta.to.global.u64 %rd6, %rd5;
|
||||
; PTX-DEFAULT-NEXT: cvta.param.u64 %rd2, %rd4;
|
||||
; PTX-DEFAULT-NEXT: add.s64 %rd3, %rd2, 4;
|
||||
; PTX-DEFAULT-NEXT: // begin inline asm
|
||||
; PTX-DEFAULT-NEXT: add.s64 %rd1, %rd2, %rd3;
|
||||
; PTX-DEFAULT-NEXT: // end inline asm
|
||||
; PTX-DEFAULT-NEXT: st.global.b64 [%rd6], %rd1;
|
||||
; PTX-DEFAULT-NEXT: ret;
|
||||
;
|
||||
; PTX-SHORT-PTR-LABEL: grid_const_inlineasm_escape(
|
||||
; PTX-SHORT-PTR: {
|
||||
; PTX-SHORT-PTR-NEXT: .reg .b32 %r<2>;
|
||||
; PTX-SHORT-PTR-NEXT: .reg .b64 %rd<7>;
|
||||
; PTX-SHORT-PTR-EMPTY:
|
||||
; PTX-SHORT-PTR-NEXT: // %bb.0:
|
||||
; PTX-SHORT-PTR-NEXT: mov.b32 %r1, grid_const_inlineasm_escape_param_0;
|
||||
; PTX-SHORT-PTR-NEXT: ld.param.b64 %rd4, [grid_const_inlineasm_escape_param_1];
|
||||
; PTX-SHORT-PTR-NEXT: cvta.to.global.u64 %rd5, %rd4;
|
||||
; PTX-SHORT-PTR-NEXT: cvt.u64.u32 %rd6, %r1;
|
||||
; PTX-SHORT-PTR-NEXT: cvta.param.u64 %rd2, %rd6;
|
||||
; PTX-SHORT-PTR-NEXT: add.s64 %rd3, %rd2, 4;
|
||||
; PTX-SHORT-PTR-NEXT: // begin inline asm
|
||||
; PTX-SHORT-PTR-NEXT: add.s64 %rd1, %rd2, %rd3;
|
||||
; PTX-SHORT-PTR-NEXT: // end inline asm
|
||||
; PTX-SHORT-PTR-NEXT: st.global.b64 [%rd5], %rd1;
|
||||
; PTX-SHORT-PTR-NEXT: ret;
|
||||
; PTX-NOT .local
|
||||
; OPT-LABEL: define ptx_kernel void @grid_const_inlineasm_escape(
|
||||
; OPT-SAME: ptr addrspace(101) readonly byval([[STRUCT_S:%.*]]) align 4 "nvvm.grid_constant" [[INPUT:%.*]], ptr [[RESULT:%.*]]) #[[ATTR0]] {
|
||||
|
|
@ -238,7 +352,6 @@ define ptx_kernel void @grid_const_inlineasm_escape(ptr byval(%struct.s) align 4
|
|||
; OPT-NEXT: [[TMP2:%.*]] = call i64 asm "add.s64 $0, $1, $2
|
||||
; OPT-NEXT: store i64 [[TMP2]], ptr [[RESULT]], align 8
|
||||
; OPT-NEXT: ret void
|
||||
;
|
||||
%tmpptr1 = getelementptr inbounds %struct.s, ptr %input, i32 0, i32 0
|
||||
%tmpptr2 = getelementptr inbounds %struct.s, ptr %input, i32 0, i32 1
|
||||
%1 = call i64 asm "add.s64 $0, $1, $2;", "=l,l,l"(ptr %tmpptr1, ptr %tmpptr2) #1
|
||||
|
|
@ -247,27 +360,50 @@ define ptx_kernel void @grid_const_inlineasm_escape(ptr byval(%struct.s) align 4
|
|||
}
|
||||
|
||||
define ptx_kernel void @grid_const_partial_escape(ptr byval(i32) align 4 "nvvm.grid_constant" %input, ptr %output) {
|
||||
; PTX-LABEL: grid_const_partial_escape(
|
||||
; PTX: {
|
||||
; PTX-NEXT: .reg .b32 %r<3>;
|
||||
; PTX-NEXT: .reg .b64 %rd<6>;
|
||||
; PTX-NEXT: prototype_2 : .callprototype (.param .b32 _) _ (.param .b64 _);
|
||||
; PTX-NEXT: // %bb.0:
|
||||
; PTX-NEXT: mov.b64 %rd1, grid_const_partial_escape_param_0;
|
||||
; PTX-NEXT: ld.param.b64 %rd2, [grid_const_partial_escape_param_1];
|
||||
; PTX-NEXT: cvta.to.global.u64 %rd3, %rd2;
|
||||
; PTX-NEXT: cvta.param.u64 %rd4, %rd1;
|
||||
; PTX-NEXT: ld.param.b32 %r1, [grid_const_partial_escape_param_0];
|
||||
; PTX-NEXT: add.s32 %r2, %r1, %r1;
|
||||
; PTX-NEXT: st.global.b32 [%rd3], %r2;
|
||||
; PTX-NEXT: { // callseq 2, 0
|
||||
; PTX-NEXT: .param .b64 param0;
|
||||
; PTX-NEXT: .param .b32 retval0;
|
||||
; PTX-NEXT: st.param.b64 [param0], %rd4;
|
||||
; PTX-NEXT: mov.b64 %rd5, escape;
|
||||
; PTX-NEXT: call (retval0), %rd5, (param0), prototype_2;
|
||||
; PTX-NEXT: } // callseq 2
|
||||
; PTX-NEXT: ret;
|
||||
; PTX-DEFAULT-LABEL: grid_const_partial_escape(
|
||||
; PTX-DEFAULT: {
|
||||
; PTX-DEFAULT-NEXT: .reg .b32 %r<3>;
|
||||
; PTX-DEFAULT-NEXT: .reg .b64 %rd<6>;
|
||||
; PTX-DEFAULT-NEXT: prototype_2 : .callprototype (.param .b32 _) _ (.param .b64 _);
|
||||
; PTX-DEFAULT-NEXT: // %bb.0:
|
||||
; PTX-DEFAULT-NEXT: mov.b64 %rd1, grid_const_partial_escape_param_0;
|
||||
; PTX-DEFAULT-NEXT: ld.param.b64 %rd2, [grid_const_partial_escape_param_1];
|
||||
; PTX-DEFAULT-NEXT: cvta.to.global.u64 %rd3, %rd2;
|
||||
; PTX-DEFAULT-NEXT: cvta.param.u64 %rd4, %rd1;
|
||||
; PTX-DEFAULT-NEXT: ld.param.b32 %r1, [grid_const_partial_escape_param_0];
|
||||
; PTX-DEFAULT-NEXT: add.s32 %r2, %r1, %r1;
|
||||
; PTX-DEFAULT-NEXT: st.global.b32 [%rd3], %r2;
|
||||
; PTX-DEFAULT-NEXT: { // callseq 2, 0
|
||||
; PTX-DEFAULT-NEXT: .param .b64 param0;
|
||||
; PTX-DEFAULT-NEXT: .param .b32 retval0;
|
||||
; PTX-DEFAULT-NEXT: st.param.b64 [param0], %rd4;
|
||||
; PTX-DEFAULT-NEXT: mov.b64 %rd5, escape;
|
||||
; PTX-DEFAULT-NEXT: call (retval0), %rd5, (param0), prototype_2;
|
||||
; PTX-DEFAULT-NEXT: } // callseq 2
|
||||
; PTX-DEFAULT-NEXT: ret;
|
||||
;
|
||||
; PTX-SHORT-PTR-LABEL: grid_const_partial_escape(
|
||||
; PTX-SHORT-PTR: {
|
||||
; PTX-SHORT-PTR-NEXT: .reg .b32 %r<4>;
|
||||
; PTX-SHORT-PTR-NEXT: .reg .b64 %rd<6>;
|
||||
; PTX-SHORT-PTR-NEXT: prototype_2 : .callprototype (.param .b32 _) _ (.param .b64 _);
|
||||
; PTX-SHORT-PTR-NEXT: // %bb.0:
|
||||
; PTX-SHORT-PTR-NEXT: mov.b32 %r1, grid_const_partial_escape_param_0;
|
||||
; PTX-SHORT-PTR-NEXT: ld.param.b64 %rd1, [grid_const_partial_escape_param_1];
|
||||
; PTX-SHORT-PTR-NEXT: cvta.to.global.u64 %rd2, %rd1;
|
||||
; PTX-SHORT-PTR-NEXT: cvt.u64.u32 %rd3, %r1;
|
||||
; PTX-SHORT-PTR-NEXT: cvta.param.u64 %rd4, %rd3;
|
||||
; PTX-SHORT-PTR-NEXT: ld.param.b32 %r2, [grid_const_partial_escape_param_0];
|
||||
; PTX-SHORT-PTR-NEXT: add.s32 %r3, %r2, %r2;
|
||||
; PTX-SHORT-PTR-NEXT: st.global.b32 [%rd2], %r3;
|
||||
; PTX-SHORT-PTR-NEXT: { // callseq 2, 0
|
||||
; PTX-SHORT-PTR-NEXT: .param .b64 param0;
|
||||
; PTX-SHORT-PTR-NEXT: .param .b32 retval0;
|
||||
; PTX-SHORT-PTR-NEXT: st.param.b64 [param0], %rd4;
|
||||
; PTX-SHORT-PTR-NEXT: mov.b64 %rd5, escape;
|
||||
; PTX-SHORT-PTR-NEXT: call (retval0), %rd5, (param0), prototype_2;
|
||||
; PTX-SHORT-PTR-NEXT: } // callseq 2
|
||||
; PTX-SHORT-PTR-NEXT: ret;
|
||||
; OPT-LABEL: define ptx_kernel void @grid_const_partial_escape(
|
||||
; OPT-SAME: ptr addrspace(101) readonly byval(i32) align 4 "nvvm.grid_constant" [[INPUT:%.*]], ptr [[OUTPUT:%.*]]) #[[ATTR0]] {
|
||||
; OPT-NEXT: [[INPUT1_GEN:%.*]] = addrspacecast ptr addrspace(101) [[INPUT]] to ptr
|
||||
|
|
@ -276,7 +412,6 @@ define ptx_kernel void @grid_const_partial_escape(ptr byval(i32) align 4 "nvvm.g
|
|||
; OPT-NEXT: store i32 [[TWICE]], ptr [[OUTPUT]], align 4
|
||||
; OPT-NEXT: [[CALL:%.*]] = call i32 @escape(ptr [[INPUT1_GEN]])
|
||||
; OPT-NEXT: ret void
|
||||
;
|
||||
%val = load i32, ptr %input
|
||||
%twice = add i32 %val, %val
|
||||
store i32 %twice, ptr %output
|
||||
|
|
@ -285,29 +420,54 @@ define ptx_kernel void @grid_const_partial_escape(ptr byval(i32) align 4 "nvvm.g
|
|||
}
|
||||
|
||||
define ptx_kernel i32 @grid_const_partial_escapemem(ptr byval(%struct.s) align 4 "nvvm.grid_constant" %input, ptr %output) {
|
||||
; PTX-LABEL: grid_const_partial_escapemem(
|
||||
; PTX: {
|
||||
; PTX-NEXT: .reg .b32 %r<4>;
|
||||
; PTX-NEXT: .reg .b64 %rd<6>;
|
||||
; PTX-NEXT: prototype_3 : .callprototype (.param .b32 _) _ (.param .b64 _);
|
||||
; PTX-NEXT: // %bb.0:
|
||||
; PTX-NEXT: mov.b64 %rd1, grid_const_partial_escapemem_param_0;
|
||||
; PTX-NEXT: ld.param.b64 %rd2, [grid_const_partial_escapemem_param_1];
|
||||
; PTX-NEXT: cvta.to.global.u64 %rd3, %rd2;
|
||||
; PTX-NEXT: cvta.param.u64 %rd4, %rd1;
|
||||
; PTX-NEXT: ld.param.b32 %r1, [grid_const_partial_escapemem_param_0];
|
||||
; PTX-NEXT: ld.param.b32 %r2, [grid_const_partial_escapemem_param_0+4];
|
||||
; PTX-NEXT: st.global.b64 [%rd3], %rd4;
|
||||
; PTX-NEXT: add.s32 %r3, %r1, %r2;
|
||||
; PTX-NEXT: { // callseq 3, 0
|
||||
; PTX-NEXT: .param .b64 param0;
|
||||
; PTX-NEXT: .param .b32 retval0;
|
||||
; PTX-NEXT: st.param.b64 [param0], %rd4;
|
||||
; PTX-NEXT: mov.b64 %rd5, escape;
|
||||
; PTX-NEXT: call (retval0), %rd5, (param0), prototype_3;
|
||||
; PTX-NEXT: } // callseq 3
|
||||
; PTX-NEXT: st.param.b32 [func_retval0], %r3;
|
||||
; PTX-NEXT: ret;
|
||||
; PTX-DEFAULT-LABEL: grid_const_partial_escapemem(
|
||||
; PTX-DEFAULT: {
|
||||
; PTX-DEFAULT-NEXT: .reg .b32 %r<4>;
|
||||
; PTX-DEFAULT-NEXT: .reg .b64 %rd<6>;
|
||||
; PTX-DEFAULT-NEXT: prototype_3 : .callprototype (.param .b32 _) _ (.param .b64 _);
|
||||
; PTX-DEFAULT-NEXT: // %bb.0:
|
||||
; PTX-DEFAULT-NEXT: mov.b64 %rd1, grid_const_partial_escapemem_param_0;
|
||||
; PTX-DEFAULT-NEXT: ld.param.b64 %rd2, [grid_const_partial_escapemem_param_1];
|
||||
; PTX-DEFAULT-NEXT: cvta.to.global.u64 %rd3, %rd2;
|
||||
; PTX-DEFAULT-NEXT: cvta.param.u64 %rd4, %rd1;
|
||||
; PTX-DEFAULT-NEXT: ld.param.b32 %r1, [grid_const_partial_escapemem_param_0];
|
||||
; PTX-DEFAULT-NEXT: ld.param.b32 %r2, [grid_const_partial_escapemem_param_0+4];
|
||||
; PTX-DEFAULT-NEXT: st.global.b64 [%rd3], %rd4;
|
||||
; PTX-DEFAULT-NEXT: add.s32 %r3, %r1, %r2;
|
||||
; PTX-DEFAULT-NEXT: { // callseq 3, 0
|
||||
; PTX-DEFAULT-NEXT: .param .b64 param0;
|
||||
; PTX-DEFAULT-NEXT: .param .b32 retval0;
|
||||
; PTX-DEFAULT-NEXT: st.param.b64 [param0], %rd4;
|
||||
; PTX-DEFAULT-NEXT: mov.b64 %rd5, escape;
|
||||
; PTX-DEFAULT-NEXT: call (retval0), %rd5, (param0), prototype_3;
|
||||
; PTX-DEFAULT-NEXT: } // callseq 3
|
||||
; PTX-DEFAULT-NEXT: st.param.b32 [func_retval0], %r3;
|
||||
; PTX-DEFAULT-NEXT: ret;
|
||||
;
|
||||
; PTX-SHORT-PTR-LABEL: grid_const_partial_escapemem(
|
||||
; PTX-SHORT-PTR: {
|
||||
; PTX-SHORT-PTR-NEXT: .reg .b32 %r<5>;
|
||||
; PTX-SHORT-PTR-NEXT: .reg .b64 %rd<6>;
|
||||
; PTX-SHORT-PTR-NEXT: prototype_3 : .callprototype (.param .b32 _) _ (.param .b64 _);
|
||||
; PTX-SHORT-PTR-NEXT: // %bb.0:
|
||||
; PTX-SHORT-PTR-NEXT: mov.b32 %r1, grid_const_partial_escapemem_param_0;
|
||||
; PTX-SHORT-PTR-NEXT: ld.param.b64 %rd1, [grid_const_partial_escapemem_param_1];
|
||||
; PTX-SHORT-PTR-NEXT: cvta.to.global.u64 %rd2, %rd1;
|
||||
; PTX-SHORT-PTR-NEXT: cvt.u64.u32 %rd3, %r1;
|
||||
; PTX-SHORT-PTR-NEXT: cvta.param.u64 %rd4, %rd3;
|
||||
; PTX-SHORT-PTR-NEXT: ld.param.b32 %r2, [grid_const_partial_escapemem_param_0];
|
||||
; PTX-SHORT-PTR-NEXT: ld.param.b32 %r3, [grid_const_partial_escapemem_param_0+4];
|
||||
; PTX-SHORT-PTR-NEXT: st.global.b64 [%rd2], %rd4;
|
||||
; PTX-SHORT-PTR-NEXT: add.s32 %r4, %r2, %r3;
|
||||
; PTX-SHORT-PTR-NEXT: { // callseq 3, 0
|
||||
; PTX-SHORT-PTR-NEXT: .param .b64 param0;
|
||||
; PTX-SHORT-PTR-NEXT: .param .b32 retval0;
|
||||
; PTX-SHORT-PTR-NEXT: st.param.b64 [param0], %rd4;
|
||||
; PTX-SHORT-PTR-NEXT: mov.b64 %rd5, escape;
|
||||
; PTX-SHORT-PTR-NEXT: call (retval0), %rd5, (param0), prototype_3;
|
||||
; PTX-SHORT-PTR-NEXT: } // callseq 3
|
||||
; PTX-SHORT-PTR-NEXT: st.param.b32 [func_retval0], %r4;
|
||||
; PTX-SHORT-PTR-NEXT: ret;
|
||||
; OPT-LABEL: define ptx_kernel i32 @grid_const_partial_escapemem(
|
||||
; OPT-SAME: ptr addrspace(101) readonly byval([[STRUCT_S:%.*]]) align 4 "nvvm.grid_constant" [[INPUT:%.*]], ptr [[OUTPUT:%.*]]) #[[ATTR0]] {
|
||||
; OPT-NEXT: [[INPUT1:%.*]] = addrspacecast ptr addrspace(101) [[INPUT]] to ptr
|
||||
|
|
@ -319,7 +479,6 @@ define ptx_kernel i32 @grid_const_partial_escapemem(ptr byval(%struct.s) align 4
|
|||
; OPT-NEXT: [[ADD:%.*]] = add i32 [[VAL1]], [[VAL2]]
|
||||
; OPT-NEXT: [[CALL2:%.*]] = call i32 @escape(ptr [[PTR1]])
|
||||
; OPT-NEXT: ret i32 [[ADD]]
|
||||
;
|
||||
%ptr1 = getelementptr inbounds %struct.s, ptr %input, i32 0, i32 0
|
||||
%val1 = load i32, ptr %ptr1
|
||||
%ptr2 = getelementptr inbounds %struct.s, ptr %input, i32 0, i32 1
|
||||
|
|
@ -331,25 +490,45 @@ define ptx_kernel i32 @grid_const_partial_escapemem(ptr byval(%struct.s) align 4
|
|||
}
|
||||
|
||||
define ptx_kernel void @grid_const_phi(ptr byval(%struct.s) align 4 "nvvm.grid_constant" %input1, ptr %inout) {
|
||||
; PTX-LABEL: grid_const_phi(
|
||||
; PTX: {
|
||||
; PTX-NEXT: .reg .pred %p<2>;
|
||||
; PTX-NEXT: .reg .b32 %r<3>;
|
||||
; PTX-NEXT: .reg .b64 %rd<4>;
|
||||
; PTX-EMPTY:
|
||||
; PTX-NEXT: // %bb.0:
|
||||
; PTX-NEXT: mov.b64 %rd3, grid_const_phi_param_0;
|
||||
; PTX-NEXT: ld.param.b64 %rd2, [grid_const_phi_param_1];
|
||||
; PTX-NEXT: cvta.to.global.u64 %rd1, %rd2;
|
||||
; PTX-NEXT: ld.global.b32 %r1, [%rd1];
|
||||
; PTX-NEXT: setp.lt.s32 %p1, %r1, 0;
|
||||
; PTX-NEXT: @%p1 bra $L__BB9_2;
|
||||
; PTX-NEXT: // %bb.1: // %second
|
||||
; PTX-NEXT: add.s64 %rd3, %rd3, 4;
|
||||
; PTX-NEXT: $L__BB9_2: // %merge
|
||||
; PTX-NEXT: ld.param.b32 %r2, [%rd3];
|
||||
; PTX-NEXT: st.global.b32 [%rd1], %r2;
|
||||
; PTX-NEXT: ret;
|
||||
; PTX-DEFAULT-LABEL: grid_const_phi(
|
||||
; PTX-DEFAULT: {
|
||||
; PTX-DEFAULT-NEXT: .reg .pred %p<2>;
|
||||
; PTX-DEFAULT-NEXT: .reg .b32 %r<3>;
|
||||
; PTX-DEFAULT-NEXT: .reg .b64 %rd<4>;
|
||||
; PTX-DEFAULT-EMPTY:
|
||||
; PTX-DEFAULT-NEXT: // %bb.0:
|
||||
; PTX-DEFAULT-NEXT: mov.b64 %rd3, grid_const_phi_param_0;
|
||||
; PTX-DEFAULT-NEXT: ld.param.b64 %rd2, [grid_const_phi_param_1];
|
||||
; PTX-DEFAULT-NEXT: cvta.to.global.u64 %rd1, %rd2;
|
||||
; PTX-DEFAULT-NEXT: ld.global.b32 %r1, [%rd1];
|
||||
; PTX-DEFAULT-NEXT: setp.lt.s32 %p1, %r1, 0;
|
||||
; PTX-DEFAULT-NEXT: @%p1 bra $L__BB9_2;
|
||||
; PTX-DEFAULT-NEXT: // %bb.1: // %second
|
||||
; PTX-DEFAULT-NEXT: add.s64 %rd3, %rd3, 4;
|
||||
; PTX-DEFAULT-NEXT: $L__BB9_2: // %merge
|
||||
; PTX-DEFAULT-NEXT: ld.param.b32 %r2, [%rd3];
|
||||
; PTX-DEFAULT-NEXT: st.global.b32 [%rd1], %r2;
|
||||
; PTX-DEFAULT-NEXT: ret;
|
||||
;
|
||||
; PTX-SHORT-PTR-LABEL: grid_const_phi(
|
||||
; PTX-SHORT-PTR: {
|
||||
; PTX-SHORT-PTR-NEXT: .reg .pred %p<2>;
|
||||
; PTX-SHORT-PTR-NEXT: .reg .b32 %r<4>;
|
||||
; PTX-SHORT-PTR-NEXT: .reg .b64 %rd<3>;
|
||||
; PTX-SHORT-PTR-EMPTY:
|
||||
; PTX-SHORT-PTR-NEXT: // %bb.0:
|
||||
; PTX-SHORT-PTR-NEXT: mov.b32 %r3, grid_const_phi_param_0;
|
||||
; PTX-SHORT-PTR-NEXT: ld.param.b64 %rd2, [grid_const_phi_param_1];
|
||||
; PTX-SHORT-PTR-NEXT: cvta.to.global.u64 %rd1, %rd2;
|
||||
; PTX-SHORT-PTR-NEXT: ld.global.b32 %r1, [%rd1];
|
||||
; PTX-SHORT-PTR-NEXT: setp.lt.s32 %p1, %r1, 0;
|
||||
; PTX-SHORT-PTR-NEXT: @%p1 bra $L__BB9_2;
|
||||
; PTX-SHORT-PTR-NEXT: // %bb.1: // %second
|
||||
; PTX-SHORT-PTR-NEXT: add.s32 %r3, %r3, 4;
|
||||
; PTX-SHORT-PTR-NEXT: $L__BB9_2: // %merge
|
||||
; PTX-SHORT-PTR-NEXT: ld.param.b32 %r2, [%r3];
|
||||
; PTX-SHORT-PTR-NEXT: st.global.b32 [%rd1], %r2;
|
||||
; PTX-SHORT-PTR-NEXT: ret;
|
||||
; OPT-LABEL: define ptx_kernel void @grid_const_phi(
|
||||
; OPT-SAME: ptr addrspace(101) readonly byval([[STRUCT_S:%.*]]) align 4 "nvvm.grid_constant" [[INPUT1:%.*]], ptr [[INOUT:%.*]]) #[[ATTR0]] {
|
||||
; OPT-NEXT: [[INPUT1_PARAM_GEN:%.*]] = addrspacecast ptr addrspace(101) [[INPUT1]] to ptr
|
||||
|
|
@ -367,7 +546,6 @@ define ptx_kernel void @grid_const_phi(ptr byval(%struct.s) align 4 "nvvm.grid_c
|
|||
; OPT-NEXT: [[VALLOADED:%.*]] = load i32, ptr [[PTRNEW]], align 4
|
||||
; OPT-NEXT: store i32 [[VALLOADED]], ptr [[INOUT]], align 4
|
||||
; OPT-NEXT: ret void
|
||||
;
|
||||
|
||||
%val = load i32, ptr %inout
|
||||
%less = icmp slt i32 %val, 0
|
||||
|
|
@ -387,26 +565,47 @@ merge:
|
|||
|
||||
; NOTE: %input2 is *not* grid_constant
|
||||
define ptx_kernel void @grid_const_phi_ngc(ptr byval(%struct.s) align 4 "nvvm.grid_constant" %input1, ptr byval(%struct.s) align 4 %input2, ptr %inout) {
|
||||
; PTX-LABEL: grid_const_phi_ngc(
|
||||
; PTX: {
|
||||
; PTX-NEXT: .reg .pred %p<2>;
|
||||
; PTX-NEXT: .reg .b32 %r<3>;
|
||||
; PTX-NEXT: .reg .b64 %rd<5>;
|
||||
; PTX-EMPTY:
|
||||
; PTX-NEXT: // %bb.0:
|
||||
; PTX-NEXT: mov.b64 %rd4, grid_const_phi_ngc_param_0;
|
||||
; PTX-NEXT: ld.param.b64 %rd3, [grid_const_phi_ngc_param_2];
|
||||
; PTX-NEXT: cvta.to.global.u64 %rd1, %rd3;
|
||||
; PTX-NEXT: ld.global.b32 %r1, [%rd1];
|
||||
; PTX-NEXT: setp.lt.s32 %p1, %r1, 0;
|
||||
; PTX-NEXT: @%p1 bra $L__BB10_2;
|
||||
; PTX-NEXT: // %bb.1: // %second
|
||||
; PTX-NEXT: mov.b64 %rd2, grid_const_phi_ngc_param_1;
|
||||
; PTX-NEXT: add.s64 %rd4, %rd2, 4;
|
||||
; PTX-NEXT: $L__BB10_2: // %merge
|
||||
; PTX-NEXT: ld.param.b32 %r2, [%rd4];
|
||||
; PTX-NEXT: st.global.b32 [%rd1], %r2;
|
||||
; PTX-NEXT: ret;
|
||||
; PTX-DEFAULT-LABEL: grid_const_phi_ngc(
|
||||
; PTX-DEFAULT: {
|
||||
; PTX-DEFAULT-NEXT: .reg .pred %p<2>;
|
||||
; PTX-DEFAULT-NEXT: .reg .b32 %r<3>;
|
||||
; PTX-DEFAULT-NEXT: .reg .b64 %rd<5>;
|
||||
; PTX-DEFAULT-EMPTY:
|
||||
; PTX-DEFAULT-NEXT: // %bb.0:
|
||||
; PTX-DEFAULT-NEXT: mov.b64 %rd4, grid_const_phi_ngc_param_0;
|
||||
; PTX-DEFAULT-NEXT: ld.param.b64 %rd3, [grid_const_phi_ngc_param_2];
|
||||
; PTX-DEFAULT-NEXT: cvta.to.global.u64 %rd1, %rd3;
|
||||
; PTX-DEFAULT-NEXT: ld.global.b32 %r1, [%rd1];
|
||||
; PTX-DEFAULT-NEXT: setp.lt.s32 %p1, %r1, 0;
|
||||
; PTX-DEFAULT-NEXT: @%p1 bra $L__BB10_2;
|
||||
; PTX-DEFAULT-NEXT: // %bb.1: // %second
|
||||
; PTX-DEFAULT-NEXT: mov.b64 %rd2, grid_const_phi_ngc_param_1;
|
||||
; PTX-DEFAULT-NEXT: add.s64 %rd4, %rd2, 4;
|
||||
; PTX-DEFAULT-NEXT: $L__BB10_2: // %merge
|
||||
; PTX-DEFAULT-NEXT: ld.param.b32 %r2, [%rd4];
|
||||
; PTX-DEFAULT-NEXT: st.global.b32 [%rd1], %r2;
|
||||
; PTX-DEFAULT-NEXT: ret;
|
||||
;
|
||||
; PTX-SHORT-PTR-LABEL: grid_const_phi_ngc(
|
||||
; PTX-SHORT-PTR: {
|
||||
; PTX-SHORT-PTR-NEXT: .reg .pred %p<2>;
|
||||
; PTX-SHORT-PTR-NEXT: .reg .b32 %r<5>;
|
||||
; PTX-SHORT-PTR-NEXT: .reg .b64 %rd<3>;
|
||||
; PTX-SHORT-PTR-EMPTY:
|
||||
; PTX-SHORT-PTR-NEXT: // %bb.0:
|
||||
; PTX-SHORT-PTR-NEXT: mov.b32 %r4, grid_const_phi_ngc_param_0;
|
||||
; PTX-SHORT-PTR-NEXT: ld.param.b64 %rd2, [grid_const_phi_ngc_param_2];
|
||||
; PTX-SHORT-PTR-NEXT: cvta.to.global.u64 %rd1, %rd2;
|
||||
; PTX-SHORT-PTR-NEXT: ld.global.b32 %r2, [%rd1];
|
||||
; PTX-SHORT-PTR-NEXT: setp.lt.s32 %p1, %r2, 0;
|
||||
; PTX-SHORT-PTR-NEXT: @%p1 bra $L__BB10_2;
|
||||
; PTX-SHORT-PTR-NEXT: // %bb.1: // %second
|
||||
; PTX-SHORT-PTR-NEXT: mov.b32 %r1, grid_const_phi_ngc_param_1;
|
||||
; PTX-SHORT-PTR-NEXT: add.s32 %r4, %r1, 4;
|
||||
; PTX-SHORT-PTR-NEXT: $L__BB10_2: // %merge
|
||||
; PTX-SHORT-PTR-NEXT: ld.param.b32 %r3, [%r4];
|
||||
; PTX-SHORT-PTR-NEXT: st.global.b32 [%rd1], %r3;
|
||||
; PTX-SHORT-PTR-NEXT: ret;
|
||||
; OPT-LABEL: define ptx_kernel void @grid_const_phi_ngc(
|
||||
; OPT-SAME: ptr addrspace(101) readonly byval([[STRUCT_S:%.*]]) align 4 "nvvm.grid_constant" [[INPUT1:%.*]], ptr addrspace(101) readonly byval([[STRUCT_S]]) align 4 [[INPUT2:%.*]], ptr [[INOUT:%.*]]) #[[ATTR0]] {
|
||||
; OPT-NEXT: [[INPUT2_PARAM_GEN:%.*]] = addrspacecast ptr addrspace(101) [[INPUT2]] to ptr
|
||||
|
|
@ -425,7 +624,6 @@ define ptx_kernel void @grid_const_phi_ngc(ptr byval(%struct.s) align 4 "nvvm.gr
|
|||
; OPT-NEXT: [[VALLOADED:%.*]] = load i32, ptr [[PTRNEW]], align 4
|
||||
; OPT-NEXT: store i32 [[VALLOADED]], ptr [[INOUT]], align 4
|
||||
; OPT-NEXT: ret void
|
||||
;
|
||||
%val = load i32, ptr %inout
|
||||
%less = icmp slt i32 %val, 0
|
||||
br i1 %less, label %first, label %second
|
||||
|
|
@ -444,23 +642,41 @@ merge:
|
|||
|
||||
; NOTE: %input2 is *not* grid_constant
|
||||
define ptx_kernel void @grid_const_select(ptr byval(i32) align 4 "nvvm.grid_constant" %input1, ptr byval(i32) align 4 %input2, ptr %inout) {
|
||||
; PTX-LABEL: grid_const_select(
|
||||
; PTX: {
|
||||
; PTX-NEXT: .reg .pred %p<2>;
|
||||
; PTX-NEXT: .reg .b32 %r<3>;
|
||||
; PTX-NEXT: .reg .b64 %rd<6>;
|
||||
; PTX-EMPTY:
|
||||
; PTX-NEXT: // %bb.0:
|
||||
; PTX-NEXT: mov.b64 %rd1, grid_const_select_param_0;
|
||||
; PTX-NEXT: ld.param.b64 %rd2, [grid_const_select_param_2];
|
||||
; PTX-NEXT: cvta.to.global.u64 %rd3, %rd2;
|
||||
; PTX-NEXT: mov.b64 %rd4, grid_const_select_param_1;
|
||||
; PTX-NEXT: ld.global.b32 %r1, [%rd3];
|
||||
; PTX-NEXT: setp.lt.s32 %p1, %r1, 0;
|
||||
; PTX-NEXT: selp.b64 %rd5, %rd1, %rd4, %p1;
|
||||
; PTX-NEXT: ld.param.b32 %r2, [%rd5];
|
||||
; PTX-NEXT: st.global.b32 [%rd3], %r2;
|
||||
; PTX-NEXT: ret;
|
||||
; PTX-DEFAULT-LABEL: grid_const_select(
|
||||
; PTX-DEFAULT: {
|
||||
; PTX-DEFAULT-NEXT: .reg .pred %p<2>;
|
||||
; PTX-DEFAULT-NEXT: .reg .b32 %r<3>;
|
||||
; PTX-DEFAULT-NEXT: .reg .b64 %rd<6>;
|
||||
; PTX-DEFAULT-EMPTY:
|
||||
; PTX-DEFAULT-NEXT: // %bb.0:
|
||||
; PTX-DEFAULT-NEXT: mov.b64 %rd1, grid_const_select_param_0;
|
||||
; PTX-DEFAULT-NEXT: ld.param.b64 %rd2, [grid_const_select_param_2];
|
||||
; PTX-DEFAULT-NEXT: cvta.to.global.u64 %rd3, %rd2;
|
||||
; PTX-DEFAULT-NEXT: mov.b64 %rd4, grid_const_select_param_1;
|
||||
; PTX-DEFAULT-NEXT: ld.global.b32 %r1, [%rd3];
|
||||
; PTX-DEFAULT-NEXT: setp.lt.s32 %p1, %r1, 0;
|
||||
; PTX-DEFAULT-NEXT: selp.b64 %rd5, %rd1, %rd4, %p1;
|
||||
; PTX-DEFAULT-NEXT: ld.param.b32 %r2, [%rd5];
|
||||
; PTX-DEFAULT-NEXT: st.global.b32 [%rd3], %r2;
|
||||
; PTX-DEFAULT-NEXT: ret;
|
||||
;
|
||||
; PTX-SHORT-PTR-LABEL: grid_const_select(
|
||||
; PTX-SHORT-PTR: {
|
||||
; PTX-SHORT-PTR-NEXT: .reg .pred %p<2>;
|
||||
; PTX-SHORT-PTR-NEXT: .reg .b32 %r<6>;
|
||||
; PTX-SHORT-PTR-NEXT: .reg .b64 %rd<3>;
|
||||
; PTX-SHORT-PTR-EMPTY:
|
||||
; PTX-SHORT-PTR-NEXT: // %bb.0:
|
||||
; PTX-SHORT-PTR-NEXT: mov.b32 %r1, grid_const_select_param_0;
|
||||
; PTX-SHORT-PTR-NEXT: ld.param.b64 %rd1, [grid_const_select_param_2];
|
||||
; PTX-SHORT-PTR-NEXT: cvta.to.global.u64 %rd2, %rd1;
|
||||
; PTX-SHORT-PTR-NEXT: mov.b32 %r2, grid_const_select_param_1;
|
||||
; PTX-SHORT-PTR-NEXT: ld.global.b32 %r3, [%rd2];
|
||||
; PTX-SHORT-PTR-NEXT: setp.lt.s32 %p1, %r3, 0;
|
||||
; PTX-SHORT-PTR-NEXT: selp.b32 %r4, %r1, %r2, %p1;
|
||||
; PTX-SHORT-PTR-NEXT: ld.param.b32 %r5, [%r4];
|
||||
; PTX-SHORT-PTR-NEXT: st.global.b32 [%rd2], %r5;
|
||||
; PTX-SHORT-PTR-NEXT: ret;
|
||||
; OPT-LABEL: define ptx_kernel void @grid_const_select(
|
||||
; OPT-SAME: ptr addrspace(101) readonly byval(i32) align 4 "nvvm.grid_constant" [[INPUT1:%.*]], ptr addrspace(101) readonly byval(i32) align 4 [[INPUT2:%.*]], ptr [[INOUT:%.*]]) #[[ATTR0]] {
|
||||
; OPT-NEXT: [[INPUT2_PARAM_GEN:%.*]] = addrspacecast ptr addrspace(101) [[INPUT2]] to ptr
|
||||
|
|
@ -471,7 +687,6 @@ define ptx_kernel void @grid_const_select(ptr byval(i32) align 4 "nvvm.grid_cons
|
|||
; OPT-NEXT: [[VALLOADED:%.*]] = load i32, ptr [[PTRNEW]], align 4
|
||||
; OPT-NEXT: store i32 [[VALLOADED]], ptr [[INOUT]], align 4
|
||||
; OPT-NEXT: ret void
|
||||
;
|
||||
%val = load i32, ptr %inout
|
||||
%less = icmp slt i32 %val, 0
|
||||
%ptrnew = select i1 %less, ptr %input1, ptr %input2
|
||||
|
|
@ -481,19 +696,34 @@ define ptx_kernel void @grid_const_select(ptr byval(i32) align 4 "nvvm.grid_cons
|
|||
}
|
||||
|
||||
define ptx_kernel i32 @grid_const_ptrtoint(ptr byval(i32) align 4 "nvvm.grid_constant" %input) {
|
||||
; PTX-LABEL: grid_const_ptrtoint(
|
||||
; PTX: {
|
||||
; PTX-NEXT: .reg .b32 %r<4>;
|
||||
; PTX-NEXT: .reg .b64 %rd<3>;
|
||||
; PTX-EMPTY:
|
||||
; PTX-NEXT: // %bb.0:
|
||||
; PTX-NEXT: mov.b64 %rd1, grid_const_ptrtoint_param_0;
|
||||
; PTX-NEXT: cvta.param.u64 %rd2, %rd1;
|
||||
; PTX-NEXT: ld.param.b32 %r1, [grid_const_ptrtoint_param_0];
|
||||
; PTX-NEXT: cvt.u32.u64 %r2, %rd2;
|
||||
; PTX-NEXT: add.s32 %r3, %r1, %r2;
|
||||
; PTX-NEXT: st.param.b32 [func_retval0], %r3;
|
||||
; PTX-NEXT: ret;
|
||||
; PTX-DEFAULT-LABEL: grid_const_ptrtoint(
|
||||
; PTX-DEFAULT: {
|
||||
; PTX-DEFAULT-NEXT: .reg .b32 %r<4>;
|
||||
; PTX-DEFAULT-NEXT: .reg .b64 %rd<3>;
|
||||
; PTX-DEFAULT-EMPTY:
|
||||
; PTX-DEFAULT-NEXT: // %bb.0:
|
||||
; PTX-DEFAULT-NEXT: mov.b64 %rd1, grid_const_ptrtoint_param_0;
|
||||
; PTX-DEFAULT-NEXT: cvta.param.u64 %rd2, %rd1;
|
||||
; PTX-DEFAULT-NEXT: ld.param.b32 %r1, [grid_const_ptrtoint_param_0];
|
||||
; PTX-DEFAULT-NEXT: cvt.u32.u64 %r2, %rd2;
|
||||
; PTX-DEFAULT-NEXT: add.s32 %r3, %r1, %r2;
|
||||
; PTX-DEFAULT-NEXT: st.param.b32 [func_retval0], %r3;
|
||||
; PTX-DEFAULT-NEXT: ret;
|
||||
;
|
||||
; PTX-SHORT-PTR-LABEL: grid_const_ptrtoint(
|
||||
; PTX-SHORT-PTR: {
|
||||
; PTX-SHORT-PTR-NEXT: .reg .b32 %r<5>;
|
||||
; PTX-SHORT-PTR-NEXT: .reg .b64 %rd<3>;
|
||||
; PTX-SHORT-PTR-EMPTY:
|
||||
; PTX-SHORT-PTR-NEXT: // %bb.0:
|
||||
; PTX-SHORT-PTR-NEXT: mov.b32 %r1, grid_const_ptrtoint_param_0;
|
||||
; PTX-SHORT-PTR-NEXT: cvt.u64.u32 %rd1, %r1;
|
||||
; PTX-SHORT-PTR-NEXT: cvta.param.u64 %rd2, %rd1;
|
||||
; PTX-SHORT-PTR-NEXT: ld.param.b32 %r2, [grid_const_ptrtoint_param_0];
|
||||
; PTX-SHORT-PTR-NEXT: cvt.u32.u64 %r3, %rd2;
|
||||
; PTX-SHORT-PTR-NEXT: add.s32 %r4, %r2, %r3;
|
||||
; PTX-SHORT-PTR-NEXT: st.param.b32 [func_retval0], %r4;
|
||||
; PTX-SHORT-PTR-NEXT: ret;
|
||||
; OPT-LABEL: define ptx_kernel i32 @grid_const_ptrtoint(
|
||||
; OPT-SAME: ptr addrspace(101) readonly byval(i32) align 4 "nvvm.grid_constant" [[INPUT:%.*]]) #[[ATTR0]] {
|
||||
; OPT-NEXT: [[INPUT1:%.*]] = addrspacecast ptr addrspace(101) [[INPUT]] to ptr
|
||||
|
|
@ -501,7 +731,6 @@ define ptx_kernel i32 @grid_const_ptrtoint(ptr byval(i32) align 4 "nvvm.grid_con
|
|||
; OPT-NEXT: [[PTRVAL:%.*]] = ptrtoint ptr [[INPUT1]] to i32
|
||||
; OPT-NEXT: [[KEEPALIVE:%.*]] = add i32 [[INPUT3]], [[PTRVAL]]
|
||||
; OPT-NEXT: ret i32 [[KEEPALIVE]]
|
||||
;
|
||||
%val = load i32, ptr %input
|
||||
%ptrval = ptrtoint ptr %input to i32
|
||||
%keepalive = add i32 %val, %ptrval
|
||||
|
|
|
|||
|
|
@ -7,6 +7,8 @@
|
|||
//===----------------------------------------------------------------------===//
|
||||
|
||||
#include "llvm/TargetParser/Triple.h"
|
||||
#include "llvm/ADT/SmallVector.h"
|
||||
#include "llvm/ADT/StringExtras.h"
|
||||
#include "llvm/Support/CodeGen.h"
|
||||
#include "llvm/Support/FormatVariadic.h"
|
||||
#include "llvm/Support/VersionTuple.h"
|
||||
|
|
@ -3821,6 +3823,36 @@ TEST(TripleTest, DefaultWCharSize) {
|
|||
EXPECT_EQ(1u, Triple("xcore-unknown-unknown").getDefaultWCharSize());
|
||||
}
|
||||
|
||||
TEST(DataLayoutTest, NVPTX) {
|
||||
Triple TT32 = Triple("nvptx-nvidia-cuda");
|
||||
Triple TT64 = Triple("nvptx64-nvidia-cuda");
|
||||
|
||||
auto PointerLayoutSpecs = [](StringRef DataLayout) {
|
||||
SmallVector<std::string, 8> Specs;
|
||||
for (StringRef Spec : split(DataLayout, '-'))
|
||||
if (Spec.starts_with("p"))
|
||||
Specs.emplace_back(Spec);
|
||||
return Specs;
|
||||
};
|
||||
|
||||
// The 32-bit target uses a single 32-bit pointer specification and is
|
||||
// unaffected by the ABI name.
|
||||
EXPECT_THAT(PointerLayoutSpecs(TT32.computeDataLayout("")),
|
||||
testing::ElementsAre("p:32:32"));
|
||||
EXPECT_THAT(PointerLayoutSpecs(TT32.computeDataLayout("shortptr")),
|
||||
testing::ElementsAre("p:32:32"));
|
||||
|
||||
// The default 64-bit target only shrinks Tensor Memory (addrspace:6).
|
||||
EXPECT_THAT(PointerLayoutSpecs(TT64.computeDataLayout("")),
|
||||
testing::ElementsAre("p6:32:32"));
|
||||
|
||||
// In shortptr mode the extra address spaces become 32-bit. The pointer
|
||||
// specifications must remain sorted by address space.
|
||||
EXPECT_THAT(PointerLayoutSpecs(TT64.computeDataLayout("shortptr")),
|
||||
testing::ElementsAre("p3:32:32", "p4:32:32", "p5:32:32",
|
||||
"p6:32:32", "p7:32:32", "p101:32:32"));
|
||||
}
|
||||
|
||||
TEST(DataLayoutTest, CheriRISCV32) {
|
||||
Triple TT = Triple("riscv32-unknown-unknown");
|
||||
|
||||
|
|
|
|||
Loading…
Reference in a new issue