Changeset View
Changeset View
Standalone View
Standalone View
test/CodeGen/NVPTX/nvptx-function-data-sharing.ll
- This file was added.
| ; Test case for NVPTXFunctionDataSharing pass. | |||||
| ; RUN: opt < %s -S -nvptx-function-data-sharing -infer-address-spaces | FileCheck %s | |||||
| ; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 | FileCheck %s --check-prefix PTX | |||||
| target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64" | |||||
| target triple = "nvptx64-unknown-unknown" | |||||
| define void @kernel() #0 { | |||||
| ; PTX-LABEL: .visible .entry kernel( | |||||
| %A = alloca i32 | |||||
| ; CHECK: addrspacecast i32* %A to i32 addrspace(3)* | |||||
| ; CHECK: addrspacecast i32 addrspace(3)* %A1 to i32* | |||||
| ; CHECK: store i32 0, i32 addrspace(3)* {{%.+}} | |||||
| ; PTX: add.u64 {{%rd[0-9]+}}, %SPS, 0; | |||||
| ; PTX: cvta.to.shared.u64 {{%rd[0-9]+}}, {{%rd[0-9]+}}; | |||||
| ; PTX: st.shared.u32 [{{%rd[0-9]+}}], {{%r[0-9]+}} | |||||
| %shared_args = alloca i32** | |||||
| call void @callee(i32*** %shared_args) | |||||
| %1 = load i32**, i32*** %shared_args | |||||
| %2 = getelementptr inbounds i32*, i32** %1, i64 0 | |||||
| store i32* %A, i32** %2 | |||||
| store i32 0, i32* %A | |||||
| ret void | |||||
| } | |||||
| declare void @callee(i32***) | |||||
| attributes #0 = {"has-nvptx-shared-depot"} | |||||
| !nvvm.annotations = !{!0} | |||||
| !0 = !{void ()* @kernel, !"kernel", i32 1} | |||||