Changeset View
Standalone View
test/CodeGen/NVPTX/insert-shared-depot.ll
- This file was added.
; Check that the shared depot is emitted when the "has-nvptx-shared-depot" is used. | |||||
; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefixes=PTX32,CHECK | |||||
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefixes=PTX64,CHECK | |||||
; CHECK: {{.*}}kernel() | |||||
tra: You could put common checks under the same label (e.g. `CHECK`) and run tests with `-check… | |||||
; PTX32: .local .align 8{{.*}}.b8{{.*}}__local_depot0 | |||||
; PTX64: .local .align 8{{.*}}.b8{{.*}}__local_depot0 | |||||
; PTX32: .shared .align 8{{.*}}.b8{{.*}}__shared_depot0 | |||||
; PTX64: .shared .align 8{{.*}}.b8{{.*}}__shared_depot0 | |||||
; PTX32: .reg .b32{{.*}}%SPS; | |||||
; PTX64: .reg .b64{{.*}}%SPS; | |||||
; PTX32: .reg .b32{{.*}}%SPSH; | |||||
; PTX64: .reg .b64{{.*}}%SPSH; | |||||
; PTX32: mov.u32{{.*}}%SPSH, __shared_depot0; | |||||
; PTX64: mov.u64{{.*}}%SPSH, __shared_depot0; | |||||
; PTX32: cvta.shared.u32{{.*}}%SPS, %SPSH; | |||||
; PTX64: 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() #0 { | |||||
%A = alloca i32, align 4 | |||||
'LABEL' is not a check-prefix and @linsert_shared_depot is not this function's name, so I'm puzzled what this line is supposed to do. Did you intend <prefix>-LABEL: @kernel ? This appears in all the test cases in the patch. tra: 'LABEL' is not a check-prefix and `@linsert_shared_depot` is not this function's name, so I'm… | |||||
Not Done ReplyInline ActionsThis is modeled after the lower-alloca.ll test which has a similar label. The label is always equal to the name of the test file. In this particular case there is a typo, it should be "insert_shared_depot" not "linsert_shared_depot" gtbercea: This is modeled after the lower-alloca.ll test which has a similar label. The label is always… | |||||
Not Done ReplyInline Actions
lower-alloca.ll indeed has the same problem.
I don't think FileCheck has such a feature. Nor do I see anything matching this description in the FileCheck documentation. Nor does it work. See below.
The line does not check *anything* right now. In this test FileCheck only pays attention to lines that have CHECK or PTX64/PTX32. This line contains neither and is ignored. You can do an experiment -- replace the line with ; LABEL: this should never match and run the test. I've tried that on lower-alloca.ll and the test, as expected, passes regardless of the nonsense I put after the LABEL:. tra: > This is modeled after the lower-alloca.ll test which has a similar label.
lower-alloca.ll… | |||||
%shared_args = alloca i8**, align 8 | |||||
call void @callee(i8*** %shared_args) | |||||
store i32 10, i32* %A | |||||
ret void | |||||
} | |||||
declare void @callee(i8***) | |||||
attributes #0 = {"has-nvptx-shared-depot"} | |||||
!nvvm.annotations = !{!0} | |||||
!0 = !{void ()* @kernel, !"kernel", i32 1} |
You could put common checks under the same label (e.g. CHECK) and run tests with -check-prefixes=PTX32,CHECK