Changeset View
Changeset View
Standalone View
Standalone View
test/CodeGen/NVPTX/no-shared-depot.ll
- This file was added.
| ; Test case for when the shared depot does not need to be emitted. | |||||
| ; The shared depot is not emitted if the function does not have the | |||||
| ; "has-nvptx-shared-depot" attribute. | |||||
| ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix=PTX32 | |||||
| ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefix=PTX64 | |||||
| ; PTX32: {{.*}}kernel() | |||||
| ; PTX64: {{.*}}kernel() | |||||
| ; PTX32: .local .align 8{{.*}}.b8{{.*}}__local_depot0 | |||||
| ; PTX64: .local .align 8{{.*}}.b8{{.*}}__local_depot0 | |||||
| ; PTX32-NOT: .shared .align 8{{.*}}.b8{{.*}}__shared_depot0 | |||||
| ; PTX64-NOT: .shared .align 8{{.*}}.b8{{.*}}__shared_depot0 | |||||
| ; PTX32-NOT: .reg .b32{{.*}}%SPS; | |||||
| ; PTX64-NOT: .reg .b64{{.*}}%SPS; | |||||
| ; PTX32-NOT: .reg .b32{{.*}}%SPSH; | |||||
| ; PTX64-NOT: .reg .b64{{.*}}%SPSH; | |||||
| ; PTX32-NOT: mov.u32{{.*}}%SPSH, __shared_depot0; | |||||
| ; PTX64-NOT: mov.u64{{.*}}%SPSH, __shared_depot0; | |||||
| ; PTX32-NOT: cvta.shared.u32{{.*}}%SPS, %SPSH; | |||||
| ; PTX64-NOT: cvta.shared.u64{{.*}}%SPS, %SPSH; | |||||
| target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64" | |||||
| target triple = "nvptx64-unknown-unknown" | |||||
| define void @kernel() { | |||||
| %A = alloca i32, align 4 | |||||
| %shared_args = alloca i8**, align 8 | |||||
| call void @callee(i8*** %shared_args) | |||||
| store i32 10, i32* %A | |||||
| ret void | |||||
| } | |||||
| declare void @callee(i8***) | |||||
| !nvvm.annotations = !{!0} | |||||
| !0 = !{void ()* @kernel, !"kernel", i32 1} | |||||