This is an archive of the discontinued LLVM Phabricator instance.

Differential D152199

[NVPTX] Adapt tests to make them usable with CUDA-12.x
ClosedPublic

Authored by tra on Jun 5 2023, 2:34 PM.

Details

Reviewers
jlebar
Commits
rGef8655adc8e0: [NVPTX] Adapt tests to make them usable with CUDA-12.x
Summary

CUDA-12 no longer supports 32-bit compilation, and we ned to adjust NVPTX tests to deal with that.
LLVM compilation itself is not affected, but optional ptxas tests no longer work in 32-bit mode.

Tests that were run with -march=nvptx only were switched to use nvptx64 and had the checks adjusted as necessary.
Tests that were targeting both nvptx and nvptx64 now have 32-bit ptxas compilation disabled with cuda-12+.

Diff Detail

Event Timeline

tra created this revision.Jun 5 2023, 2:34 PM
Herald added a project: Restricted Project. · View Herald TranscriptJun 5 2023, 2:34 PM
Herald added subscribers: mattd, gchakrabarti, asavonic and 2 others. · View Herald Transcript
tra edited the summary of this revision. (Show Details)Jun 5 2023, 2:52 PM
tra updated this revision to Diff 528587.Jun 5 2023, 2:53 PM
tra edited the summary of this revision. (Show Details)

Test fix.

tra published this revision for review.Jun 5 2023, 2:53 PM
tra added a reviewer: jlebar.
Herald added a project: Restricted Project. · View Herald TranscriptJun 5 2023, 2:53 PM
Herald added subscribers: llvm-commits, • pcwang-thead, jholewinski. · View Herald Transcript
Harbormaster completed remote builds in B236735: Diff 528587.Jun 5 2023, 3:45 PM
jlebar accepted this revision.Jun 5 2023, 4:18 PM
This revision is now accepted and ready to land.Jun 5 2023, 4:18 PM
tra updated this revision to Diff 528966.Jun 6 2023, 11:46 AM

Changed EOL style to unix on few stray tests still using DOS EOL.

Harbormaster completed remote builds in B237022: Diff 528966.Jun 6 2023, 1:40 PM
Closed by commit rGef8655adc8e0: [NVPTX] Adapt tests to make them usable with CUDA-12.x (authored by tra). · Explain WhyJun 6 2023, 2:22 PM
This revision was automatically updated to reflect the committed changes.
tra added a commit: rGef8655adc8e0: [NVPTX] Adapt tests to make them usable with CUDA-12.x.

Revision Contents

PathSize
llvm/
test/
CodeGen/
NVPTX/
2 lines
6 lines
24 lines
2 lines
4 lines
2 lines
4 lines
2 lines
2 lines
2 lines
2 lines
2 lines
4 lines
16 lines
4 lines
54 lines
4 lines
4 lines
4 lines
4 lines
4 lines
4 lines
4 lines
4 lines
2 lines
4 lines
4 lines
2 lines
4 lines
2 lines
2 lines
4 lines
4 lines
4 lines
4 lines
4 lines
8 lines
4 lines
4 lines
8 lines
4 lines
86 lines
8 lines
4 lines
4 lines
4 lines
4 lines
2 lines
2 lines
4 lines
4 lines
4 lines
16 lines
4 lines
4 lines
6 lines
4 lines
48 lines
2 lines
4 lines
4 lines
4 lines
2 lines
2 lines
4 lines
2 lines
2 lines
4 lines
6 lines
4 lines
8 lines
4 lines
4 lines
2 lines
4 lines
4 lines
8 lines
2 lines
4 lines
4 lines
8 lines
2 lines
2 lines
8 lines
4 lines
2 lines
8 lines
28 lines
52 lines
4 lines
2 lines
4 lines
4 lines
4 lines
4 lines
2 lines
8 lines
4 lines
4 lines
4 lines
4 lines
4 lines
4 lines
2 lines
52 lines
4 lines
2 lines
2 lines
4 lines
16 lines
6 lines
8 lines
4 lines
2 lines
20 lines
6 lines
8 lines
2 lines
4 lines
6 lines
4 lines
62 lines
2 lines
4 lines
15 lines
4 lines
4 lines
4 lines
1 line

Diff 529044

llvm/test/CodeGen/NVPTX/access-non-generic.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix PTX ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix PTX
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefix PTX ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefix PTX
; RUN: opt -mtriple=nvptx-- < %s -S -passes=infer-address-spaces | FileCheck %s --check-prefix IR ; RUN: opt -mtriple=nvptx-- < %s -S -passes=infer-address-spaces | FileCheck %s --check-prefix IR
; RUN: opt -mtriple=nvptx64-- < %s -S -passes=infer-address-spaces | FileCheck %s --check-prefix IR ; RUN: opt -mtriple=nvptx64-- < %s -S -passes=infer-address-spaces | FileCheck %s --check-prefix IR
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
@array = internal addrspace(3) global [10 x float] zeroinitializer, align 4 @array = internal addrspace(3) global [10 x float] zeroinitializer, align 4
@scalar = internal addrspace(3) global float 0.000000e+00, align 4 @scalar = internal addrspace(3) global float 0.000000e+00, align 4
; Verifies nvptx-favor-non-generic correctly optimizes generic address space ; Verifies nvptx-favor-non-generic correctly optimizes generic address space
; usage to non-generic address space usage for the patterns we claim to handle: ; usage to non-generic address space usage for the patterns we claim to handle:
; 1. load cast ; 1. load cast
▲ Show 20 Lines • Show All 167 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/add-sub-128bit.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefixes=COMMON,NOCARRY ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefixes=COMMON,NOCARRY
; RUN: llc < %s -march=nvptx -mcpu=sm_20 -mattr=+ptx43 | FileCheck %s --check-prefixes=COMMON,CARRY ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -mattr=+ptx43 | FileCheck %s --check-prefixes=COMMON,CARRY
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
target datalayout = "e-p:32:32:32-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 datalayout = "e-p:32:32:32-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"
; COMMON-LABEL: test_add ; COMMON-LABEL: test_add
define i128 @test_add(i128 %a, i128 %b) { define i128 @test_add(i128 %a, i128 %b) {
; NOCARRY: add.s64 ; NOCARRY: add.s64
; NOCARRY-NEXT: add.s64 ; NOCARRY-NEXT: add.s64
; NOCARRY-NEXT: setp.lt.u64 ; NOCARRY-NEXT: setp.lt.u64
Show All 24 Lines

llvm/test/CodeGen/NVPTX/addrspacecast-gvar.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; CHECK: .visible .global .align 4 .u32 g = 42; ; CHECK: .visible .global .align 4 .u32 g = 42;
; CHECK: .visible .global .align 1 .b8 ga[4] = {0, 1, 2, 3}; ; CHECK: .visible .global .align 1 .b8 ga[4] = {0, 1, 2, 3};
; CHECK: .visible .global .align 4 .u32 g2 = generic(g); ; CHECK: .visible .global .align 8 .u64 g2 = generic(g);
; CHECK: .visible .global .align 4 .u32 g3 = g; ; CHECK: .visible .global .align 8 .u64 g3 = g;
; CHECK: .visible .global .align 8 .u32 g4[2] = {0, generic(g)}; ; CHECK: .visible .global .align 8 .u64 g4[2] = {0, generic(g)};
; CHECK: .visible .global .align 8 .u32 g5[2] = {0, generic(g)+8}; ; CHECK: .visible .global .align 8 .u64 g5[2] = {0, generic(g)+8};
@g = addrspace(1) global i32 42 @g = addrspace(1) global i32 42
@ga = addrspace(1) global [4 x i8] c"\00\01\02\03" @ga = addrspace(1) global [4 x i8] c"\00\01\02\03"
@g2 = addrspace(1) global ptr addrspacecast (ptr addrspace(1) @g to ptr) @g2 = addrspace(1) global ptr addrspacecast (ptr addrspace(1) @g to ptr)
@g3 = addrspace(1) global ptr addrspace(1) @g @g3 = addrspace(1) global ptr addrspace(1) @g
@g4 = constant {ptr, ptr} {ptr null, ptr addrspacecast (ptr addrspace(1) @g to ptr)} @g4 = constant {ptr, ptr} {ptr null, ptr addrspacecast (ptr addrspace(1) @g to ptr)}
@g5 = constant {ptr, ptr} {ptr null, ptr addrspacecast (ptr addrspace(1) getelementptr (i32, ptr addrspace(1) @g, i32 2) to ptr)} @g5 = constant {ptr, ptr} {ptr null, ptr addrspacecast (ptr addrspace(1) getelementptr (i32, ptr addrspace(1) @g, i32 2) to ptr)}
; CHECK: .visible .global .align 4 .u32 g6 = generic(ga)+2; ; CHECK: .visible .global .align 8 .u64 g6 = generic(ga)+2;
@g6 = addrspace(1) global ptr getelementptr inbounds ( @g6 = addrspace(1) global ptr getelementptr inbounds (
[4 x i8], ptr addrspacecast (ptr addrspace(1) @ga to ptr), [4 x i8], ptr addrspacecast (ptr addrspace(1) @ga to ptr),
i32 0, i32 2 i32 0, i32 2
) )
; CHECK: .visible .global .align 4 .u32 g7 = generic(g); ; CHECK: .visible .global .align 8 .u64 g7 = generic(g);
@g7 = addrspace(1) global ptr addrspacecast ( @g7 = addrspace(1) global ptr addrspacecast (
ptr addrspace(1) @g ptr addrspace(1) @g
to ptr to ptr
) )
; CHECK: .visible .global .align 4 .u32 g8[2] = {0, g}; ; CHECK: .visible .global .align 8 .u64 g8[2] = {0, g};
@g8 = addrspace(1) global [2 x ptr addrspace(1)] [ptr addrspace(1) null, ptr addrspace(1) @g] @g8 = addrspace(1) global [2 x ptr addrspace(1)] [ptr addrspace(1) null, ptr addrspace(1) @g]
; CHECK: .visible .global .align 4 .u32 g9[2] = {0, generic(g)}; ; CHECK: .visible .global .align 8 .u64 g9[2] = {0, generic(g)};
@g9 = addrspace(1) global [2 x ptr] [ @g9 = addrspace(1) global [2 x ptr] [
ptr null, ptr null,
ptr addrspacecast (ptr addrspace(1) @g to ptr) ptr addrspacecast (ptr addrspace(1) @g to ptr)
] ]
; CHECK: .visible .global .align 4 .u32 g10[2] = {0, g}; ; CHECK: .visible .global .align 8 .u64 g10[2] = {0, g};
@g10 = addrspace(1) global [2 x ptr addrspace(1)] [ @g10 = addrspace(1) global [2 x ptr addrspace(1)] [
ptr addrspace(1) null, ptr addrspace(1) null,
ptr addrspace(1) @g ptr addrspace(1) @g
] ]
; CHECK: .visible .global .align 4 .u32 g11[2] = {0, generic(g)}; ; CHECK: .visible .global .align 8 .u64 g11[2] = {0, generic(g)};
@g11 = addrspace(1) global [2 x ptr] [ @g11 = addrspace(1) global [2 x ptr] [
ptr null, ptr null,
ptr addrspacecast (ptr addrspace(1) @g to ptr) ptr addrspacecast (ptr addrspace(1) @g to ptr)
] ]

llvm/test/CodeGen/NVPTX/addrspacecast.ll

; RUN: llc -O0 < %s -march=nvptx -mcpu=sm_20 | FileCheck %s -check-prefixes=ALL,CLS32,G32 ; RUN: llc -O0 < %s -march=nvptx -mcpu=sm_20 | FileCheck %s -check-prefixes=ALL,CLS32,G32
; RUN: llc -O0 < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s -check-prefixes=ALL,NOPTRCONV,CLS64,G64 ; RUN: llc -O0 < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s -check-prefixes=ALL,NOPTRCONV,CLS64,G64
; RUN: llc -O0 < %s -march=nvptx64 -mcpu=sm_20 --nvptx-short-ptr| FileCheck %s -check-prefixes=ALL,PTRCONV,CLS64,G64 ; RUN: llc -O0 < %s -march=nvptx64 -mcpu=sm_20 --nvptx-short-ptr| FileCheck %s -check-prefixes=ALL,PTRCONV,CLS64,G64
; RUN: %if ptxas %{ llc -O0 < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc -O0 < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{ llc -O0 < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc -O0 < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{ llc -O0 < %s -march=nvptx64 -mcpu=sm_20 --nvptx-short-ptr | %ptxas-verify %} ; RUN: %if ptxas %{ llc -O0 < %s -march=nvptx64 -mcpu=sm_20 --nvptx-short-ptr | %ptxas-verify %}
; ALL-LABEL: conv1 ; ALL-LABEL: conv1
define i32 @conv1(i32 addrspace(1)* %ptr) { define i32 @conv1(i32 addrspace(1)* %ptr) {
; G32: cvta.global.u32 ; G32: cvta.global.u32
; ALL-NOT: cvt.u64.u32 ; ALL-NOT: cvt.u64.u32
; G64: cvta.global.u64 ; G64: cvta.global.u64
▲ Show 20 Lines • Show All 88 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/aggr-param.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; Make sure aggregate param types get emitted properly. ; Make sure aggregate param types get emitted properly.
%struct.float4 = type { float, float, float, float } %struct.float4 = type { float, float, float, float }
; CHECK: .visible .func bar ; CHECK: .visible .func bar
; CHECK: .param .align 4 .b8 bar_param_0[16] ; CHECK: .param .align 4 .b8 bar_param_0[16]
define void @bar(%struct.float4 %f) { define void @bar(%struct.float4 %f) {
Show All 11 Lines

llvm/test/CodeGen/NVPTX/annotations.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
@texture = internal addrspace(1) global i64 0, align 8 @texture = internal addrspace(1) global i64 0, align 8
; CHECK: .global .texref texture ; CHECK: .global .texref texture
@surface = internal addrspace(1) global i64 0, align 8 @surface = internal addrspace(1) global i64 0, align 8
; CHECK: .global .surfref surface ; CHECK: .global .surfref surface
; CHECK: .entry kernel_func_maxntid ; CHECK: .entry kernel_func_maxntid
▲ Show 20 Lines • Show All 43 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/arg-lowering.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; CHECK: .visible .func (.param .align 16 .b8 func_retval0[16]) foo0( ; CHECK: .visible .func (.param .align 16 .b8 func_retval0[16]) foo0(
; CHECK: .param .align 4 .b8 foo0_param_0[8] ; CHECK: .param .align 4 .b8 foo0_param_0[8]
define <4 x float> @foo0({float, float} %arg0) { define <4 x float> @foo0({float, float} %arg0) {
ret <4 x float> <float 1.0, float 1.0, float 1.0, float 1.0> ret <4 x float> <float 1.0, float 1.0, float 1.0, float 1.0>
} }
; CHECK: .visible .func (.param .align 8 .b8 func_retval0[8]) foo1( ; CHECK: .visible .func (.param .align 8 .b8 func_retval0[8]) foo1(
; CHECK: .param .align 8 .b8 foo1_param_0[16] ; CHECK: .param .align 8 .b8 foo1_param_0[16]
define <2 x float> @foo1({float, float, i64} %arg0) { define <2 x float> @foo1({float, float, i64} %arg0) {
ret <2 x float> <float 1.0, float 1.0> ret <2 x float> <float 1.0, float 1.0>
} }

llvm/test/CodeGen/NVPTX/arithmetic-fp-sm20.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 -fp-contract=fast | FileCheck %s ; RUN: llc < %s -march=nvptx -mcpu=sm_20 -fp-contract=fast | FileCheck %s
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -fp-contract=fast | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -fp-contract=fast | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -fp-contract=fast | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_20 -fp-contract=fast | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -fp-contract=fast | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -fp-contract=fast | %ptxas-verify %}
;; These tests should run for all targets ;; These tests should run for all targets
;;===-- Basic instruction selection tests ---------------------------------===;; ;;===-- Basic instruction selection tests ---------------------------------===;;
;;; f64 ;;; f64
▲ Show 20 Lines • Show All 63 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/arithmetic-int.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
;; These tests should run for all targets ;; These tests should run for all targets
;;===-- Basic instruction selection tests ---------------------------------===;; ;;===-- Basic instruction selection tests ---------------------------------===;;
;;; i64 ;;; i64
▲ Show 20 Lines • Show All 308 Lines • Show Last 20 Lines

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

; RUN: llc < %s -march=nvptx -mcpu=sm_80 -mattr=+ptx70 | FileCheck -check-prefixes=CHECK,CHECK_PTX32 %s ; RUN: llc < %s -march=nvptx -mcpu=sm_80 -mattr=+ptx70 | FileCheck -check-prefixes=CHECK,CHECK_PTX32 %s
; RUN: llc < %s -march=nvptx64 -mcpu=sm_80 -mattr=+ptx70 | FileCheck -check-prefixes=CHECK,CHECK_PTX64 %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_80 -mattr=+ptx70 | FileCheck -check-prefixes=CHECK,CHECK_PTX64 %s
; RUN: %if ptxas-11.0 %{ llc < %s -march=nvptx -mcpu=sm_80 -mattr=+ptx70 | %ptxas-verify -arch=sm_80 %} ; RUN: %if ptxas-11.0 && ! ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_80 -mattr=+ptx70 | %ptxas-verify -arch=sm_80 %}
; RUN: %if ptxas-11.0 %{ llc < %s -march=nvptx64 -mcpu=sm_80 -mattr=+ptx70 | %ptxas-verify -arch=sm_80 %} ; RUN: %if ptxas-11.0 %{ llc < %s -march=nvptx64 -mcpu=sm_80 -mattr=+ptx70 | %ptxas-verify -arch=sm_80 %}
declare void @llvm.nvvm.cp.async.wait.group(i32) declare void @llvm.nvvm.cp.async.wait.group(i32)
; CHECK-LABEL: asyncwaitgroup ; CHECK-LABEL: asyncwaitgroup
define void @asyncwaitgroup() { define void @asyncwaitgroup() {
; CHECK: cp.async.wait_group 8; ; CHECK: cp.async.wait_group 8;
tail call void @llvm.nvvm.cp.async.wait.group(i32 8) tail call void @llvm.nvvm.cp.async.wait.group(i32 8)
▲ Show 20 Lines • Show All 116 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/atomics-sm60.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_60 | FileCheck %s ; RUN: llc < %s -march=nvptx -mcpu=sm_60 | FileCheck %s
; RUN: llc < %s -march=nvptx64 -mcpu=sm_60 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_60 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_60 | %ptxas-verify -arch=sm_60 %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_60 | %ptxas-verify -arch=sm_60 %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_60 | %ptxas-verify -arch=sm_60 %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_60 | %ptxas-verify -arch=sm_60 %}
; CHECK-LABEL: .func test( ; CHECK-LABEL: .func test(
define void @test(ptr %dp0, ptr addrspace(1) %dp1, ptr addrspace(3) %dp3, double %d) { define void @test(ptr %dp0, ptr addrspace(1) %dp1, ptr addrspace(3) %dp3, double %d) {
; CHECK: atom.add.f64 ; CHECK: atom.add.f64
%r1 = call double @llvm.nvvm.atomic.load.add.f64.p0(ptr %dp0, double %d) %r1 = call double @llvm.nvvm.atomic.load.add.f64.p0(ptr %dp0, double %d)
; CHECK: atom.global.add.f64 ; CHECK: atom.global.add.f64
%r2 = call double @llvm.nvvm.atomic.load.add.f64.p1(ptr addrspace(1) %dp1, double %d) %r2 = call double @llvm.nvvm.atomic.load.add.f64.p1(ptr addrspace(1) %dp1, double %d)
Show All 21 Lines

llvm/test/CodeGen/NVPTX/atomics-with-scope.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_60 | FileCheck %s -check-prefixes=CHECK,CHECK32 ; RUN: llc < %s -march=nvptx -mcpu=sm_60 | FileCheck %s -check-prefixes=CHECK,CHECK32
; RUN: llc < %s -march=nvptx64 -mcpu=sm_60 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_60 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_60 | %ptxas-verify -arch=sm_60 %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_60 | %ptxas-verify -arch=sm_60 %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_60 | %ptxas-verify -arch=sm_60 %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_60 | %ptxas-verify -arch=sm_60 %}
; CHECK-LABEL: .func test_atomics_scope( ; CHECK-LABEL: .func test_atomics_scope(
define void @test_atomics_scope(ptr %fp, float %f, define void @test_atomics_scope(ptr %fp, float %f,
ptr %dfp, double %df, ptr %dfp, double %df,
ptr %ip, i32 %i, ptr %ip, i32 %i,
ptr %uip, i32 %ui, ptr %uip, i32 %ui,
ptr %llp, i64 %ll) #0 { ptr %llp, i64 %ll) #0 {
▲ Show 20 Lines • Show All 178 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/atomics.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_32 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_32 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_32 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_32 | %ptxas-verify %}
; CHECK-LABEL: atom0 ; CHECK-LABEL: atom0
define i32 @atom0(ptr %addr, i32 %val) { define i32 @atom0(ptr %addr, i32 %val) {
; CHECK: atom.add.u32 ; CHECK: atom.add.u32
%ret = atomicrmw add ptr %addr, i32 %val seq_cst %ret = atomicrmw add ptr %addr, i32 %val seq_cst
ret i32 %ret ret i32 %ret
} }
▲ Show 20 Lines • Show All 194 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/bf16.ll

; RUN: llc < %s -march=nvptx | FileCheck %s ; RUN: llc < %s -march=nvptx64 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 | %ptxas-verify %}
; LDST: .b8 bfloat_array[8] = {1, 2, 3, 4, 5, 6, 7, 8}; ; LDST: .b8 bfloat_array[8] = {1, 2, 3, 4, 5, 6, 7, 8};
@"bfloat_array" = addrspace(1) constant [4 x bfloat] @"bfloat_array" = addrspace(1) constant [4 x bfloat]
[bfloat 0xR0201, bfloat 0xR0403, bfloat 0xR0605, bfloat 0xR0807] [bfloat 0xR0201, bfloat 0xR0403, bfloat 0xR0605, bfloat 0xR0807]
define void @test_load_store(ptr addrspace(1) %in, ptr addrspace(1) %out) { define void @test_load_store(ptr addrspace(1) %in, ptr addrspace(1) %out) {
; CHECK-LABEL: @test_load_store ; CHECK-LABEL: @test_load_store
; CHECK: ld.global.b16 [[TMP:%rs[0-9]+]], [{{%r[0-9]+}}] ; CHECK: ld.global.b16 [[TMP:%rs[0-9]+]], [{{%rd[0-9]+}}]
; CHECK: st.global.b16 [{{%r[0-9]+}}], [[TMP]] ; CHECK: st.global.b16 [{{%rd[0-9]+}}], [[TMP]]
%val = load bfloat, ptr addrspace(1) %in %val = load bfloat, ptr addrspace(1) %in
store bfloat %val, ptr addrspace(1) %out store bfloat %val, ptr addrspace(1) %out
ret void ret void
} }
define void @test_bitcast_from_bfloat(ptr addrspace(1) %in, ptr addrspace(1) %out) { define void @test_bitcast_from_bfloat(ptr addrspace(1) %in, ptr addrspace(1) %out) {
; CHECK-LABEL: @test_bitcast_from_bfloat ; CHECK-LABEL: @test_bitcast_from_bfloat
; CHECK: ld.global.b16 [[TMP:%rs[0-9]+]], [{{%r[0-9]+}}] ; CHECK: ld.global.b16 [[TMP:%rs[0-9]+]], [{{%rd[0-9]+}}]
; CHECK: st.global.b16 [{{%r[0-9]+}}], [[TMP]] ; CHECK: st.global.b16 [{{%rd[0-9]+}}], [[TMP]]
%val = load bfloat, ptr addrspace(1) %in %val = load bfloat, ptr addrspace(1) %in
%val_int = bitcast bfloat %val to i16 %val_int = bitcast bfloat %val to i16
store i16 %val_int, ptr addrspace(1) %out store i16 %val_int, ptr addrspace(1) %out
ret void ret void
} }
define void @test_bitcast_to_bfloat(ptr addrspace(1) %out, ptr addrspace(1) %in) { define void @test_bitcast_to_bfloat(ptr addrspace(1) %out, ptr addrspace(1) %in) {
; CHECK-LABEL: @test_bitcast_to_bfloat ; CHECK-LABEL: @test_bitcast_to_bfloat
; CHECK: ld.global.u16 [[TMP:%rs[0-9]+]], [{{%r[0-9]+}}] ; CHECK: ld.global.u16 [[TMP:%rs[0-9]+]], [{{%rd[0-9]+}}]
; CHECK: st.global.u16 [{{%r[0-9]+}}], [[TMP]] ; CHECK: st.global.u16 [{{%rd[0-9]+}}], [[TMP]]
%val = load i16, ptr addrspace(1) %in %val = load i16, ptr addrspace(1) %in
%val_fp = bitcast i16 %val to bfloat %val_fp = bitcast i16 %val to bfloat
store bfloat %val_fp, ptr addrspace(1) %out store bfloat %val_fp, ptr addrspace(1) %out
ret void ret void
} }

llvm/test/CodeGen/NVPTX/bfe.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; CHECK: bfe0 ; CHECK: bfe0
define i32 @bfe0(i32 %a) { define i32 @bfe0(i32 %a) {
; CHECK: bfe.u32 %r{{[0-9]+}}, %r{{[0-9]+}}, 4, 4 ; CHECK: bfe.u32 %r{{[0-9]+}}, %r{{[0-9]+}}, 4, 4
; CHECK-NOT: shr ; CHECK-NOT: shr
; CHECK-NOT: and ; CHECK-NOT: and
%val0 = ashr i32 %a, 4 %val0 = ashr i32 %a, 4
Show All 23 Lines

llvm/test/CodeGen/NVPTX/bug17709.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
; ModuleID = '__kernelgen_main_module' ; ModuleID = '__kernelgen_main_module'
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 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-nvidia-cuda" target triple = "nvptx64-nvidia-cuda"
define private ptx_device { double, double } @__utils1_MOD_trace(ptr noalias %m) { define private ptx_device { double, double } @__utils1_MOD_trace(ptr noalias %m) {
entry: entry:
;unreachable ;unreachable
%t0 = insertvalue {double, double} undef, double 1.0, 0 %t0 = insertvalue {double, double} undef, double 1.0, 0
%t1 = insertvalue {double, double} %t0, double 1.0, 1 %t1 = insertvalue {double, double} %t0, double 1.0, 1
ret { double, double } %t1 ret { double, double } %t1
} }
%struct.descriptor_dimension.0.52 = type { i64, i64, i64 } %struct.descriptor_dimension.0.52 = type { i64, i64, i64 }
%"struct.array2_complex(kind=8).37.18.70" = type { ptr, i64, i64, [2 x %struct.descriptor_dimension.0.52] } %"struct.array2_complex(kind=8).37.18.70" = type { ptr, i64, i64, [2 x %struct.descriptor_dimension.0.52] }
%"struct.array2_complex(kind=8).43.5.57" = type { ptr, i64, i64, [2 x %struct.descriptor_dimension.0.52] } %"struct.array2_complex(kind=8).43.5.57" = type { ptr, i64, i64, [2 x %struct.descriptor_dimension.0.52] }
@replacementOfAlloca8 = private global %"struct.array2_complex(kind=8).37.18.70" zeroinitializer, align 4096 @replacementOfAlloca8 = private global %"struct.array2_complex(kind=8).37.18.70" zeroinitializer, align 4096
; CHECK: .visible .entry __kernelgen_main ; CHECK: .visible .entry __kernelgen_main
define ptx_kernel void @__kernelgen_main(ptr nocapture %args, ptr) { define ptx_kernel void @__kernelgen_main(ptr nocapture %args, ptr) {
entry: entry:
%1 = tail call ptx_device { double, double } @__utils1_MOD_trace(ptr noalias @replacementOfAlloca8) %1 = tail call ptx_device { double, double } @__utils1_MOD_trace(ptr noalias @replacementOfAlloca8)
ret void ret void
} }

llvm/test/CodeGen/NVPTX/bug22246.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64" target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
target triple = "nvptx64-nvidia-cuda" target triple = "nvptx64-nvidia-cuda"
; CHECK-LABEL: _Z3foobbbPb ; CHECK-LABEL: _Z3foobbbPb
define void @_Z3foobbbPb(i1 zeroext %p1, i1 zeroext %p2, i1 zeroext %p3, ptr nocapture %output) { define void @_Z3foobbbPb(i1 zeroext %p1, i1 zeroext %p2, i1 zeroext %p3, ptr nocapture %output) {
entry: entry:
; CHECK: selp.b32 %r{{[0-9]+}}, %r{{[0-9]+}}, %r{{[0-9]+}}, %p{{[0-9]+}} ; CHECK: selp.b32 %r{{[0-9]+}}, %r{{[0-9]+}}, %r{{[0-9]+}}, %p{{[0-9]+}}
%.sink.v = select i1 %p1, i1 %p2, i1 %p3 %.sink.v = select i1 %p1, i1 %p2, i1 %p3
%frombool5 = zext i1 %.sink.v to i8 %frombool5 = zext i1 %.sink.v to i8
store i8 %frombool5, ptr %output, align 1 store i8 %frombool5, ptr %output, align 1
ret void ret void
} }

llvm/test/CodeGen/NVPTX/bug22322.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64" target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
target triple = "nvptx64-nvidia-cuda" target triple = "nvptx64-nvidia-cuda"
%class.float3 = type { float, float, float } %class.float3 = type { float, float, float }
; Function Attrs: nounwind ; Function Attrs: nounwind
; CHECK-LABEL: some_kernel ; CHECK-LABEL: some_kernel
▲ Show 20 Lines • Show All 52 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/bug26185-2.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_35 -verify-machineinstrs | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 -verify-machineinstrs | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_35 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_35 | %ptxas-verify %}
; Verify that we correctly emit code for extending ldg/ldu. We do not expose ; Verify that we correctly emit code for extending ldg/ldu. We do not expose
; extending variants in the backend, but the ldg/ldu selection code may pick ; extending variants in the backend, but the ldg/ldu selection code may pick
; extending loads as candidates. We do want to support this, so make sure we ; extending loads as candidates. We do want to support this, so make sure we
; emit the necessary cvt.* instructions to implement the extension and let ptxas ; emit the necessary cvt.* instructions to implement the extension and let ptxas
; emit the real extending loads. ; emit the real extending loads.
target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64" target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
Show All 23 Lines

llvm/test/CodeGen/NVPTX/bug26185.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_35 -verify-machineinstrs | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 -verify-machineinstrs | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_35 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_35 | %ptxas-verify %}
; Verify that we correctly emit code for i8 ldg/ldu. We do not expose 8-bit ; Verify that we correctly emit code for i8 ldg/ldu. We do not expose 8-bit
; registers in the backend, so these loads need special handling. ; registers in the backend, so these loads need special handling.
target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64" target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
target triple = "nvptx64-unknown-unknown" target triple = "nvptx64-unknown-unknown"
; CHECK-LABEL: ex_zext ; CHECK-LABEL: ex_zext
▲ Show 20 Lines • Show All 48 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/bug52623.ll

; RUN: llc < %s -march=nvptx -verify-machineinstrs ; RUN: llc < %s -march=nvptx64 -verify-machineinstrs
; RUN: %if ptxas %{ llc < %s -march=nvptx | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 | %ptxas-verify %}
; Check that llc will not crash even when first MBB doesn't contain ; Check that llc will not crash even when first MBB doesn't contain
; any instruction. ; any instruction.
target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64" target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64"
target triple = "nvptx64" target triple = "nvptx64"
%printf_args.0.8 = type { ptr } %printf_args.0.8 = type { ptr }
Show All 15 Lines

llvm/test/CodeGen/NVPTX/bypass-div.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_35 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_35 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_35 | %ptxas-verify %}
; 64-bit divides and rems should be split into a fast and slow path where ; 64-bit divides and rems should be split into a fast and slow path where
; the fast path uses a 32-bit operation. ; the fast path uses a 32-bit operation.
define void @sdiv64(i64 %a, i64 %b, ptr %retptr) { define void @sdiv64(i64 %a, i64 %b, ptr %retptr) {
; CHECK-LABEL: sdiv64( ; CHECK-LABEL: sdiv64(
; CHECK: div.s64 ; CHECK: div.s64
; CHECK: div.u32 ; CHECK: div.u32
▲ Show 20 Lines • Show All 71 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/call_bitcast_byval.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_50 -verify-machineinstrs | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_50 -verify-machineinstrs | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_50 -verify-machineinstrs | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_50 -verify-machineinstrs | %ptxas-verify %}
; calls with a bitcasted function symbol should be fine, but in combination with ; calls with a bitcasted function symbol should be fine, but in combination with
; a byval attribute were causing a segfault during isel. This testcase was ; a byval attribute were causing a segfault during isel. This testcase was
; reduced from a SYCL kernel using aggregate types which ended up being passed ; reduced from a SYCL kernel using aggregate types which ended up being passed
; `byval` ; `byval`
target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64" target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64"
target triple = "nvptx64-nvidia-cuda" target triple = "nvptx64-nvidia-cuda"
Show All 37 Lines

llvm/test/CodeGen/NVPTX/callchain.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
target triple = "nvptx" target triple = "nvptx"
define void @foo(ptr %ptr) { define void @foo(ptr %ptr) {
; CHECK: prototype_0 : .callprototype ()_ () ; CHECK: prototype_0 : .callprototype ()_ ()
tail call void %ptr() tail call void %ptr()
ret void ret void
} }

llvm/test/CodeGen/NVPTX/calling-conv.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
;; Kernel function using ptx_kernel calling conv ;; Kernel function using ptx_kernel calling conv
; CHECK: .entry kernel_func ; CHECK: .entry kernel_func
define ptx_kernel void @kernel_func(ptr %a) { define ptx_kernel void @kernel_func(ptr %a) {
; CHECK: ret ; CHECK: ret
Show All 21 Lines

llvm/test/CodeGen/NVPTX/calls-with-phi.ll

; RUN: llc < %s -march=nvptx 2>&1 | FileCheck %s ; RUN: llc < %s -march=nvptx64 2>&1 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 | %ptxas-verify %}
; Make sure the example doesn't crash with segfault ; Make sure the example doesn't crash with segfault
; CHECK: .visible .func ({{.*}}) loop ; CHECK: .visible .func ({{.*}}) loop
define i32 @loop(i32, i32) { define i32 @loop(i32, i32) {
entry: entry:
br label %loop br label %loop
Show All 14 Lines

llvm/test/CodeGen/NVPTX/combine-min-max.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 -O2 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -O2 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -O2 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -O2 | %ptxas-verify %}
; ************************************* ; *************************************
; * Cases with no min/max ; * Cases with no min/max
define i32 @ab_eq_i32(i32 %a, i32 %b) { define i32 @ab_eq_i32(i32 %a, i32 %b) {
; CHECK-LABEL: @ab_eq_i32 ; CHECK-LABEL: @ab_eq_i32
; CHECK-NOT: min ; CHECK-NOT: min
; CHECK-NOT: max ; CHECK-NOT: max
▲ Show 20 Lines • Show All 418 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/compare-int.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
;; These tests should run for all targets ;; These tests should run for all targets
;;===-- Basic instruction selection tests ---------------------------------===;; ;;===-- Basic instruction selection tests ---------------------------------===;;
;;; i64 ;;; i64
▲ Show 20 Lines • Show All 378 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/constant-vectors.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
target triple = "nvptx-nvidia-cuda" target triple = "nvptx-nvidia-cuda"
; CHECK: .visible .global .align 16 .b8 testArray[8] = {0, 1, 2, 3, 4, 5, 6, 7}; ; CHECK: .visible .global .align 16 .b8 testArray[8] = {0, 1, 2, 3, 4, 5, 6, 7};
@testArray = constant [2 x <4 x i8>] [<4 x i8> <i8 0, i8 1, i8 2, i8 3>, <4 x i8> <i8 4, i8 5, i8 6, i8 7>], align 16 @testArray = constant [2 x <4 x i8>] [<4 x i8> <i8 0, i8 1, i8 2, i8 3>, <4 x i8> <i8 4, i8 5, i8 6, i8 7>], align 16

llvm/test/CodeGen/NVPTX/convert-fp.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
define i16 @cvt_u16_f32(float %x) { define i16 @cvt_u16_f32(float %x) {
; CHECK: cvt.rzi.u16.f32 %rs{{[0-9]+}}, %f{{[0-9]+}}; ; CHECK: cvt.rzi.u16.f32 %rs{{[0-9]+}}, %f{{[0-9]+}};
; CHECK: ret; ; CHECK: ret;
%a = fptoui float %x to i16 %a = fptoui float %x to i16
ret i16 %a ret i16 %a
} }
▲ Show 20 Lines • Show All 154 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/convert-int-sm20.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
;; Integer conversions happen inplicitly by loading/storing the proper types ;; Integer conversions happen inplicitly by loading/storing the proper types
; i16 ; i16
▲ Show 20 Lines • Show All 55 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/ctlz.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
target datalayout = "e-p:32:32:32-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 datalayout = "e-p:32:32:32-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"
declare i16 @llvm.ctlz.i16(i16, i1) readnone declare i16 @llvm.ctlz.i16(i16, i1) readnone
declare i32 @llvm.ctlz.i32(i32, i1) readnone declare i32 @llvm.ctlz.i32(i32, i1) readnone
declare i64 @llvm.ctlz.i64(i64, i1) readnone declare i64 @llvm.ctlz.i64(i64, i1) readnone
; There should be no difference between llvm.ctlz.i32(%a, true) and ; There should be no difference between llvm.ctlz.i32(%a, true) and
▲ Show 20 Lines • Show All 123 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/ctpop.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
target datalayout = "e-p:32:32:32-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 datalayout = "e-p:32:32:32-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"
define i32 @myctpop(i32 %a) { define i32 @myctpop(i32 %a) {
; CHECK: popc.b32 ; CHECK: popc.b32
%val = tail call i32 @llvm.ctpop.i32(i32 %a) %val = tail call i32 @llvm.ctpop.i32(i32 %a)
ret i32 %val ret i32 %val
} }
Show All 16 Lines

llvm/test/CodeGen/NVPTX/cttz.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
target datalayout = "e-p:32:32:32-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 datalayout = "e-p:32:32:32-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"
declare i16 @llvm.cttz.i16(i16, i1) readnone declare i16 @llvm.cttz.i16(i16, i1) readnone
declare i32 @llvm.cttz.i32(i32, i1) readnone declare i32 @llvm.cttz.i32(i32, i1) readnone
declare i64 @llvm.cttz.i64(i64, i1) readnone declare i64 @llvm.cttz.i64(i64, i1) readnone
define i32 @myctpop(i32 %a) { define i32 @myctpop(i32 %a) {
Show All 35 Lines

llvm/test/CodeGen/NVPTX/disable-opt.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 -O0 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -O0 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -O0 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -O0 | %ptxas-verify %}
define void @foo(ptr %output) { define void @foo(ptr %output) {
; CHECK-LABEL: .visible .func foo( ; CHECK-LABEL: .visible .func foo(
entry: entry:
%local = alloca i32 %local = alloca i32
; CHECK: __local_depot ; CHECK: __local_depot
store i32 1, ptr %local store i32 1, ptr %local
%0 = load i32, ptr %local %0 = load i32, ptr %local
store i32 %0, ptr %output store i32 %0, ptr %output
ret void ret void
} }

llvm/test/CodeGen/NVPTX/div-ri.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 -nvptx-prec-divf32=0 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -nvptx-prec-divf32=0 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -nvptx-prec-divf32=0 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -nvptx-prec-divf32=0 | %ptxas-verify %}
define float @foo(float %a) { define float @foo(float %a) {
; CHECK: div.approx.f32 ; CHECK: div.approx.f32
%div = fdiv float %a, 13.0 %div = fdiv float %a, 13.0
ret float %div ret float %div
} }

llvm/test/CodeGen/NVPTX/divrem-combine.ll

; RUN: llc -O2 < %s -march=nvptx -mcpu=sm_35 | FileCheck %s --check-prefix=O2 --check-prefix=CHECK ; RUN: llc -O2 < %s -march=nvptx64 -mcpu=sm_35 | FileCheck %s --check-prefix=O2 --check-prefix=CHECK
; RUN: llc -O0 < %s -march=nvptx -mcpu=sm_35 | FileCheck %s --check-prefix=O0 --check-prefix=CHECK ; RUN: llc -O0 < %s -march=nvptx64 -mcpu=sm_35 | FileCheck %s --check-prefix=O0 --check-prefix=CHECK
; RUN: %if ptxas %{ llc -O2 < %s -march=nvptx -mcpu=sm_35 | %ptxas-verify %} ; RUN: %if ptxas %{ llc -O2 < %s -march=nvptx64 -mcpu=sm_35 | %ptxas-verify %}
; RUN: %if ptxas %{ llc -O0 < %s -march=nvptx -mcpu=sm_35 | %ptxas-verify %} ; RUN: %if ptxas %{ llc -O0 < %s -march=nvptx64 -mcpu=sm_35 | %ptxas-verify %}
; The following IR ; The following IR
; ;
; quot = n / d ; quot = n / d
; rem = n % d ; rem = n % d
; ;
; should be transformed into ; should be transformed into
; ;
▲ Show 20 Lines • Show All 102 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/envreg.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
declare i32 @llvm.nvvm.read.ptx.sreg.envreg0() declare i32 @llvm.nvvm.read.ptx.sreg.envreg0()
declare i32 @llvm.nvvm.read.ptx.sreg.envreg1() declare i32 @llvm.nvvm.read.ptx.sreg.envreg1()
declare i32 @llvm.nvvm.read.ptx.sreg.envreg2() declare i32 @llvm.nvvm.read.ptx.sreg.envreg2()
declare i32 @llvm.nvvm.read.ptx.sreg.envreg3() declare i32 @llvm.nvvm.read.ptx.sreg.envreg3()
declare i32 @llvm.nvvm.read.ptx.sreg.envreg4() declare i32 @llvm.nvvm.read.ptx.sreg.envreg4()
declare i32 @llvm.nvvm.read.ptx.sreg.envreg5() declare i32 @llvm.nvvm.read.ptx.sreg.envreg5()
▲ Show 20 Lines • Show All 130 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/fast-math.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
declare float @llvm.sqrt.f32(float) declare float @llvm.sqrt.f32(float)
declare double @llvm.sqrt.f64(double) declare double @llvm.sqrt.f64(double)
; CHECK-LABEL: sqrt_div( ; CHECK-LABEL: sqrt_div(
; CHECK: sqrt.rn.f32 ; CHECK: sqrt.rn.f32
; CHECK: div.rn.f32 ; CHECK: div.rn.f32
define float @sqrt_div(float %a, float %b) { define float @sqrt_div(float %a, float %b) {
▲ Show 20 Lines • Show All 256 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/fma-assoc.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 -fp-contract=fast | FileCheck %s -check-prefix=CHECK ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -fp-contract=fast | FileCheck %s -check-prefix=CHECK
; RUN: llc < %s -march=nvptx -mcpu=sm_20 -fp-contract=fast -enable-unsafe-fp-math | FileCheck %s -check-prefix=CHECK -check-prefix=CHECK-UNSAFE ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -fp-contract=fast -enable-unsafe-fp-math | FileCheck %s -check-prefix=CHECK -check-prefix=CHECK-UNSAFE
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -fp-contract=fast | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -fp-contract=fast | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -fp-contract=fast -enable-unsafe-fp-math | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -fp-contract=fast -enable-unsafe-fp-math | %ptxas-verify %}
define ptx_device float @t1_f32(float %x, float %y, float %z, define ptx_device float @t1_f32(float %x, float %y, float %z,
float %u, float %v) { float %u, float %v) {
; CHECK-UNSAFE: fma.rn.f32 %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}; ; CHECK-UNSAFE: fma.rn.f32 %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}};
; CHECK-UNSAFE: fma.rn.f32 %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}; ; CHECK-UNSAFE: fma.rn.f32 %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}};
; CHECK: ret; ; CHECK: ret;
%a = fmul float %x, %y %a = fmul float %x, %y
%b = fmul float %u, %v %b = fmul float %u, %v
Show All 29 Lines

llvm/test/CodeGen/NVPTX/fma-disable.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 -nvptx-fma-level=1 | FileCheck %s -check-prefix=FMA ; RUN: llc < %s -march=nvptx -mcpu=sm_20 -nvptx-fma-level=1 | FileCheck %s -check-prefix=FMA
; RUN: llc < %s -march=nvptx -mcpu=sm_20 -nvptx-fma-level=0 | FileCheck %s -check-prefix=MUL ; RUN: llc < %s -march=nvptx -mcpu=sm_20 -nvptx-fma-level=0 | FileCheck %s -check-prefix=MUL
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -nvptx-fma-level=1 | FileCheck %s -check-prefix=FMA ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -nvptx-fma-level=1 | FileCheck %s -check-prefix=FMA
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -nvptx-fma-level=0 | FileCheck %s -check-prefix=MUL ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -nvptx-fma-level=0 | FileCheck %s -check-prefix=MUL
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -nvptx-fma-level=1 | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_20 -nvptx-fma-level=1 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -nvptx-fma-level=0 | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_20 -nvptx-fma-level=0 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -nvptx-fma-level=1 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -nvptx-fma-level=1 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -nvptx-fma-level=0 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -nvptx-fma-level=0 | %ptxas-verify %}
define ptx_device float @test_mul_add_f(float %x, float %y, float %z) { define ptx_device float @test_mul_add_f(float %x, float %y, float %z) {
entry: entry:
; FMA: fma.rn.f32 ; FMA: fma.rn.f32
; MUL: mul.rn.f32 ; MUL: mul.rn.f32
; MUL: add.rn.f32 ; MUL: add.rn.f32
Show All 14 Lines

llvm/test/CodeGen/NVPTX/fma.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 -fp-contract=fast -verify-machineinstrs | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -fp-contract=fast -verify-machineinstrs | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -fp-contract=fast -verify-machineinstrs | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -fp-contract=fast -verify-machineinstrs | %ptxas-verify %}
declare float @dummy_f32(float, float) #0 declare float @dummy_f32(float, float) #0
declare double @dummy_f64(double, double) #0 declare double @dummy_f64(double, double) #0
define ptx_device float @t1_f32(float %x, float %y, float %z) { define ptx_device float @t1_f32(float %x, float %y, float %z) {
; CHECK: fma.rn.f32 %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}; ; CHECK: fma.rn.f32 %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}};
; CHECK: ret; ; CHECK: ret;
%a = fmul float %x, %y %a = fmul float %x, %y
%b = fadd float %a, %z %b = fadd float %a, %z
ret float %b ret float %b
} }
define ptx_device float @t2_f32(float %x, float %y, float %z, float %w) { define ptx_device float @t2_f32(float %x, float %y, float %z, float %w) {
; CHECK: fma.rn.f32 %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}; ; CHECK: fma.rn.f32 %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}};
; CHECK: fma.rn.f32 %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}; ; CHECK: fma.rn.f32 %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}}, %f{{[0-9]+}};
; CHECK: ret; ; CHECK: ret;
%a = fmul float %x, %y %a = fmul float %x, %y
%b = fadd float %a, %z %b = fadd float %a, %z
%c = fadd float %a, %w %c = fadd float %a, %w
%d = call float @dummy_f32(float %b, float %c) %d = call float @dummy_f32(float %b, float %c)
ret float %d ret float %d
} }
define ptx_device double @t1_f64(double %x, double %y, double %z) { define ptx_device double @t1_f64(double %x, double %y, double %z) {
; CHECK: fma.rn.f64 %fd{{[0-9]+}}, %fd{{[0-9]+}}, %fd{{[0-9]+}}, %fd{{[0-9]+}}; ; CHECK: fma.rn.f64 %fd{{[0-9]+}}, %fd{{[0-9]+}}, %fd{{[0-9]+}}, %fd{{[0-9]+}};
; CHECK: ret; ; CHECK: ret;
%a = fmul double %x, %y %a = fmul double %x, %y
%b = fadd double %a, %z %b = fadd double %a, %z
ret double %b ret double %b
} }
define ptx_device double @t2_f64(double %x, double %y, double %z, double %w) { define ptx_device double @t2_f64(double %x, double %y, double %z, double %w) {
; CHECK: fma.rn.f64 %fd{{[0-9]+}}, %fd{{[0-9]+}}, %fd{{[0-9]+}}, %fd{{[0-9]+}}; ; CHECK: fma.rn.f64 %fd{{[0-9]+}}, %fd{{[0-9]+}}, %fd{{[0-9]+}}, %fd{{[0-9]+}};
; CHECK: fma.rn.f64 %fd{{[0-9]+}}, %fd{{[0-9]+}}, %fd{{[0-9]+}}, %fd{{[0-9]+}}; ; CHECK: fma.rn.f64 %fd{{[0-9]+}}, %fd{{[0-9]+}}, %fd{{[0-9]+}}, %fd{{[0-9]+}};
; CHECK: ret; ; CHECK: ret;
%a = fmul double %x, %y %a = fmul double %x, %y
%b = fadd double %a, %z %b = fadd double %a, %z
%c = fadd double %a, %w %c = fadd double %a, %w
%d = call double @dummy_f64(double %b, double %c) %d = call double @dummy_f64(double %b, double %c)
ret double %d ret double %d
} }

llvm/test/CodeGen/NVPTX/fminimum-fmaximum.ll

; RUN: llc < %s -march=nvptx | FileCheck %s --check-prefixes=CHECK ; RUN: llc < %s -march=nvptx64 | FileCheck %s --check-prefixes=CHECK
; RUN: llc < %s -march=nvptx -mcpu=sm_80 | FileCheck %s --check-prefixes=CHECK ; RUN: llc < %s -march=nvptx64 -mcpu=sm_80 | FileCheck %s --check-prefixes=CHECK
; RUN: %if ptxas %{ llc < %s -march=nvptx | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 | %ptxas-verify %}
; RUN: %if ptxas-11.0 %{ llc < %s -march=nvptx -mcpu=sm_80 | %ptxas-verify -arch=sm_80 %} ; RUN: %if ptxas-11.0 %{ llc < %s -march=nvptx64 -mcpu=sm_80 | %ptxas-verify -arch=sm_80 %}
; ---- minimum ---- ; ---- minimum ----
; CHECK-LABEL: minimum_half ; CHECK-LABEL: minimum_half
define half @minimum_half(half %a) #0 { define half @minimum_half(half %a) #0 {
; CHECK: setp ; CHECK: setp
; CHECK: selp.b16 ; CHECK: selp.b16
%p = fcmp ult half %a, 0.0 %p = fcmp ult half %a, 0.0
▲ Show 20 Lines • Show All 70 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/fp-literals.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 -fp-contract=fast | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -fp-contract=fast | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -fp-contract=fast | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -fp-contract=fast | %ptxas-verify %}
target triple = "nvptx64-unknown-cuda" target triple = "nvptx64-unknown-cuda"
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 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"
; Make sure we can properly differentiate between single-precision and ; Make sure we can properly differentiate between single-precision and
; double-precision FP literals. ; double-precision FP literals.
; CHECK: myaddf ; CHECK: myaddf
Show All 12 Lines

llvm/test/CodeGen/NVPTX/fp16.ll

; RUN: llc -march=nvptx -verify-machineinstrs < %s | FileCheck %s ; RUN: llc -march=nvptx64 -verify-machineinstrs < %s | FileCheck %s
; RUN: %if ptxas %{ llc -march=nvptx -verify-machineinstrs < %s | %ptxas-verify %} ; RUN: %if ptxas %{ llc -march=nvptx64 -verify-machineinstrs < %s | %ptxas-verify %}
declare float @llvm.convert.from.fp16.f32(i16) nounwind readnone declare float @llvm.convert.from.fp16.f32(i16) nounwind readnone
declare double @llvm.convert.from.fp16.f64(i16) nounwind readnone declare double @llvm.convert.from.fp16.f64(i16) nounwind readnone
declare i16 @llvm.convert.to.fp16.f32(float) nounwind readnone declare i16 @llvm.convert.to.fp16.f32(float) nounwind readnone
declare i16 @llvm.convert.to.fp16.f64(double) nounwind readnone declare i16 @llvm.convert.to.fp16.f64(double) nounwind readnone
; CHECK-LABEL: @test_convert_fp16_to_fp32 ; CHECK-LABEL: @test_convert_fp16_to_fp32
; CHECK: cvt.f32.f16 ; CHECK: cvt.f32.f16
Show All 36 Lines

llvm/test/CodeGen/NVPTX/function-align.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; CHECK-NOT: .align 2 ; CHECK-NOT: .align 2
define ptx_device void @foo() align 2 { define ptx_device void @foo() align 2 {
; CHECK-LABEL: .func foo ; CHECK-LABEL: .func foo
ret void ret void
} }

llvm/test/CodeGen/NVPTX/generic-to-nvvm.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
target datalayout = "e-p:32:32:32-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 datalayout = "e-p:32:32:32-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 = "nvptx-nvidia-cuda" target triple = "nvptx-nvidia-cuda"
; Ensure global variables in address space 0 are promoted to address space 1 ; Ensure global variables in address space 0 are promoted to address space 1
; CHECK: .global .align 4 .u32 myglobal = 42; ; CHECK: .global .align 4 .u32 myglobal = 42;
@myglobal = internal global i32 42, align 4 @myglobal = internal global i32 42, align 4
Show All 20 Lines

llvm/test/CodeGen/NVPTX/global-addrspace.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix=PTX32 ; 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 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefix=PTX64
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; PTX32: .visible .global .align 4 .u32 i; ; PTX32: .visible .global .align 4 .u32 i;
; PTX32: .visible .const .align 4 .u32 j; ; PTX32: .visible .const .align 4 .u32 j;
; PTX32: .visible .shared .align 4 .u32 k; ; PTX32: .visible .shared .align 4 .u32 k;
; PTX64: .visible .global .align 4 .u32 i; ; PTX64: .visible .global .align 4 .u32 i;
; PTX64: .visible .const .align 4 .u32 j; ; PTX64: .visible .const .align 4 .u32 j;
; PTX64: .visible .shared .align 4 .u32 k; ; PTX64: .visible .shared .align 4 .u32 k;
@i = addrspace(1) externally_initialized global i32 0, align 4 @i = addrspace(1) externally_initialized global i32 0, align 4
@j = addrspace(4) externally_initialized global i32 0, align 4 @j = addrspace(4) externally_initialized global i32 0, align 4
@k = addrspace(3) global i32 undef, align 4 @k = addrspace(3) global i32 undef, align 4

llvm/test/CodeGen/NVPTX/global-ordering.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix=PTX32 ; 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 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefix=PTX64
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; Make sure we emit these globals in def-use order ; Make sure we emit these globals in def-use order
; PTX32: .visible .global .align 1 .u8 a = 2; ; PTX32: .visible .global .align 1 .u8 a = 2;
; PTX32-NEXT: .visible .global .align 4 .u32 a2 = a; ; PTX32-NEXT: .visible .global .align 4 .u32 a2 = a;
; PTX64: .visible .global .align 1 .u8 a = 2; ; PTX64: .visible .global .align 1 .u8 a = 2;
Show All 11 Lines

llvm/test/CodeGen/NVPTX/global-visibility.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; PTX does not support .hidden or .protected. ; PTX does not support .hidden or .protected.
; Make sure we do not emit them. ; Make sure we do not emit them.
define hidden void @f_hidden() { define hidden void @f_hidden() {
ret void ret void
} }
; CHECK-NOT: .hidden ; CHECK-NOT: .hidden
; CHECK: .visible .func f_hidden ; CHECK: .visible .func f_hidden
define protected void @f_protected() { define protected void @f_protected() {
ret void ret void
} }
; CHECK-NOT: .protected ; CHECK-NOT: .protected
; CHECK: .visible .func f_protected ; CHECK: .visible .func f_protected

llvm/test/CodeGen/NVPTX/globals_init.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; Make sure the globals constant initializers are not prone to host endianess ; Make sure the globals constant initializers are not prone to host endianess
; issues. ; issues.
; CHECK-DAG: .b8 Gbli08[2] = {171, 205}; ; CHECK-DAG: .b8 Gbli08[2] = {171, 205};
@Gbli08 = global [2 x i8] [i8 171, i8 205] @Gbli08 = global [2 x i8] [i8 171, i8 205]
; CHECK-DAG: .b8 Gbli16[4] = {205, 171, 1, 239}; ; CHECK-DAG: .b8 Gbli16[4] = {205, 171, 1, 239};
Show All 21 Lines

llvm/test/CodeGen/NVPTX/globals_lowering.ll

; RUN: llc < %s -mtriple=nvptx -mcpu=sm_20 -relocation-model=static | FileCheck %s --check-prefix CHK ; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -relocation-model=static | FileCheck %s --check-prefix CHK
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx -mcpu=sm_20 -relocation-model=static | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 -relocation-model=static | %ptxas-verify %}
%MyStruct = type { i32, i32, float } %MyStruct = type { i32, i32, float }
@Gbl = internal addrspace(3) global [1024 x %MyStruct] zeroinitializer @Gbl = internal addrspace(3) global [1024 x %MyStruct] zeroinitializer
; CHK-LABEL: foo ; CHK-LABEL: foo
define void @foo(float %f) { define void @foo(float %f) {
entry: entry:
; CHK: ld.shared.f32 %{{[a-zA-Z0-9]+}}, [Gbl+8]; ; CHK: ld.shared.f32 %{{[a-zA-Z0-9]+}}, [Gbl+8];
%0 = load float, ptr addrspace(3) getelementptr inbounds ([1024 x %MyStruct], ptr addrspace(3) @Gbl, i32 0, i32 0, i32 2) %0 = load float, ptr addrspace(3) getelementptr inbounds ([1024 x %MyStruct], ptr addrspace(3) @Gbl, i32 0, i32 0, i32 2)
%add = fadd float %0, %f %add = fadd float %0, %f
; CHK: st.shared.f32 [Gbl+8], %{{[a-zA-Z0-9]+}}; ; CHK: st.shared.f32 [Gbl+8], %{{[a-zA-Z0-9]+}};
store float %add, ptr addrspace(3) getelementptr inbounds ([1024 x %MyStruct], ptr addrspace(3) @Gbl, i32 0, i32 0, i32 2) store float %add, ptr addrspace(3) getelementptr inbounds ([1024 x %MyStruct], ptr addrspace(3) @Gbl, i32 0, i32 0, i32 2)
ret void ret void
} }

llvm/test/CodeGen/NVPTX/half.ll

; RUN: llc < %s -march=nvptx | FileCheck %s ; RUN: llc < %s -march=nvptx64 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 | %ptxas-verify %}
; CHECK: .b8 half_array[8] = {1, 2, 3, 4, 5, 6, 7, 8}; ; CHECK: .b8 half_array[8] = {1, 2, 3, 4, 5, 6, 7, 8};
@"half_array" = addrspace(1) constant [4 x half] @"half_array" = addrspace(1) constant [4 x half]
[half 0xH0201, half 0xH0403, half 0xH0605, half 0xH0807] [half 0xH0201, half 0xH0403, half 0xH0605, half 0xH0807]
define void @test_load_store(ptr addrspace(1) %in, ptr addrspace(1) %out) { define void @test_load_store(ptr addrspace(1) %in, ptr addrspace(1) %out) {
; CHECK-LABEL: @test_load_store ; CHECK-LABEL: @test_load_store
; CHECK: ld.global.b16 [[TMP:%rs[0-9]+]], [{{%r[0-9]+}}] ; CHECK: ld.global.b16 [[TMP:%rs[0-9]+]], [{{%rd[0-9]+}}]
; CHECK: st.global.b16 [{{%r[0-9]+}}], [[TMP]] ; CHECK: st.global.b16 [{{%rd[0-9]+}}], [[TMP]]
%val = load half, ptr addrspace(1) %in %val = load half, ptr addrspace(1) %in
store half %val, ptr addrspace(1) %out store half %val, ptr addrspace(1) %out
ret void ret void
} }
define void @test_bitcast_from_half(ptr addrspace(1) %in, ptr addrspace(1) %out) { define void @test_bitcast_from_half(ptr addrspace(1) %in, ptr addrspace(1) %out) {
; CHECK-LABEL: @test_bitcast_from_half ; CHECK-LABEL: @test_bitcast_from_half
; CHECK: ld.global.b16 [[TMP:%rs[0-9]+]], [{{%r[0-9]+}}] ; CHECK: ld.global.b16 [[TMP:%rs[0-9]+]], [{{%rd[0-9]+}}]
; CHECK: st.global.b16 [{{%r[0-9]+}}], [[TMP]] ; CHECK: st.global.b16 [{{%rd[0-9]+}}], [[TMP]]
%val = load half, ptr addrspace(1) %in %val = load half, ptr addrspace(1) %in
%val_int = bitcast half %val to i16 %val_int = bitcast half %val to i16
store i16 %val_int, ptr addrspace(1) %out store i16 %val_int, ptr addrspace(1) %out
ret void ret void
} }
define void @test_bitcast_to_half(ptr addrspace(1) %out, ptr addrspace(1) %in) { define void @test_bitcast_to_half(ptr addrspace(1) %out, ptr addrspace(1) %in) {
; CHECK-LABEL: @test_bitcast_to_half ; CHECK-LABEL: @test_bitcast_to_half
; CHECK: ld.global.u16 [[TMP:%rs[0-9]+]], [{{%r[0-9]+}}] ; CHECK: ld.global.u16 [[TMP:%rs[0-9]+]], [{{%rd[0-9]+}}]
; CHECK: st.global.u16 [{{%r[0-9]+}}], [[TMP]] ; CHECK: st.global.u16 [{{%rd[0-9]+}}], [[TMP]]
%val = load i16, ptr addrspace(1) %in %val = load i16, ptr addrspace(1) %in
%val_fp = bitcast i16 %val to half %val_fp = bitcast i16 %val to half
store half %val_fp, ptr addrspace(1) %out store half %val_fp, ptr addrspace(1) %out
ret void ret void
} }
define void @test_extend32(ptr addrspace(1) %in, ptr addrspace(1) %out) { define void @test_extend32(ptr addrspace(1) %in, ptr addrspace(1) %out) {
; CHECK-LABEL: @test_extend32 ; CHECK-LABEL: @test_extend32
Show All 37 Lines

llvm/test/CodeGen/NVPTX/i1-global.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
target datalayout = "e-p:32:32:32-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 datalayout = "e-p:32:32:32-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 = "nvptx-nvidia-cuda" target triple = "nvptx-nvidia-cuda"
; CHECK: .visible .global .align 1 .u8 mypred ; CHECK: .visible .global .align 1 .u8 mypred
@mypred = addrspace(1) global i1 true, align 1 @mypred = addrspace(1) global i1 true, align 1
Show All 10 Lines

llvm/test/CodeGen/NVPTX/i1-int-to-fp.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; CHECK-LABEL: foo ; CHECK-LABEL: foo
; CHECK: setp ; CHECK: setp
; CHECK: selp ; CHECK: selp
; CHECK: cvt.rn.f32.u32 ; CHECK: cvt.rn.f32.u32
define float @foo(i1 %a) { define float @foo(i1 %a) {
%ret = uitofp i1 %a to float %ret = uitofp i1 %a to float
ret float %ret ret float %ret
Show All 28 Lines

llvm/test/CodeGen/NVPTX/i1-param.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
target datalayout = "e-p:32:32:32-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 datalayout = "e-p:32:32:32-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 = "nvptx-nvidia-cuda" target triple = "nvptx-nvidia-cuda"
; Make sure predicate (i1) operands to kernels get expanded out to .u8 ; Make sure predicate (i1) operands to kernels get expanded out to .u8
; CHECK: .entry foo ; CHECK: .entry foo
; CHECK: .param .u8 foo_param_0 ; CHECK: .param .u8 foo_param_0
; CHECK: .param .u32 foo_param_1 ; CHECK: .param .u64 foo_param_1
define void @foo(i1 %p, ptr %out) { define void @foo(i1 %p, ptr %out) {
%val = zext i1 %p to i32 %val = zext i1 %p to i32
store i32 %val, ptr %out store i32 %val, ptr %out
ret void ret void
} }
!nvvm.annotations = !{!0} !nvvm.annotations = !{!0}
!0 = !{ptr @foo, !"kernel", i32 1} !0 = !{ptr @foo, !"kernel", i32 1}

llvm/test/CodeGen/NVPTX/i128-param.ll

; RUN: llc < %s -O0 -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -O0 -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -O0 -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -O0 -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; CHECK-LABEL: .visible .func callee( ; CHECK-LABEL: .visible .func callee(
; CHECK-NEXT: .param .align 16 .b8 callee_param_0[16], ; CHECK-NEXT: .param .align 16 .b8 callee_param_0[16],
; CHECK-NEXT: .param .align 16 .b8 callee_param_1[16], ; CHECK-NEXT: .param .align 16 .b8 callee_param_1[16],
define void @callee(i128, i128, ptr) { define void @callee(i128, i128, ptr) {
; CHECK-DAG: ld.param.v2.u64 {%[[REG0:rd[0-9]+]], %[[REG1:rd[0-9]+]]}, [callee_param_0]; ; CHECK-DAG: ld.param.v2.u64 {%[[REG0:rd[0-9]+]], %[[REG1:rd[0-9]+]]}, [callee_param_0];
; CHECK-DAG: ld.param.v2.u64 {%[[REG2:rd[0-9]+]], %[[REG3:rd[0-9]+]]}, [callee_param_1]; ; CHECK-DAG: ld.param.v2.u64 {%[[REG2:rd[0-9]+]], %[[REG3:rd[0-9]+]]}, [callee_param_1];
▲ Show 20 Lines • Show All 49 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/i8-param.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
target datalayout = "e-p:32:32:32-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 datalayout = "e-p:32:32:32-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"
; CHECK: .visible .func (.param .b32 func_retval0) callee ; CHECK: .visible .func (.param .b32 func_retval0) callee
define i8 @callee(i8 %a) { define i8 @callee(i8 %a) {
; CHECK: ld.param.u8 ; CHECK: ld.param.u8
%ret = add i8 %a, 42 %ret = add i8 %a, 42
; CHECK: st.param.b32 ; CHECK: st.param.b32
ret i8 %ret ret i8 %ret
} }
; CHECK: .visible .func caller ; CHECK: .visible .func caller
define void @caller(ptr %a) { define void @caller(ptr %a) {
; CHECK: ld.u8 ; CHECK: ld.u8
%val = load i8, ptr %a %val = load i8, ptr %a
%ret = tail call i8 @callee(i8 %val) %ret = tail call i8 @callee(i8 %val)
; CHECK: ld.param.b32 ; CHECK: ld.param.b32
store i8 %ret, ptr %a store i8 %ret, ptr %a
ret void ret void
} }

llvm/test/CodeGen/NVPTX/idioms.ll

; Check that various LLVM idioms get lowered to NVPTX as expected. ; Check that various LLVM idioms get lowered to NVPTX as expected.
; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
%struct.S16 = type { i16, i16 } %struct.S16 = type { i16, i16 }
%struct.S32 = type { i32, i32 } %struct.S32 = type { i32, i32 }
; CHECK-LABEL: abs_i16( ; CHECK-LABEL: abs_i16(
define i16 @abs_i16(i16 %a) { define i16 @abs_i16(i16 %a) {
; CHECK: abs.s16 ; CHECK: abs.s16
▲ Show 20 Lines • Show All 111 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/imad.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; CHECK: imad ; CHECK: imad
define i32 @imad(i32 %a, i32 %b, i32 %c) { define i32 @imad(i32 %a, i32 %b, i32 %c) {
; CHECK: mad.lo.s32 ; CHECK: mad.lo.s32
%val0 = mul i32 %a, %b %val0 = mul i32 %a, %b
%val1 = add i32 %val0, %c %val1 = add i32 %val0, %c
ret i32 %val1 ret i32 %val1
} }

llvm/test/CodeGen/NVPTX/inline-asm.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
define float @test(float %x) { define float @test(float %x) {
entry: entry:
; CHECK: ex2.approx.ftz.f32 %f{{[0-9]+}}, %f{{[0-9]+}} ; CHECK: ex2.approx.ftz.f32 %f{{[0-9]+}}, %f{{[0-9]+}}
%0 = call float asm "ex2.approx.ftz.f32 $0, $1;", "=f,f"(float %x) %0 = call float asm "ex2.approx.ftz.f32 $0, $1;", "=f,f"(float %x)
ret float %0 ret float %0
} }
define i32 @foo(i1 signext %cond, i32 %a, i32 %b) #0 { define i32 @foo(i1 signext %cond, i32 %a, i32 %b) #0 {
entry: entry:
; CHECK: selp.b32 %r{{[0-9]+}}, %r{{[0-9]+}}, %r{{[0-9]+}}, %p{{[0-9]+}} ; CHECK: selp.b32 %r{{[0-9]+}}, %r{{[0-9]+}}, %r{{[0-9]+}}, %p{{[0-9]+}}
%0 = tail call i32 asm "selp.b32 $0, $1, $2, $3;", "=r,r,r,b"(i32 %a, i32 %b, i1 %cond) %0 = tail call i32 asm "selp.b32 $0, $1, $2, $3;", "=r,r,r,b"(i32 %a, i32 %b, i1 %cond)
ret i32 %0 ret i32 %0
} }

llvm/test/CodeGen/NVPTX/inlineasm-output-template.ll

; RUN: llc -march=nvptx < %s | FileCheck %s ; RUN: llc -march=nvptx64 < %s | FileCheck %s
; RUN: %if ptxas %{ llc -march=nvptx < %s | %ptxas-verify %} ; RUN: %if ptxas %{ llc -march=nvptx64 < %s | %ptxas-verify %}
; Test that %c works with immediates ; Test that %c works with immediates
; CHECK-LABEL: test_inlineasm_c_output_template0 ; CHECK-LABEL: test_inlineasm_c_output_template0
; CHECK: //TEST 42 ; CHECK: //TEST 42
define dso_local i32 @test_inlineasm_c_output_template0() { define dso_local i32 @test_inlineasm_c_output_template0() {
tail call void asm sideeffect "//TEST ${0:c}", "i"(i32 42) tail call void asm sideeffect "//TEST ${0:c}", "i"(i32 42)
ret i32 42 ret i32 42
} }
Show All 19 Lines

llvm/test/CodeGen/NVPTX/intrinsic-old.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck -allow-deprecated-dag-overlap %s ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck -allow-deprecated-dag-overlap %s
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck -allow-deprecated-dag-overlap %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck -allow-deprecated-dag-overlap %s
; RUN: opt < %s -S -mtriple=nvptx-nvidia-cuda -passes=nvvm-intr-range \ ; RUN: opt < %s -S -mtriple=nvptx-nvidia-cuda -passes=nvvm-intr-range \
; RUN: | FileCheck -allow-deprecated-dag-overlap --check-prefix=RANGE --check-prefix=RANGE_20 %s ; RUN: | FileCheck -allow-deprecated-dag-overlap --check-prefix=RANGE --check-prefix=RANGE_20 %s
; RUN: opt < %s -S -mtriple=nvptx-nvidia-cuda -passes=nvvm-intr-range \ ; RUN: opt < %s -S -mtriple=nvptx-nvidia-cuda -passes=nvvm-intr-range \
; RUN: | FileCheck -allow-deprecated-dag-overlap --check-prefix=RANGE --check-prefix=RANGE_20 %s ; RUN: | FileCheck -allow-deprecated-dag-overlap --check-prefix=RANGE --check-prefix=RANGE_20 %s
; RUN: opt < %s -S -mtriple=nvptx-nvidia-cuda \ ; RUN: opt < %s -S -mtriple=nvptx-nvidia-cuda \
; RUN: -passes=nvvm-intr-range -nvvm-intr-range-sm=30 \ ; RUN: -passes=nvvm-intr-range -nvvm-intr-range-sm=30 \
; RUN: | FileCheck -allow-deprecated-dag-overlap --check-prefix=RANGE --check-prefix=RANGE_30 %s ; RUN: | FileCheck -allow-deprecated-dag-overlap --check-prefix=RANGE --check-prefix=RANGE_30 %s
; RUN: opt < %s -S -mtriple=nvptx-nvidia-cuda \ ; RUN: opt < %s -S -mtriple=nvptx-nvidia-cuda \
; RUN: -passes=nvvm-intr-range -nvvm-intr-range-sm=30 \ ; RUN: -passes=nvvm-intr-range -nvvm-intr-range-sm=30 \
; RUN: | FileCheck -allow-deprecated-dag-overlap --check-prefix=RANGE --check-prefix=RANGE_30 %s ; RUN: | FileCheck -allow-deprecated-dag-overlap --check-prefix=RANGE --check-prefix=RANGE_30 %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
define ptx_device i32 @test_tid_x() { define ptx_device i32 @test_tid_x() {
; CHECK: mov.u32 %r{{[0-9]+}}, %tid.x; ; CHECK: mov.u32 %r{{[0-9]+}}, %tid.x;
; RANGE: call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !range ![[BLK_IDX_XY:[0-9]+]] ; RANGE: call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !range ![[BLK_IDX_XY:[0-9]+]]
; CHECK: ret; ; CHECK: ret;
%x = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() %x = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
ret i32 %x ret i32 %x
▲ Show 20 Lines • Show All 319 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/intrinsics.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; CHECK-LABEL: test_fabsf( ; CHECK-LABEL: test_fabsf(
define float @test_fabsf(float %f) { define float @test_fabsf(float %f) {
; CHECK: abs.f32 ; CHECK: abs.f32
%x = call float @llvm.fabs.f32(float %f) %x = call float @llvm.fabs.f32(float %f)
ret float %x ret float %x
} }
▲ Show 20 Lines • Show All 137 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/isspacep.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
declare i1 @llvm.nvvm.isspacep.const(ptr) readnone noinline declare i1 @llvm.nvvm.isspacep.const(ptr) readnone noinline
declare i1 @llvm.nvvm.isspacep.global(ptr) readnone noinline declare i1 @llvm.nvvm.isspacep.global(ptr) readnone noinline
declare i1 @llvm.nvvm.isspacep.local(ptr) readnone noinline declare i1 @llvm.nvvm.isspacep.local(ptr) readnone noinline
declare i1 @llvm.nvvm.isspacep.shared(ptr) readnone noinline declare i1 @llvm.nvvm.isspacep.shared(ptr) readnone noinline
; CHECK: is_const ; CHECK: is_const
define i1 @is_const(ptr %addr) { define i1 @is_const(ptr %addr) {
Show All 26 Lines

llvm/test/CodeGen/NVPTX/ld-addrspace.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefixes=ALL,G32,LS32 ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefixes=ALL,G32,LS32
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefixes=ALL,G64,LS64 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefixes=ALL,G64,LS64
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 --nvptx-short-ptr | FileCheck %s --check-prefixes=G64,LS32 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 --nvptx-short-ptr | FileCheck %s --check-prefixes=G64,LS32
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 --nvptx-short-ptr | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 --nvptx-short-ptr | %ptxas-verify %}
;; i8 ;; i8
define i8 @ld_global_i8(i8 addrspace(1)* %ptr) { define i8 @ld_global_i8(i8 addrspace(1)* %ptr) {
; ALL-LABEL: ld_global_i8 ; ALL-LABEL: ld_global_i8
; G32: ld.global.u8 %{{.*}}, [%r{{[0-9]+}}] ; G32: ld.global.u8 %{{.*}}, [%r{{[0-9]+}}]
▲ Show 20 Lines • Show All 151 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/ld-generic.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix=PTX32 ; 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 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefix=PTX64
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
;; i8 ;; i8
define i8 @ld_global_i8(ptr addrspace(0) %ptr) { define i8 @ld_global_i8(ptr addrspace(0) %ptr) {
; PTX32: ld.u8 %r{{[0-9]+}}, [%r{{[0-9]+}}] ; PTX32: ld.u8 %r{{[0-9]+}}, [%r{{[0-9]+}}]
; PTX32: ret ; PTX32: ret
; PTX64: ld.u8 %r{{[0-9]+}}, [%rd{{[0-9]+}}] ; PTX64: ld.u8 %r{{[0-9]+}}, [%rd{{[0-9]+}}]
▲ Show 20 Lines • Show All 54 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/ld-st-addrrspace.py

# This test generates all variants of load/store instructions and verifies that # This test generates all variants of load/store instructions and verifies that
# LLVM generates correct PTX for them. # LLVM generates correct PTX for them.
# RUN: %python %s > %t.ll # RUN: %python %s > %t.ll
# RUN: llc < %t.ll -march=nvptx64 -mcpu=sm_30 | FileCheck -check-prefixes=CHECK,CHECK_P64 %t.ll
# RUN: llc < %t.ll -march=nvptx -mcpu=sm_30 | FileCheck -check-prefixes=CHECK,CHECK_P32 %t.ll # RUN: llc < %t.ll -march=nvptx -mcpu=sm_30 | FileCheck -check-prefixes=CHECK,CHECK_P32 %t.ll
# RUN: llc < %t.ll -march=nvptx64 -mcpu=sm_30 | FileCheck -check-prefixes=CHECK,CHECK_P64 %t.ll
# RUN: %if ptxas && !ptxas-12.0 %{ llc < %t.ll -march=nvptx -mcpu=sm_30 | %ptxas-verify %}
# RUN: %if ptxas %{ llc < %t.ll -march=nvptx64 -mcpu=sm_30 | %ptxas-verify %} # RUN: %if ptxas %{ llc < %t.ll -march=nvptx64 -mcpu=sm_30 | %ptxas-verify %}
# RUN: %if ptxas %{ llc < %t.ll -march=nvptx -mcpu=sm_30 | %ptxas-verify %}
from __future__ import print_function from __future__ import print_function
from itertools import product from itertools import product
from string import Template from string import Template
llvm_type_to_ptx_type = { llvm_type_to_ptx_type = {
"i8": "u8", "i8": "u8",
▲ Show 20 Lines • Show All 91 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/ldparam-v4.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
declare <4 x float> @bar() declare <4 x float> @bar()
; CHECK-LABEL: .func foo( ; CHECK-LABEL: .func foo(
define void @foo(ptr %ptr) { define void @foo(ptr %ptr) {
; CHECK: ld.param.u32 %[[PTR:r[0-9]+]], [foo_param_0]; ; CHECK: ld.param.u64 %[[PTR:rd[0-9]+]], [foo_param_0];
; CHECK: ld.param.v4.f32 {[[E0:%f[0-9]+]], [[E1:%f[0-9]+]], [[E2:%f[0-9]+]], [[E3:%f[0-9]+]]}, [retval0+0]; ; CHECK: ld.param.v4.f32 {[[E0:%f[0-9]+]], [[E1:%f[0-9]+]], [[E2:%f[0-9]+]], [[E3:%f[0-9]+]]}, [retval0+0];
; CHECK: st.v4.f32 [%[[PTR]]], {[[E0]], [[E1]], [[E2]], [[E3]]} ; CHECK: st.v4.f32 [%[[PTR]]], {[[E0]], [[E1]], [[E2]], [[E3]]}
%val = tail call <4 x float> @bar() %val = tail call <4 x float> @bar()
store <4 x float> %val, ptr %ptr store <4 x float> %val, ptr %ptr
ret void ret void
} }

llvm/test/CodeGen/NVPTX/ldu-i8.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
target datalayout = "e-p:32:32:32-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 datalayout = "e-p:32:32:32-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"
declare i8 @llvm.nvvm.ldu.global.i.i8.p0(ptr, i32) declare i8 @llvm.nvvm.ldu.global.i.i8.p0(ptr, i32)
define i8 @foo(ptr %a) { define i8 @foo(ptr %a) {
; Ensure we properly truncate off the high-order 24 bits ; Ensure we properly truncate off the high-order 24 bits
; CHECK: ldu.global.u8 ; CHECK: ldu.global.u8
; CHECK: cvt.u32.u16 ; CHECK: cvt.u32.u16
; CHECK: and.b32 %r{{[0-9]+}}, %r{{[0-9]+}}, 255 ; CHECK: and.b32 %r{{[0-9]+}}, %r{{[0-9]+}}, 255
%val = tail call i8 @llvm.nvvm.ldu.global.i.i8.p0(ptr %a, i32 4) %val = tail call i8 @llvm.nvvm.ldu.global.i.i8.p0(ptr %a, i32 4)
ret i8 %val ret i8 %val
} }

llvm/test/CodeGen/NVPTX/ldu-reg-plus-offset.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
target datalayout = "e-p:32:32:32-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 datalayout = "e-p:32:32:32-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"
define void @reg_plus_offset(ptr %a) { define void @reg_plus_offset(ptr %a) {
; CHECK: ldu.global.u32 %r{{[0-9]+}}, [%r{{[0-9]+}}+32]; ; CHECK: ldu.global.u32 %r{{[0-9]+}}, [%rd{{[0-9]+}}+32];
; CHECK: ldu.global.u32 %r{{[0-9]+}}, [%r{{[0-9]+}}+36]; ; CHECK: ldu.global.u32 %r{{[0-9]+}}, [%rd{{[0-9]+}}+36];
%p2 = getelementptr i32, ptr %a, i32 8 %p2 = getelementptr i32, ptr %a, i32 8
%t1 = call i32 @llvm.nvvm.ldu.global.i.i32.p0(ptr %p2, i32 4) %t1 = call i32 @llvm.nvvm.ldu.global.i.i32.p0(ptr %p2, i32 4)
%p3 = getelementptr i32, ptr %a, i32 9 %p3 = getelementptr i32, ptr %a, i32 9
%t2 = call i32 @llvm.nvvm.ldu.global.i.i32.p0(ptr %p3, i32 4) %t2 = call i32 @llvm.nvvm.ldu.global.i.i32.p0(ptr %p3, i32 4)
%t3 = mul i32 %t1, %t2 %t3 = mul i32 %t1, %t2
store i32 %t3, ptr %a store i32 %t3, ptr %a
ret void ret void
} }
declare i32 @llvm.nvvm.ldu.global.i.i32.p0(ptr, i32) declare i32 @llvm.nvvm.ldu.global.i.i32.p0(ptr, i32)
declare i32 @llvm.nvvm.read.ptx.sreg.tid.x() declare i32 @llvm.nvvm.read.ptx.sreg.tid.x()

llvm/test/CodeGen/NVPTX/libcall-fulfilled.ll

; RUN: llc < %s -march=nvptx 2>&1 | FileCheck %s ; RUN: llc < %s -march=nvptx64 2>&1 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 | %ptxas-verify %}
; Allow to make libcalls that are defined in the current module ; Allow to make libcalls that are defined in the current module
declare ptr @malloc(i64) declare ptr @malloc(i64)
declare void @free(ptr) declare void @free(ptr)
; Underlying libcall declaration ; Underlying libcall declaration
; CHECK: .visible .func (.param .align 16 .b8 func_retval0[16]) __umodti3 ; CHECK: .visible .func (.param .align 16 .b8 func_retval0[16]) __umodti3
Show All 37 Lines

llvm/test/CodeGen/NVPTX/load-sext-i1.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
target datalayout = "e-p:32:32:32-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 datalayout = "e-p:32:32:32-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 = "nvptx-nvidia-cuda" target triple = "nvptx-nvidia-cuda"
define void @main(ptr %a1, i32 %a2, ptr %arg3) { define void @main(ptr %a1, i32 %a2, ptr %arg3) {
; CHECK: ld.u8 ; CHECK: ld.u8
; CHECK-NOT: ld.u1 ; CHECK-NOT: ld.u1
%t1 = getelementptr i1, ptr %a1, i32 %a2 %t1 = getelementptr i1, ptr %a1, i32 %a2
%t2 = load i1, ptr %t1 %t2 = load i1, ptr %t1
%t3 = sext i1 %t2 to i32 %t3 = sext i1 %t2 to i32
store i32 %t3, ptr %arg3 store i32 %t3, ptr %arg3
ret void ret void
} }

llvm/test/CodeGen/NVPTX/local-stack-frame.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s --check-prefix=PTX32 ; RUN: llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s --check-prefix=PTX32
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s --check-prefix=PTX64 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s --check-prefix=PTX64
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
; Ensure we access the local stack properly ; Ensure we access the local stack properly
; PTX32: mov.u32 %SPL, __local_depot{{[0-9]+}}; ; PTX32: mov.u32 %SPL, __local_depot{{[0-9]+}};
; PTX32: cvta.local.u32 %SP, %SPL; ; PTX32: cvta.local.u32 %SP, %SPL;
; PTX32: ld.param.u32 %r{{[0-9]+}}, [foo_param_0]; ; PTX32: ld.param.u32 %r{{[0-9]+}}, [foo_param_0];
; PTX32: st.volatile.u32 [%SP+0], %r{{[0-9]+}}; ; PTX32: st.volatile.u32 [%SP+0], %r{{[0-9]+}};
▲ Show 20 Lines • Show All 72 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/lower-byval-args.ll

; RUN: llc < %s -mtriple nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s --check-prefixes=CHECK,CHECK64
; RUN: llc < %s -mtriple nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s --check-prefixes=CHECK,CHECK32 ; RUN: llc < %s -mtriple nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s --check-prefixes=CHECK,CHECK32
; RUN: llc < %s -mtriple nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s --check-prefixes=CHECK,CHECK64
; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -mtriple nvptx -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -mtriple nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -mtriple nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -mtriple nvptx -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
%struct.ham = type { [4 x i32] } %struct.ham = type { [4 x i32] }
; // Verify that load with static offset into parameter is done directly. ; // Verify that load with static offset into parameter is done directly.
; CHECK-LABEL: .visible .entry static_offset ; CHECK-LABEL: .visible .entry static_offset
; CHECK-NOT: .local ; CHECK-NOT: .local
; CHECK64: ld.param.u64 [[result_addr:%rd[0-9]+]], [{{.*}}_param_0] ; CHECK64: ld.param.u64 [[result_addr:%rd[0-9]+]], [{{.*}}_param_0]
; CHECK64: mov.b64 %[[param_addr:rd[0-9]+]], {{.*}}_param_1 ; CHECK64: mov.b64 %[[param_addr:rd[0-9]+]], {{.*}}_param_1
▲ Show 20 Lines • Show All 154 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/machine-sink.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
target datalayout = "e-p:32:32:32-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 datalayout = "e-p:32:32:32-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"
@scalar1 = internal addrspace(3) global float 0.000000e+00, align 4 @scalar1 = internal addrspace(3) global float 0.000000e+00, align 4
@scalar2 = internal addrspace(3) global float 0.000000e+00, align 4 @scalar2 = internal addrspace(3) global float 0.000000e+00, align 4
; We shouldn't sink mul.rn.f32 to BB %merge because BB %merge post-dominates ; We shouldn't sink mul.rn.f32 to BB %merge because BB %merge post-dominates
; BB %entry. Over-sinking created more register pressure on this example. The ; BB %entry. Over-sinking created more register pressure on this example. The
Show All 31 Lines

llvm/test/CodeGen/NVPTX/managed.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_30 -mattr=+ptx40 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_30 -mattr=+ptx40 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_30 -mattr=+ptx40 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_30 -mattr=+ptx40 | %ptxas-verify %}
; RUN: not --crash llc < %s -march=nvptx -mcpu=sm_20 2>&1 | FileCheck %s --check-prefix ERROR ; RUN: not --crash llc < %s -march=nvptx64 -mcpu=sm_20 2>&1 | FileCheck %s --check-prefix ERROR
; ERROR: LLVM ERROR: .attribute(.managed) requires PTX version >= 4.0 and sm_30 ; ERROR: LLVM ERROR: .attribute(.managed) requires PTX version >= 4.0 and sm_30
; CHECK: .visible .global .align 4 .u32 device_g; ; CHECK: .visible .global .align 4 .u32 device_g;
@device_g = addrspace(1) global i32 zeroinitializer @device_g = addrspace(1) global i32 zeroinitializer
; CHECK: .visible .global .attribute(.managed) .align 4 .u32 managed_g; ; CHECK: .visible .global .attribute(.managed) .align 4 .u32 managed_g;
@managed_g = addrspace(1) global i32 zeroinitializer @managed_g = addrspace(1) global i32 zeroinitializer
; CHECK: .extern .global .align 4 .u32 decl_g; ; CHECK: .extern .global .align 4 .u32 decl_g;
@decl_g = external addrspace(1) global i32, align 4 @decl_g = external addrspace(1) global i32, align 4
; CHECK: .extern .global .attribute(.managed) .align 8 .b32 managed_decl_g; ; CHECK: .extern .global .attribute(.managed) .align 8 .b64 managed_decl_g;
@managed_decl_g = external addrspace(1) global ptr, align 8 @managed_decl_g = external addrspace(1) global ptr, align 8
!nvvm.annotations = !{!0, !1} !nvvm.annotations = !{!0, !1}
!0 = !{ptr addrspace(1) @managed_g, !"managed", i32 1} !0 = !{ptr addrspace(1) @managed_g, !"managed", i32 1}
!1 = !{ptr addrspace(1) @managed_decl_g, !"managed", i32 1} !1 = !{ptr addrspace(1) @managed_decl_g, !"managed", i32 1}

llvm/test/CodeGen/NVPTX/mbarrier.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_80 | FileCheck %s -check-prefix=CHECK_PTX32 ; RUN: llc < %s -march=nvptx -mcpu=sm_80 | FileCheck %s -check-prefix=CHECK_PTX32
; RUN: llc < %s -march=nvptx64 -mcpu=sm_80 | FileCheck %s -check-prefix=CHECK_PTX64 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_80 | FileCheck %s -check-prefix=CHECK_PTX64
; RUN: %if ptxas-11.0 %{ llc < %s -march=nvptx -mcpu=sm_80 | %ptxas-verify -arch=sm_80 %} ; RUN: %if ptxas-11.0 && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_80 | %ptxas-verify -arch=sm_80 %}
; RUN: %if ptxas-11.0 %{ llc < %s -march=nvptx64 -mcpu=sm_80 | %ptxas-verify -arch=sm_80 %} ; RUN: %if ptxas-11.0 %{ llc < %s -march=nvptx64 -mcpu=sm_80 | %ptxas-verify -arch=sm_80 %}
declare void @llvm.nvvm.mbarrier.init(ptr %a, i32 %b) declare void @llvm.nvvm.mbarrier.init(ptr %a, i32 %b)
declare void @llvm.nvvm.mbarrier.init.shared(ptr addrspace(3) %a, i32 %b) declare void @llvm.nvvm.mbarrier.init.shared(ptr addrspace(3) %a, i32 %b)
; CHECK-LABEL: barrierinit ; CHECK-LABEL: barrierinit
define void @barrierinit(ptr %a, i32 %b) { define void @barrierinit(ptr %a, i32 %b) {
; CHECK_PTX32: mbarrier.init.b64 [%r{{[0-9]+}}], %r{{[0-9]+}}; ; CHECK_PTX32: mbarrier.init.b64 [%r{{[0-9]+}}], %r{{[0-9]+}};
▲ Show 20 Lines • Show All 136 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/minmax-negative.ll

; RUN: llc < %s -march=nvptx -O0 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -O0 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -O0 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -O0 | %ptxas-verify %}
define i16 @test1(ptr %sur1) { define i16 @test1(ptr %sur1) {
; CHECK-NOT: mov.u16 %rs{{[0-9]+}}, 32767 ; CHECK-NOT: mov.u16 %rs{{[0-9]+}}, 32767
%_tmp21.i = icmp sle i16 0, 0 %_tmp21.i = icmp sle i16 0, 0
%_tmp22.i = select i1 %_tmp21.i, i16 0, i16 32767 %_tmp22.i = select i1 %_tmp21.i, i16 0, i16 32767
store i16 %_tmp22.i, ptr %sur1 store i16 %_tmp22.i, ptr %sur1
ret i16 0 ret i16 0
} }

llvm/test/CodeGen/NVPTX/module-inline-asm.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
target datalayout = "e-p:32:32:32-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 datalayout = "e-p:32:32:32-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"
; module asm must come after PTX version/target directives. ; module asm must come after PTX version/target directives.
; CHECK-NOT: .global .b32 val; ; CHECK-NOT: .global .b32 val;
; CHECK-DAG: .version ; CHECK-DAG: .version
; CHECK-DAG: .target ; CHECK-DAG: .target
Show All 11 Lines

llvm/test/CodeGen/NVPTX/mulwide.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 -O3 | FileCheck %s --check-prefix=OPT ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -O3 | FileCheck %s --check-prefix=OPT
; RUN: llc < %s -march=nvptx -mcpu=sm_20 -O0 | FileCheck %s --check-prefix=NOOPT ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -O0 | FileCheck %s --check-prefix=NOOPT
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -O3 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -O3 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -O0 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -O0 | %ptxas-verify %}
; OPT-LABEL: @mulwide16 ; OPT-LABEL: @mulwide16
; NOOPT-LABEL: @mulwide16 ; NOOPT-LABEL: @mulwide16
define i32 @mulwide16(i16 %a, i16 %b) { define i32 @mulwide16(i16 %a, i16 %b) {
; OPT: mul.wide.s16 ; OPT: mul.wide.s16
; NOOPT: mul.lo.s32 ; NOOPT: mul.lo.s32
%val0 = sext i16 %a to i32 %val0 = sext i16 %a to i32
%val1 = sext i16 %b to i32 %val1 = sext i16 %b to i32
▲ Show 20 Lines • Show All 100 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/named-barriers.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; Use bar.sync to arrive at a pre-computed barrier number and ; Use bar.sync to arrive at a pre-computed barrier number and
; wait for all threads in CTA to also arrive: ; wait for all threads in CTA to also arrive:
define ptx_device void @test_barrier_named_cta() { define ptx_device void @test_barrier_named_cta() {
; CHECK: mov.u32 %r[[REG0:[0-9]+]], 0; ; CHECK: mov.u32 %r[[REG0:[0-9]+]], 0;
; CHECK: bar.sync %r[[REG0]]; ; CHECK: bar.sync %r[[REG0]];
; CHECK: mov.u32 %r[[REG1:[0-9]+]], 10; ; CHECK: mov.u32 %r[[REG1:[0-9]+]], 10;
Show All 31 Lines

llvm/test/CodeGen/NVPTX/nofunc.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; Test that we don't crash if we're compiling a module with function references, ; Test that we don't crash if we're compiling a module with function references,
; but without any functions in it. ; but without any functions in it.
target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64" target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64"
target triple = "nvptx64-nvidia-cuda" target triple = "nvptx64-nvidia-cuda"
@Funcs = local_unnamed_addr addrspace(1) externally_initialized @Funcs = local_unnamed_addr addrspace(1) externally_initialized
global [1 x ptr] [ptr @func], align 8 global [1 x ptr] [ptr @func], align 8
declare void @func(ptr) declare void @func(ptr)
; CHECK: Funcs[1] = {func} ; CHECK: Funcs[1] = {func}

llvm/test/CodeGen/NVPTX/nvcl-param-align.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
target triple = "nvptx-unknown-nvcl" target triple = "nvptx-unknown-nvcl"
define void @foo(i64 %img, i64 %sampler, ptr align 32 %v1, ptr %v2) { define void @foo(i64 %img, i64 %sampler, ptr align 32 %v1, ptr %v2) {
; The parameter alignment is determined by the align attribute (default 1). ; The parameter alignment is determined by the align attribute (default 1).
; CHECK-LABEL: .entry foo( ; CHECK-LABEL: .entry foo(
; CHECK: .param .u32 .ptr .align 32 foo_param_2 ; CHECK: .param .u64 .ptr .align 32 foo_param_2
; CHECK: .param .u32 .ptr .align 1 foo_param_3 ; CHECK: .param .u64 .ptr .align 1 foo_param_3
ret void ret void
} }
!nvvm.annotations = !{!1, !2, !3} !nvvm.annotations = !{!1, !2, !3}
!1 = !{ptr @foo, !"kernel", i32 1} !1 = !{ptr @foo, !"kernel", i32 1}
!2 = !{ptr @foo, !"rdoimage", i32 0} !2 = !{ptr @foo, !"rdoimage", i32 0}
!3 = !{ptr @foo, !"sampler", i32 1} !3 = !{ptr @foo, !"sampler", i32 1}

llvm/test/CodeGen/NVPTX/nvvm-annotations-D120129.ll

; RUN: llc < %s -mtriple=nvptx-unknown-unknown | FileCheck %s ; RUN: llc < %s -mtriple=nvptx64-unknown-unknown | FileCheck %s
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx-unknown-unknown | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64-unknown-unknown | %ptxas-verify %}
; ;
; NVPTXTargetLowering::getFunctionParamOptimizedAlign, which was introduces in ; NVPTXTargetLowering::getFunctionParamOptimizedAlign, which was introduces in
; D120129, contained a poorly designed assertion checking that a function with ; D120129, contained a poorly designed assertion checking that a function with
; internal or private linkage is not a kernel. It relied on invariants that ; internal or private linkage is not a kernel. It relied on invariants that
; were not actually guaranteed, and that resulted in compiler crash with some ; were not actually guaranteed, and that resulted in compiler crash with some
; CUDA versions (see discussion with @jdoerfert in D120129). This test contains ; CUDA versions (see discussion with @jdoerfert in D120129). This test contains
; metadata that caused compiler crash and a function with internal linkage ; metadata that caused compiler crash and a function with internal linkage
; which purpose is to let compiler run on path where the crash happened. ; which purpose is to let compiler run on path where the crash happened.
Show All 25 Lines

llvm/test/CodeGen/NVPTX/packed-aggr.ll

; RUN: not --crash llc < %s -march=nvptx -mcpu=sm_20 -mattr=+ptx70 2>&1 | \ ; RUN: not --crash llc < %s -march=nvptx -mcpu=sm_20 -mattr=+ptx70 2>&1 | \
; RUN: FileCheck %s --check-prefix=ERROR ; RUN: FileCheck %s --check-prefix=ERROR
; RUN: llc < %s -march=nvptx -mcpu=sm_20 -mattr=+ptx71 | \ ; RUN: llc < %s -march=nvptx -mcpu=sm_20 -mattr=+ptx71 | \
; RUN: FileCheck %s --check-prefixes=CHECK,CHECK32 ; RUN: FileCheck %s --check-prefixes=CHECK,CHECK32
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -mattr=+ptx71 | \ ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -mattr=+ptx71 | \
; RUN: FileCheck %s --check-prefixes=CHECK,CHECK64 ; RUN: FileCheck %s --check-prefixes=CHECK,CHECK64
; RUN: %if ptxas-11.1 %{ llc < %s -march=nvptx -mcpu=sm_20 -mattr=+ptx71 | %ptxas-verify %} ; RUN: %if ptxas-11.1 && !ptxas-12.0%{ llc < %s -march=nvptx -mcpu=sm_20 -mattr=+ptx71 | %ptxas-verify %}
; RUN: %if ptxas-11.1 %{ llc < %s -march=nvptx64 -mcpu=sm_20 -mattr=+ptx71 | %ptxas-verify %} ; RUN: %if ptxas-11.1 %{ llc < %s -march=nvptx64 -mcpu=sm_20 -mattr=+ptx71 | %ptxas-verify %}
;; Test that packed structs with symbol references are represented using the ;; Test that packed structs with symbol references are represented using the
;; mask() operator. ;; mask() operator.
declare void @func() declare void @func()
@p = addrspace(1) global i8 0 @p = addrspace(1) global i8 0
; CHECK: .extern .func func ; CHECK: .extern .func func
▲ Show 20 Lines • Show All 79 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/param-align.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefixes=CHECK,NOALIGN4 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefixes=CHECK,NOALIGN4
; RUN: llc < %s -march=nvptx -mcpu=sm_20 -nvptx-force-min-byval-param-align | FileCheck %s --check-prefixes=CHECK,ALIGN4 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -nvptx-force-min-byval-param-align | FileCheck %s --check-prefixes=CHECK,ALIGN4
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -nvptx-force-min-byval-param-align | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -nvptx-force-min-byval-param-align | %ptxas-verify %}
;;; Need 4-byte alignment on ptr passed byval ;;; Need 4-byte alignment on ptr passed byval
define ptx_device void @t1(ptr byval(float) %x) { define ptx_device void @t1(ptr byval(float) %x) {
; CHECK: .func t1 ; CHECK: .func t1
; CHECK: .param .align 4 .b8 t1_param_0[4] ; CHECK: .param .align 4 .b8 t1_param_0[4]
ret void ret void
} }
▲ Show 20 Lines • Show All 59 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/param-vectorize-device.ll

; RUN: llc < %s -mtriple=nvptx-unknown-unknown | FileCheck %s ; RUN: llc < %s -mtriple=nvptx64-unknown-unknown | FileCheck %s
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx-unknown-unknown | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64-unknown-unknown | %ptxas-verify %}
; ;
; Check that parameters of a __device__ function with private or internal ; Check that parameters of a __device__ function with private or internal
; linkage called from a __global__ (kernel) function get increased alignment, ; linkage called from a __global__ (kernel) function get increased alignment,
; and additional vectorization is performed on loads/stores with that ; and additional vectorization is performed on loads/stores with that
; parameters. ; parameters.
; ;
; Test IR is a minimized version of IR generated with the following command ; Test IR is a minimized version of IR generated with the following command
; from the source code below: ; from the source code below:
▲ Show 20 Lines • Show All 65 Lines • ▼ Show 20 Lines
; Section 1 - checking that: ; Section 1 - checking that:
; - function argument (including retval) vectorization is done with internal linkage; ; - function argument (including retval) vectorization is done with internal linkage;
; - caller and callee specify correct alignment for callee's params. ; - caller and callee specify correct alignment for callee's params.
define dso_local void @caller_St4x1(ptr nocapture noundef readonly byval(%struct.St4x1) align 4 %in, ptr nocapture noundef writeonly %ret) { define dso_local void @caller_St4x1(ptr nocapture noundef readonly byval(%struct.St4x1) align 4 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func caller_St4x1( ; CHECK-LABEL: .visible .func caller_St4x1(
; CHECK: .param .align 4 .b8 caller_St4x1_param_0[4], ; CHECK: .param .align 4 .b8 caller_St4x1_param_0[4],
; CHECK: .param .b32 caller_St4x1_param_1 ; CHECK: .param .b64 caller_St4x1_param_1
; CHECK: ) ; CHECK: )
; CHECK: .param .b32 param0; ; CHECK: .param .b32 param0;
; CHECK: st.param.b32 [param0+0], {{%r[0-9]+}}; ; CHECK: st.param.b32 [param0+0], {{%r[0-9]+}};
; CHECK: .param .align 16 .b8 retval0[4]; ; CHECK: .param .align 16 .b8 retval0[4];
; CHECK: call.uni (retval0), ; CHECK: call.uni (retval0),
; CHECK-NEXT: callee_St4x1, ; CHECK-NEXT: callee_St4x1,
; CHECK-NEXT: ( ; CHECK-NEXT: (
; CHECK-NEXT: param0 ; CHECK-NEXT: param0
Show All 15 Linesdefine internal fastcc [1 x i32] @callee_St4x1(i32 %in.0.val) {
; CHECK-NEXT: ret; ; CHECK-NEXT: ret;
%oldret = insertvalue [1 x i32] poison, i32 %in.0.val, 0 %oldret = insertvalue [1 x i32] poison, i32 %in.0.val, 0
ret [1 x i32] %oldret ret [1 x i32] %oldret
} }
define dso_local void @caller_St4x2(ptr nocapture noundef readonly byval(%struct.St4x2) align 4 %in, ptr nocapture noundef writeonly %ret) { define dso_local void @caller_St4x2(ptr nocapture noundef readonly byval(%struct.St4x2) align 4 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func caller_St4x2( ; CHECK-LABEL: .visible .func caller_St4x2(
; CHECK: .param .align 4 .b8 caller_St4x2_param_0[8], ; CHECK: .param .align 4 .b8 caller_St4x2_param_0[8],
; CHECK: .param .b32 caller_St4x2_param_1 ; CHECK: .param .b64 caller_St4x2_param_1
; CHECK: ) ; CHECK: )
; CHECK: .param .align 16 .b8 param0[8]; ; CHECK: .param .align 16 .b8 param0[8];
; CHECK: st.param.v2.b32 [param0+0], {{{%r[0-9]+}}, {{%r[0-9]+}}}; ; CHECK: st.param.v2.b32 [param0+0], {{{%r[0-9]+}}, {{%r[0-9]+}}};
; CHECK: .param .align 16 .b8 retval0[8]; ; CHECK: .param .align 16 .b8 retval0[8];
; CHECK: call.uni (retval0), ; CHECK: call.uni (retval0),
; CHECK-NEXT: callee_St4x2, ; CHECK-NEXT: callee_St4x2,
; CHECK-NEXT: ( ; CHECK-NEXT: (
; CHECK-NEXT: param0 ; CHECK-NEXT: param0
Show All 24 Linesdefine internal fastcc [2 x i32] @callee_St4x2(ptr nocapture noundef readonly byval(%struct.St4x2) align 4 %in) {
%3 = insertvalue [2 x i32] poison, i32 %1, 0 %3 = insertvalue [2 x i32] poison, i32 %1, 0
%oldret = insertvalue [2 x i32] %3, i32 %2, 1 %oldret = insertvalue [2 x i32] %3, i32 %2, 1
ret [2 x i32] %oldret ret [2 x i32] %oldret
} }
define dso_local void @caller_St4x3(ptr nocapture noundef readonly byval(%struct.St4x3) align 4 %in, ptr nocapture noundef writeonly %ret) { define dso_local void @caller_St4x3(ptr nocapture noundef readonly byval(%struct.St4x3) align 4 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func caller_St4x3( ; CHECK-LABEL: .visible .func caller_St4x3(
; CHECK: .param .align 4 .b8 caller_St4x3_param_0[12], ; CHECK: .param .align 4 .b8 caller_St4x3_param_0[12],
; CHECK: .param .b32 caller_St4x3_param_1 ; CHECK: .param .b64 caller_St4x3_param_1
; CHECK: ) ; CHECK: )
; CHECK: .param .align 16 .b8 param0[12]; ; CHECK: .param .align 16 .b8 param0[12];
; CHECK: st.param.v2.b32 [param0+0], {{{%r[0-9]+}}, {{%r[0-9]+}}}; ; CHECK: st.param.v2.b32 [param0+0], {{{%r[0-9]+}}, {{%r[0-9]+}}};
; CHECK: st.param.b32 [param0+8], {{%r[0-9]+}}; ; CHECK: st.param.b32 [param0+8], {{%r[0-9]+}};
; CHECK: .param .align 16 .b8 retval0[12]; ; CHECK: .param .align 16 .b8 retval0[12];
; CHECK: call.uni (retval0), ; CHECK: call.uni (retval0),
; CHECK-NEXT: callee_St4x3, ; CHECK-NEXT: callee_St4x3,
; CHECK-NEXT: ( ; CHECK-NEXT: (
Show All 31 Linesdefine internal fastcc [3 x i32] @callee_St4x3(ptr nocapture noundef readonly byval(%struct.St4x3) align 4 %in) {
%5 = insertvalue [3 x i32] %4, i32 %2, 1 %5 = insertvalue [3 x i32] %4, i32 %2, 1
%oldret = insertvalue [3 x i32] %5, i32 %3, 2 %oldret = insertvalue [3 x i32] %5, i32 %3, 2
ret [3 x i32] %oldret ret [3 x i32] %oldret
} }
define dso_local void @caller_St4x4(ptr nocapture noundef readonly byval(%struct.St4x4) align 4 %in, ptr nocapture noundef writeonly %ret) { define dso_local void @caller_St4x4(ptr nocapture noundef readonly byval(%struct.St4x4) align 4 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func caller_St4x4( ; CHECK-LABEL: .visible .func caller_St4x4(
; CHECK: .param .align 4 .b8 caller_St4x4_param_0[16], ; CHECK: .param .align 4 .b8 caller_St4x4_param_0[16],
; CHECK: .param .b32 caller_St4x4_param_1 ; CHECK: .param .b64 caller_St4x4_param_1
; CHECK: ) ; CHECK: )
; CHECK: .param .align 16 .b8 param0[16]; ; CHECK: .param .align 16 .b8 param0[16];
; CHECK: st.param.v4.b32 [param0+0], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}; ; CHECK: st.param.v4.b32 [param0+0], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}};
; CHECK: .param .align 16 .b8 retval0[16]; ; CHECK: .param .align 16 .b8 retval0[16];
; CHECK: call.uni (retval0), ; CHECK: call.uni (retval0),
; CHECK-NEXT: callee_St4x4, ; CHECK-NEXT: callee_St4x4,
; CHECK-NEXT: ( ; CHECK-NEXT: (
; CHECK-NEXT: param0 ; CHECK-NEXT: param0
Show All 33 Linesdefine internal fastcc [4 x i32] @callee_St4x4(ptr nocapture noundef readonly byval(%struct.St4x4) align 4 %in) {
%7 = insertvalue [4 x i32] %6, i32 %3, 2 %7 = insertvalue [4 x i32] %6, i32 %3, 2
%oldret = insertvalue [4 x i32] %7, i32 %4, 3 %oldret = insertvalue [4 x i32] %7, i32 %4, 3
ret [4 x i32] %oldret ret [4 x i32] %oldret
} }
define dso_local void @caller_St4x5(ptr nocapture noundef readonly byval(%struct.St4x5) align 4 %in, ptr nocapture noundef writeonly %ret) { define dso_local void @caller_St4x5(ptr nocapture noundef readonly byval(%struct.St4x5) align 4 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func caller_St4x5( ; CHECK-LABEL: .visible .func caller_St4x5(
; CHECK: .param .align 4 .b8 caller_St4x5_param_0[20], ; CHECK: .param .align 4 .b8 caller_St4x5_param_0[20],
; CHECK: .param .b32 caller_St4x5_param_1 ; CHECK: .param .b64 caller_St4x5_param_1
; CHECK: ) ; CHECK: )
; CHECK: .param .align 16 .b8 param0[20]; ; CHECK: .param .align 16 .b8 param0[20];
; CHECK: st.param.v4.b32 [param0+0], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}; ; CHECK: st.param.v4.b32 [param0+0], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}};
; CHECK: st.param.b32 [param0+16], {{%r[0-9]+}}; ; CHECK: st.param.b32 [param0+16], {{%r[0-9]+}};
; CHECK: .param .align 16 .b8 retval0[20]; ; CHECK: .param .align 16 .b8 retval0[20];
; CHECK: call.uni (retval0), ; CHECK: call.uni (retval0),
; CHECK-NEXT: callee_St4x5, ; CHECK-NEXT: callee_St4x5,
; CHECK-NEXT: ( ; CHECK-NEXT: (
▲ Show 20 Lines • Show All 43 Lines • ▼ Show 20 Linesdefine internal fastcc [5 x i32] @callee_St4x5(ptr nocapture noundef readonly byval(%struct.St4x5) align 4 %in) {
%9 = insertvalue [5 x i32] %8, i32 %4, 3 %9 = insertvalue [5 x i32] %8, i32 %4, 3
%oldret = insertvalue [5 x i32] %9, i32 %5, 4 %oldret = insertvalue [5 x i32] %9, i32 %5, 4
ret [5 x i32] %oldret ret [5 x i32] %oldret
} }
define dso_local void @caller_St4x6(ptr nocapture noundef readonly byval(%struct.St4x6) align 4 %in, ptr nocapture noundef writeonly %ret) { define dso_local void @caller_St4x6(ptr nocapture noundef readonly byval(%struct.St4x6) align 4 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func caller_St4x6( ; CHECK-LABEL: .visible .func caller_St4x6(
; CHECK: .param .align 4 .b8 caller_St4x6_param_0[24], ; CHECK: .param .align 4 .b8 caller_St4x6_param_0[24],
; CHECK: .param .b32 caller_St4x6_param_1 ; CHECK: .param .b64 caller_St4x6_param_1
; CHECK: ) ; CHECK: )
; CHECK: .param .align 16 .b8 param0[24]; ; CHECK: .param .align 16 .b8 param0[24];
; CHECK: st.param.v4.b32 [param0+0], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}; ; CHECK: st.param.v4.b32 [param0+0], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}};
; CHECK: st.param.v2.b32 [param0+16], {{{%r[0-9]+}}, {{%r[0-9]+}}}; ; CHECK: st.param.v2.b32 [param0+16], {{{%r[0-9]+}}, {{%r[0-9]+}}};
; CHECK: .param .align 16 .b8 retval0[24]; ; CHECK: .param .align 16 .b8 retval0[24];
; CHECK: call.uni (retval0), ; CHECK: call.uni (retval0),
; CHECK-NEXT: callee_St4x6, ; CHECK-NEXT: callee_St4x6,
; CHECK-NEXT: ( ; CHECK-NEXT: (
▲ Show 20 Lines • Show All 49 Lines • ▼ Show 20 Linesdefine internal fastcc [6 x i32] @callee_St4x6(ptr nocapture noundef readonly byval(%struct.St4x6) align 4 %in) {
%11 = insertvalue [6 x i32] %10, i32 %5, 4 %11 = insertvalue [6 x i32] %10, i32 %5, 4
%oldret = insertvalue [6 x i32] %11, i32 %6, 5 %oldret = insertvalue [6 x i32] %11, i32 %6, 5
ret [6 x i32] %oldret ret [6 x i32] %oldret
} }
define dso_local void @caller_St4x7(ptr nocapture noundef readonly byval(%struct.St4x7) align 4 %in, ptr nocapture noundef writeonly %ret) { define dso_local void @caller_St4x7(ptr nocapture noundef readonly byval(%struct.St4x7) align 4 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func caller_St4x7( ; CHECK-LABEL: .visible .func caller_St4x7(
; CHECK: .param .align 4 .b8 caller_St4x7_param_0[28], ; CHECK: .param .align 4 .b8 caller_St4x7_param_0[28],
; CHECK: .param .b32 caller_St4x7_param_1 ; CHECK: .param .b64 caller_St4x7_param_1
; CHECK: ) ; CHECK: )
; CHECK: .param .align 16 .b8 param0[28]; ; CHECK: .param .align 16 .b8 param0[28];
; CHECK: st.param.v4.b32 [param0+0], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}; ; CHECK: st.param.v4.b32 [param0+0], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}};
; CHECK: st.param.v2.b32 [param0+16], {{{%r[0-9]+}}, {{%r[0-9]+}}}; ; CHECK: st.param.v2.b32 [param0+16], {{{%r[0-9]+}}, {{%r[0-9]+}}};
; CHECK: st.param.b32 [param0+24], {{%r[0-9]+}}; ; CHECK: st.param.b32 [param0+24], {{%r[0-9]+}};
; CHECK: .param .align 16 .b8 retval0[28]; ; CHECK: .param .align 16 .b8 retval0[28];
; CHECK: call.uni (retval0), ; CHECK: call.uni (retval0),
; CHECK-NEXT: callee_St4x7, ; CHECK-NEXT: callee_St4x7,
▲ Show 20 Lines • Show All 59 Lines • ▼ Show 20 Linesdefine internal fastcc [7 x i32] @callee_St4x7(ptr nocapture noundef readonly byval(%struct.St4x7) align 4 %in) {
%13 = insertvalue [7 x i32] %12, i32 %6, 5 %13 = insertvalue [7 x i32] %12, i32 %6, 5
%oldret = insertvalue [7 x i32] %13, i32 %7, 6 %oldret = insertvalue [7 x i32] %13, i32 %7, 6
ret [7 x i32] %oldret ret [7 x i32] %oldret
} }
define dso_local void @caller_St4x8(ptr nocapture noundef readonly byval(%struct.St4x8) align 4 %in, ptr nocapture noundef writeonly %ret) { define dso_local void @caller_St4x8(ptr nocapture noundef readonly byval(%struct.St4x8) align 4 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func caller_St4x8( ; CHECK-LABEL: .visible .func caller_St4x8(
; CHECK: .param .align 4 .b8 caller_St4x8_param_0[32], ; CHECK: .param .align 4 .b8 caller_St4x8_param_0[32],
; CHECK: .param .b32 caller_St4x8_param_1 ; CHECK: .param .b64 caller_St4x8_param_1
; CHECK: ) ; CHECK: )
; CHECK: .param .align 16 .b8 param0[32]; ; CHECK: .param .align 16 .b8 param0[32];
; CHECK: st.param.v4.b32 [param0+0], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}; ; CHECK: st.param.v4.b32 [param0+0], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}};
; CHECK: st.param.v4.b32 [param0+16], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}}; ; CHECK: st.param.v4.b32 [param0+16], {{{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}, {{%r[0-9]+}}};
; CHECK: .param .align 16 .b8 retval0[32]; ; CHECK: .param .align 16 .b8 retval0[32];
; CHECK: call.uni (retval0), ; CHECK: call.uni (retval0),
; CHECK-NEXT: callee_St4x8, ; CHECK-NEXT: callee_St4x8,
; CHECK-NEXT: ( ; CHECK-NEXT: (
▲ Show 20 Lines • Show All 61 Lines • ▼ Show 20 Linesdefine internal fastcc [8 x i32] @callee_St4x8(ptr nocapture noundef readonly byval(%struct.St4x8) align 4 %in) {
%15 = insertvalue [8 x i32] %14, i32 %7, 6 %15 = insertvalue [8 x i32] %14, i32 %7, 6
%oldret = insertvalue [8 x i32] %15, i32 %8, 7 %oldret = insertvalue [8 x i32] %15, i32 %8, 7
ret [8 x i32] %oldret ret [8 x i32] %oldret
} }
define dso_local void @caller_St8x1(ptr nocapture noundef readonly byval(%struct.St8x1) align 8 %in, ptr nocapture noundef writeonly %ret) { define dso_local void @caller_St8x1(ptr nocapture noundef readonly byval(%struct.St8x1) align 8 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func caller_St8x1( ; CHECK-LABEL: .visible .func caller_St8x1(
; CHECK: .param .align 8 .b8 caller_St8x1_param_0[8], ; CHECK: .param .align 8 .b8 caller_St8x1_param_0[8],
; CHECK: .param .b32 caller_St8x1_param_1 ; CHECK: .param .b64 caller_St8x1_param_1
; CHECK: ) ; CHECK: )
; CHECK: .param .b64 param0; ; CHECK: .param .b64 param0;
; CHECK: st.param.b64 [param0+0], {{%rd[0-9]+}}; ; CHECK: st.param.b64 [param0+0], {{%rd[0-9]+}};
; CHECK: .param .align 16 .b8 retval0[8]; ; CHECK: .param .align 16 .b8 retval0[8];
; CHECK: call.uni (retval0), ; CHECK: call.uni (retval0),
; CHECK-NEXT: callee_St8x1, ; CHECK-NEXT: callee_St8x1,
; CHECK-NEXT: ( ; CHECK-NEXT: (
; CHECK-NEXT: param0 ; CHECK-NEXT: param0
Show All 15 Linesdefine internal fastcc [1 x i64] @callee_St8x1(i64 %in.0.val) {
; CHECK-NEXT: ret; ; CHECK-NEXT: ret;
%oldret = insertvalue [1 x i64] poison, i64 %in.0.val, 0 %oldret = insertvalue [1 x i64] poison, i64 %in.0.val, 0
ret [1 x i64] %oldret ret [1 x i64] %oldret
} }
define dso_local void @caller_St8x2(ptr nocapture noundef readonly byval(%struct.St8x2) align 8 %in, ptr nocapture noundef writeonly %ret) { define dso_local void @caller_St8x2(ptr nocapture noundef readonly byval(%struct.St8x2) align 8 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func caller_St8x2( ; CHECK-LABEL: .visible .func caller_St8x2(
; CHECK: .param .align 8 .b8 caller_St8x2_param_0[16], ; CHECK: .param .align 8 .b8 caller_St8x2_param_0[16],
; CHECK: .param .b32 caller_St8x2_param_1 ; CHECK: .param .b64 caller_St8x2_param_1
; CHECK: ) ; CHECK: )
; CHECK: .param .align 16 .b8 param0[16]; ; CHECK: .param .align 16 .b8 param0[16];
; CHECK: st.param.v2.b64 [param0+0], {{{%rd[0-9]+}}, {{%rd[0-9]+}}}; ; CHECK: st.param.v2.b64 [param0+0], {{{%rd[0-9]+}}, {{%rd[0-9]+}}};
; CHECK: .param .align 16 .b8 retval0[16]; ; CHECK: .param .align 16 .b8 retval0[16];
; CHECK: call.uni (retval0), ; CHECK: call.uni (retval0),
; CHECK-NEXT: callee_St8x2, ; CHECK-NEXT: callee_St8x2,
; CHECK-NEXT: ( ; CHECK-NEXT: (
; CHECK-NEXT: param0 ; CHECK-NEXT: param0
Show All 21 Linesdefine internal fastcc [2 x i64] @callee_St8x2(ptr nocapture noundef readonly byval(%struct.St8x2) align 8 %in) {
%3 = insertvalue [2 x i64] poison, i64 %1, 0 %3 = insertvalue [2 x i64] poison, i64 %1, 0
%oldret = insertvalue [2 x i64] %3, i64 %2, 1 %oldret = insertvalue [2 x i64] %3, i64 %2, 1
ret [2 x i64] %oldret ret [2 x i64] %oldret
} }
define dso_local void @caller_St8x3(ptr nocapture noundef readonly byval(%struct.St8x3) align 8 %in, ptr nocapture noundef writeonly %ret) { define dso_local void @caller_St8x3(ptr nocapture noundef readonly byval(%struct.St8x3) align 8 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func caller_St8x3( ; CHECK-LABEL: .visible .func caller_St8x3(
; CHECK: .param .align 8 .b8 caller_St8x3_param_0[24], ; CHECK: .param .align 8 .b8 caller_St8x3_param_0[24],
; CHECK: .param .b32 caller_St8x3_param_1 ; CHECK: .param .b64 caller_St8x3_param_1
; CHECK: ) ; CHECK: )
; CHECK: .param .align 16 .b8 param0[24]; ; CHECK: .param .align 16 .b8 param0[24];
; CHECK: st.param.v2.b64 [param0+0], {{{%rd[0-9]+}}, {{%rd[0-9]+}}}; ; CHECK: st.param.v2.b64 [param0+0], {{{%rd[0-9]+}}, {{%rd[0-9]+}}};
; CHECK: st.param.b64 [param0+16], {{%rd[0-9]+}}; ; CHECK: st.param.b64 [param0+16], {{%rd[0-9]+}};
; CHECK: .param .align 16 .b8 retval0[24]; ; CHECK: .param .align 16 .b8 retval0[24];
; CHECK: call.uni (retval0), ; CHECK: call.uni (retval0),
; CHECK-NEXT: callee_St8x3, ; CHECK-NEXT: callee_St8x3,
; CHECK-NEXT: ( ; CHECK-NEXT: (
Show All 31 Linesdefine internal fastcc [3 x i64] @callee_St8x3(ptr nocapture noundef readonly byval(%struct.St8x3) align 8 %in) {
%5 = insertvalue [3 x i64] %4, i64 %2, 1 %5 = insertvalue [3 x i64] %4, i64 %2, 1
%oldret = insertvalue [3 x i64] %5, i64 %3, 2 %oldret = insertvalue [3 x i64] %5, i64 %3, 2
ret [3 x i64] %oldret ret [3 x i64] %oldret
} }
define dso_local void @caller_St8x4(ptr nocapture noundef readonly byval(%struct.St8x4) align 8 %in, ptr nocapture noundef writeonly %ret) { define dso_local void @caller_St8x4(ptr nocapture noundef readonly byval(%struct.St8x4) align 8 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func caller_St8x4( ; CHECK-LABEL: .visible .func caller_St8x4(
; CHECK: .param .align 8 .b8 caller_St8x4_param_0[32], ; CHECK: .param .align 8 .b8 caller_St8x4_param_0[32],
; CHECK: .param .b32 caller_St8x4_param_1 ; CHECK: .param .b64 caller_St8x4_param_1
; CHECK: ) ; CHECK: )
; CHECK: .param .align 16 .b8 param0[32]; ; CHECK: .param .align 16 .b8 param0[32];
; CHECK: st.param.v2.b64 [param0+0], {{{%rd[0-9]+}}, {{%rd[0-9]+}}}; ; CHECK: st.param.v2.b64 [param0+0], {{{%rd[0-9]+}}, {{%rd[0-9]+}}};
; CHECK: st.param.v2.b64 [param0+16], {{{%rd[0-9]+}}, {{%rd[0-9]+}}}; ; CHECK: st.param.v2.b64 [param0+16], {{{%rd[0-9]+}}, {{%rd[0-9]+}}};
; CHECK: .param .align 16 .b8 retval0[32]; ; CHECK: .param .align 16 .b8 retval0[32];
; CHECK: call.uni (retval0), ; CHECK: call.uni (retval0),
; CHECK-NEXT: callee_St8x4, ; CHECK-NEXT: callee_St8x4,
; CHECK-NEXT: ( ; CHECK-NEXT: (
▲ Show 20 Lines • Show All 95 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/param-vectorize-kernel.ll

; RUN: llc < %s -mtriple=nvptx-unknown-unknown | FileCheck %s ; RUN: llc < %s -mtriple=nvptx64-unknown-unknown | FileCheck %s
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx-unknown-unknown | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64-unknown-unknown | %ptxas-verify %}
; ;
; Check that parameters of a __global__ (kernel) function do not get increased ; Check that parameters of a __global__ (kernel) function do not get increased
; alignment, and no additional vectorization is performed on loads/stores with ; alignment, and no additional vectorization is performed on loads/stores with
; that parameters. ; that parameters.
; ;
; Test IR is a minimized version of IR generated with the following command ; Test IR is a minimized version of IR generated with the following command
; from the source code below: ; from the source code below:
; $ clang++ -O3 --cuda-gpu-arch=sm_35 -S -emit-llvm src.cu ; $ clang++ -O3 --cuda-gpu-arch=sm_35 -S -emit-llvm src.cu
▲ Show 20 Lines • Show All 47 Lines • ▼ Show 20 Lines
%struct.St8x1 = type { [1 x i64] } %struct.St8x1 = type { [1 x i64] }
%struct.St8x2 = type { [2 x i64] } %struct.St8x2 = type { [2 x i64] }
%struct.St8x3 = type { [3 x i64] } %struct.St8x3 = type { [3 x i64] }
%struct.St8x4 = type { [4 x i64] } %struct.St8x4 = type { [4 x i64] }
define dso_local void @foo_St4x1(ptr nocapture noundef readonly byval(%struct.St4x1) align 4 %in, ptr nocapture noundef writeonly %ret) { define dso_local void @foo_St4x1(ptr nocapture noundef readonly byval(%struct.St4x1) align 4 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func foo_St4x1( ; CHECK-LABEL: .visible .func foo_St4x1(
; CHECK: .param .align 4 .b8 foo_St4x1_param_0[4], ; CHECK: .param .align 4 .b8 foo_St4x1_param_0[4],
; CHECK: .param .b32 foo_St4x1_param_1 ; CHECK: .param .b64 foo_St4x1_param_1
; CHECK: ) ; CHECK: )
; CHECK: ld.param.u32 [[R1:%r[0-9]+]], [foo_St4x1_param_1]; ; CHECK: ld.param.u64 [[R1:%rd[0-9]+]], [foo_St4x1_param_1];
; CHECK: ld.param.u32 [[R2:%r[0-9]+]], [foo_St4x1_param_0]; ; CHECK: ld.param.u32 [[R2:%r[0-9]+]], [foo_St4x1_param_0];
; CHECK: st.u32 [[[R1]]], [[R2]]; ; CHECK: st.u32 [[[R1]]], [[R2]];
; CHECK: ret; ; CHECK: ret;
%1 = load i32, ptr %in, align 4 %1 = load i32, ptr %in, align 4
store i32 %1, ptr %ret, align 4 store i32 %1, ptr %ret, align 4
ret void ret void
} }
define dso_local void @foo_St4x2(ptr nocapture noundef readonly byval(%struct.St4x2) align 4 %in, ptr nocapture noundef writeonly %ret) { define dso_local void @foo_St4x2(ptr nocapture noundef readonly byval(%struct.St4x2) align 4 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func foo_St4x2( ; CHECK-LABEL: .visible .func foo_St4x2(
; CHECK: .param .align 4 .b8 foo_St4x2_param_0[8], ; CHECK: .param .align 4 .b8 foo_St4x2_param_0[8],
; CHECK: .param .b32 foo_St4x2_param_1 ; CHECK: .param .b64 foo_St4x2_param_1
; CHECK: ) ; CHECK: )
; CHECK: ld.param.u32 [[R1:%r[0-9]+]], [foo_St4x2_param_1]; ; CHECK: ld.param.u64 [[R1:%rd[0-9]+]], [foo_St4x2_param_1];
; CHECK: ld.param.u32 [[R2:%r[0-9]+]], [foo_St4x2_param_0]; ; CHECK: ld.param.u32 [[R2:%r[0-9]+]], [foo_St4x2_param_0];
; CHECK: st.u32 [[[R1]]], [[R2]]; ; CHECK: st.u32 [[[R1]]], [[R2]];
; CHECK: ld.param.u32 [[R3:%r[0-9]+]], [foo_St4x2_param_0+4]; ; CHECK: ld.param.u32 [[R3:%r[0-9]+]], [foo_St4x2_param_0+4];
; CHECK: st.u32 [[[R1]]+4], [[R3]]; ; CHECK: st.u32 [[[R1]]+4], [[R3]];
; CHECK: ret; ; CHECK: ret;
%1 = load i32, ptr %in, align 4 %1 = load i32, ptr %in, align 4
store i32 %1, ptr %ret, align 4 store i32 %1, ptr %ret, align 4
%arrayidx.1 = getelementptr inbounds [2 x i32], ptr %in, i64 0, i64 1 %arrayidx.1 = getelementptr inbounds [2 x i32], ptr %in, i64 0, i64 1
%2 = load i32, ptr %arrayidx.1, align 4 %2 = load i32, ptr %arrayidx.1, align 4
%arrayidx3.1 = getelementptr inbounds [2 x i32], ptr %ret, i64 0, i64 1 %arrayidx3.1 = getelementptr inbounds [2 x i32], ptr %ret, i64 0, i64 1
store i32 %2, ptr %arrayidx3.1, align 4 store i32 %2, ptr %arrayidx3.1, align 4
ret void ret void
} }
define dso_local void @foo_St4x3(ptr nocapture noundef readonly byval(%struct.St4x3) align 4 %in, ptr nocapture noundef writeonly %ret) { define dso_local void @foo_St4x3(ptr nocapture noundef readonly byval(%struct.St4x3) align 4 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func foo_St4x3( ; CHECK-LABEL: .visible .func foo_St4x3(
; CHECK: .param .align 4 .b8 foo_St4x3_param_0[12], ; CHECK: .param .align 4 .b8 foo_St4x3_param_0[12],
; CHECK: .param .b32 foo_St4x3_param_1 ; CHECK: .param .b64 foo_St4x3_param_1
; CHECK: ) ; CHECK: )
; CHECK: ld.param.u32 [[R1:%r[0-9]+]], [foo_St4x3_param_1]; ; CHECK: ld.param.u64 [[R1:%rd[0-9]+]], [foo_St4x3_param_1];
; CHECK: ld.param.u32 [[R2:%r[0-9]+]], [foo_St4x3_param_0]; ; CHECK: ld.param.u32 [[R2:%r[0-9]+]], [foo_St4x3_param_0];
; CHECK: st.u32 [[[R1]]], [[R2]]; ; CHECK: st.u32 [[[R1]]], [[R2]];
; CHECK: ld.param.u32 [[R3:%r[0-9]+]], [foo_St4x3_param_0+4]; ; CHECK: ld.param.u32 [[R3:%r[0-9]+]], [foo_St4x3_param_0+4];
; CHECK: st.u32 [[[R1]]+4], [[R3]]; ; CHECK: st.u32 [[[R1]]+4], [[R3]];
; CHECK: ld.param.u32 [[R4:%r[0-9]+]], [foo_St4x3_param_0+8]; ; CHECK: ld.param.u32 [[R4:%r[0-9]+]], [foo_St4x3_param_0+8];
; CHECK: st.u32 [[[R1]]+8], [[R4]]; ; CHECK: st.u32 [[[R1]]+8], [[R4]];
; CHECK: ret; ; CHECK: ret;
%1 = load i32, ptr %in, align 4 %1 = load i32, ptr %in, align 4
store i32 %1, ptr %ret, align 4 store i32 %1, ptr %ret, align 4
%arrayidx.1 = getelementptr inbounds [3 x i32], ptr %in, i64 0, i64 1 %arrayidx.1 = getelementptr inbounds [3 x i32], ptr %in, i64 0, i64 1
%2 = load i32, ptr %arrayidx.1, align 4 %2 = load i32, ptr %arrayidx.1, align 4
%arrayidx3.1 = getelementptr inbounds [3 x i32], ptr %ret, i64 0, i64 1 %arrayidx3.1 = getelementptr inbounds [3 x i32], ptr %ret, i64 0, i64 1
store i32 %2, ptr %arrayidx3.1, align 4 store i32 %2, ptr %arrayidx3.1, align 4
%arrayidx.2 = getelementptr inbounds [3 x i32], ptr %in, i64 0, i64 2 %arrayidx.2 = getelementptr inbounds [3 x i32], ptr %in, i64 0, i64 2
%3 = load i32, ptr %arrayidx.2, align 4 %3 = load i32, ptr %arrayidx.2, align 4
%arrayidx3.2 = getelementptr inbounds [3 x i32], ptr %ret, i64 0, i64 2 %arrayidx3.2 = getelementptr inbounds [3 x i32], ptr %ret, i64 0, i64 2
store i32 %3, ptr %arrayidx3.2, align 4 store i32 %3, ptr %arrayidx3.2, align 4
ret void ret void
} }
define dso_local void @foo_St4x4(ptr nocapture noundef readonly byval(%struct.St4x4) align 4 %in, ptr nocapture noundef writeonly %ret) { define dso_local void @foo_St4x4(ptr nocapture noundef readonly byval(%struct.St4x4) align 4 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func foo_St4x4( ; CHECK-LABEL: .visible .func foo_St4x4(
; CHECK: .param .align 4 .b8 foo_St4x4_param_0[16], ; CHECK: .param .align 4 .b8 foo_St4x4_param_0[16],
; CHECK: .param .b32 foo_St4x4_param_1 ; CHECK: .param .b64 foo_St4x4_param_1
; CHECK: ) ; CHECK: )
; CHECK: ld.param.u32 [[R1:%r[0-9]+]], [foo_St4x4_param_1]; ; CHECK: ld.param.u64 [[R1:%rd[0-9]+]], [foo_St4x4_param_1];
; CHECK: ld.param.u32 [[R2:%r[0-9]+]], [foo_St4x4_param_0]; ; CHECK: ld.param.u32 [[R2:%r[0-9]+]], [foo_St4x4_param_0];
; CHECK: st.u32 [[[R1]]], [[R2]]; ; CHECK: st.u32 [[[R1]]], [[R2]];
; CHECK: ld.param.u32 [[R3:%r[0-9]+]], [foo_St4x4_param_0+4]; ; CHECK: ld.param.u32 [[R3:%r[0-9]+]], [foo_St4x4_param_0+4];
; CHECK: st.u32 [[[R1]]+4], [[R3]]; ; CHECK: st.u32 [[[R1]]+4], [[R3]];
; CHECK: ld.param.u32 [[R4:%r[0-9]+]], [foo_St4x4_param_0+8]; ; CHECK: ld.param.u32 [[R4:%r[0-9]+]], [foo_St4x4_param_0+8];
; CHECK: st.u32 [[[R1]]+8], [[R4]]; ; CHECK: st.u32 [[[R1]]+8], [[R4]];
; CHECK: ld.param.u32 [[R5:%r[0-9]+]], [foo_St4x4_param_0+12]; ; CHECK: ld.param.u32 [[R5:%r[0-9]+]], [foo_St4x4_param_0+12];
; CHECK: st.u32 [[[R1]]+12], [[R5]]; ; CHECK: st.u32 [[[R1]]+12], [[R5]];
Show All 13 Linesdefine dso_local void @foo_St4x4(ptr nocapture noundef readonly byval(%struct.St4x4) align 4 %in, ptr nocapture noundef writeonly %ret) {
%arrayidx3.3 = getelementptr inbounds [4 x i32], ptr %ret, i64 0, i64 3 %arrayidx3.3 = getelementptr inbounds [4 x i32], ptr %ret, i64 0, i64 3
store i32 %4, ptr %arrayidx3.3, align 4 store i32 %4, ptr %arrayidx3.3, align 4
ret void ret void
} }
define dso_local void @foo_St4x5(ptr nocapture noundef readonly byval(%struct.St4x5) align 4 %in, ptr nocapture noundef writeonly %ret) { define dso_local void @foo_St4x5(ptr nocapture noundef readonly byval(%struct.St4x5) align 4 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func foo_St4x5( ; CHECK-LABEL: .visible .func foo_St4x5(
; CHECK: .param .align 4 .b8 foo_St4x5_param_0[20], ; CHECK: .param .align 4 .b8 foo_St4x5_param_0[20],
; CHECK: .param .b32 foo_St4x5_param_1 ; CHECK: .param .b64 foo_St4x5_param_1
; CHECK: ) ; CHECK: )
; CHECK: ld.param.u32 [[R1:%r[0-9]+]], [foo_St4x5_param_1]; ; CHECK: ld.param.u64 [[R1:%rd[0-9]+]], [foo_St4x5_param_1];
; CHECK: ld.param.u32 [[R2:%r[0-9]+]], [foo_St4x5_param_0]; ; CHECK: ld.param.u32 [[R2:%r[0-9]+]], [foo_St4x5_param_0];
; CHECK: st.u32 [[[R1]]], [[R2]]; ; CHECK: st.u32 [[[R1]]], [[R2]];
; CHECK: ld.param.u32 [[R3:%r[0-9]+]], [foo_St4x5_param_0+4]; ; CHECK: ld.param.u32 [[R3:%r[0-9]+]], [foo_St4x5_param_0+4];
; CHECK: st.u32 [[[R1]]+4], [[R3]]; ; CHECK: st.u32 [[[R1]]+4], [[R3]];
; CHECK: ld.param.u32 [[R4:%r[0-9]+]], [foo_St4x5_param_0+8]; ; CHECK: ld.param.u32 [[R4:%r[0-9]+]], [foo_St4x5_param_0+8];
; CHECK: st.u32 [[[R1]]+8], [[R4]]; ; CHECK: st.u32 [[[R1]]+8], [[R4]];
; CHECK: ld.param.u32 [[R5:%r[0-9]+]], [foo_St4x5_param_0+12]; ; CHECK: ld.param.u32 [[R5:%r[0-9]+]], [foo_St4x5_param_0+12];
; CHECK: st.u32 [[[R1]]+12], [[R5]]; ; CHECK: st.u32 [[[R1]]+12], [[R5]];
Show All 19 Linesdefine dso_local void @foo_St4x5(ptr nocapture noundef readonly byval(%struct.St4x5) align 4 %in, ptr nocapture noundef writeonly %ret) {
%arrayidx3.4 = getelementptr inbounds [5 x i32], ptr %ret, i64 0, i64 4 %arrayidx3.4 = getelementptr inbounds [5 x i32], ptr %ret, i64 0, i64 4
store i32 %5, ptr %arrayidx3.4, align 4 store i32 %5, ptr %arrayidx3.4, align 4
ret void ret void
} }
define dso_local void @foo_St4x6(ptr nocapture noundef readonly byval(%struct.St4x6) align 4 %in, ptr nocapture noundef writeonly %ret) { define dso_local void @foo_St4x6(ptr nocapture noundef readonly byval(%struct.St4x6) align 4 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func foo_St4x6( ; CHECK-LABEL: .visible .func foo_St4x6(
; CHECK: .param .align 4 .b8 foo_St4x6_param_0[24], ; CHECK: .param .align 4 .b8 foo_St4x6_param_0[24],
; CHECK: .param .b32 foo_St4x6_param_1 ; CHECK: .param .b64 foo_St4x6_param_1
; CHECK: ) ; CHECK: )
; CHECK: ld.param.u32 [[R1:%r[0-9]+]], [foo_St4x6_param_1]; ; CHECK: ld.param.u64 [[R1:%rd[0-9]+]], [foo_St4x6_param_1];
; CHECK: ld.param.u32 [[R2:%r[0-9]+]], [foo_St4x6_param_0]; ; CHECK: ld.param.u32 [[R2:%r[0-9]+]], [foo_St4x6_param_0];
; CHECK: st.u32 [[[R1]]], [[R2]]; ; CHECK: st.u32 [[[R1]]], [[R2]];
; CHECK: ld.param.u32 [[R3:%r[0-9]+]], [foo_St4x6_param_0+4]; ; CHECK: ld.param.u32 [[R3:%r[0-9]+]], [foo_St4x6_param_0+4];
; CHECK: st.u32 [[[R1]]+4], [[R3]]; ; CHECK: st.u32 [[[R1]]+4], [[R3]];
; CHECK: ld.param.u32 [[R4:%r[0-9]+]], [foo_St4x6_param_0+8]; ; CHECK: ld.param.u32 [[R4:%r[0-9]+]], [foo_St4x6_param_0+8];
; CHECK: st.u32 [[[R1]]+8], [[R4]]; ; CHECK: st.u32 [[[R1]]+8], [[R4]];
; CHECK: ld.param.u32 [[R5:%r[0-9]+]], [foo_St4x6_param_0+12]; ; CHECK: ld.param.u32 [[R5:%r[0-9]+]], [foo_St4x6_param_0+12];
; CHECK: st.u32 [[[R1]]+12], [[R5]]; ; CHECK: st.u32 [[[R1]]+12], [[R5]];
Show All 25 Linesdefine dso_local void @foo_St4x6(ptr nocapture noundef readonly byval(%struct.St4x6) align 4 %in, ptr nocapture noundef writeonly %ret) {
%arrayidx3.5 = getelementptr inbounds [6 x i32], ptr %ret, i64 0, i64 5 %arrayidx3.5 = getelementptr inbounds [6 x i32], ptr %ret, i64 0, i64 5
store i32 %6, ptr %arrayidx3.5, align 4 store i32 %6, ptr %arrayidx3.5, align 4
ret void ret void
} }
define dso_local void @foo_St4x7(ptr nocapture noundef readonly byval(%struct.St4x7) align 4 %in, ptr nocapture noundef writeonly %ret) { define dso_local void @foo_St4x7(ptr nocapture noundef readonly byval(%struct.St4x7) align 4 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func foo_St4x7( ; CHECK-LABEL: .visible .func foo_St4x7(
; CHECK: .param .align 4 .b8 foo_St4x7_param_0[28], ; CHECK: .param .align 4 .b8 foo_St4x7_param_0[28],
; CHECK: .param .b32 foo_St4x7_param_1 ; CHECK: .param .b64 foo_St4x7_param_1
; CHECK: ) ; CHECK: )
; CHECK: ld.param.u32 [[R1:%r[0-9]+]], [foo_St4x7_param_1]; ; CHECK: ld.param.u64 [[R1:%rd[0-9]+]], [foo_St4x7_param_1];
; CHECK: ld.param.u32 [[R2:%r[0-9]+]], [foo_St4x7_param_0]; ; CHECK: ld.param.u32 [[R2:%r[0-9]+]], [foo_St4x7_param_0];
; CHECK: st.u32 [[[R1]]], [[R2]]; ; CHECK: st.u32 [[[R1]]], [[R2]];
; CHECK: ld.param.u32 [[R3:%r[0-9]+]], [foo_St4x7_param_0+4]; ; CHECK: ld.param.u32 [[R3:%r[0-9]+]], [foo_St4x7_param_0+4];
; CHECK: st.u32 [[[R1]]+4], [[R3]]; ; CHECK: st.u32 [[[R1]]+4], [[R3]];
; CHECK: ld.param.u32 [[R4:%r[0-9]+]], [foo_St4x7_param_0+8]; ; CHECK: ld.param.u32 [[R4:%r[0-9]+]], [foo_St4x7_param_0+8];
; CHECK: st.u32 [[[R1]]+8], [[R4]]; ; CHECK: st.u32 [[[R1]]+8], [[R4]];
; CHECK: ld.param.u32 [[R5:%r[0-9]+]], [foo_St4x7_param_0+12]; ; CHECK: ld.param.u32 [[R5:%r[0-9]+]], [foo_St4x7_param_0+12];
; CHECK: st.u32 [[[R1]]+12], [[R5]]; ; CHECK: st.u32 [[[R1]]+12], [[R5]];
Show All 31 Linesdefine dso_local void @foo_St4x7(ptr nocapture noundef readonly byval(%struct.St4x7) align 4 %in, ptr nocapture noundef writeonly %ret) {
%arrayidx3.6 = getelementptr inbounds [7 x i32], ptr %ret, i64 0, i64 6 %arrayidx3.6 = getelementptr inbounds [7 x i32], ptr %ret, i64 0, i64 6
store i32 %7, ptr %arrayidx3.6, align 4 store i32 %7, ptr %arrayidx3.6, align 4
ret void ret void
} }
define dso_local void @foo_St4x8(ptr nocapture noundef readonly byval(%struct.St4x8) align 4 %in, ptr nocapture noundef writeonly %ret) { define dso_local void @foo_St4x8(ptr nocapture noundef readonly byval(%struct.St4x8) align 4 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func foo_St4x8( ; CHECK-LABEL: .visible .func foo_St4x8(
; CHECK: .param .align 4 .b8 foo_St4x8_param_0[32], ; CHECK: .param .align 4 .b8 foo_St4x8_param_0[32],
; CHECK: .param .b32 foo_St4x8_param_1 ; CHECK: .param .b64 foo_St4x8_param_1
; CHECK: ) ; CHECK: )
; CHECK: ld.param.u32 [[R1:%r[0-9]+]], [foo_St4x8_param_1]; ; CHECK: ld.param.u64 [[R1:%rd[0-9]+]], [foo_St4x8_param_1];
; CHECK: ld.param.u32 [[R2:%r[0-9]+]], [foo_St4x8_param_0]; ; CHECK: ld.param.u32 [[R2:%r[0-9]+]], [foo_St4x8_param_0];
; CHECK: st.u32 [[[R1]]], [[R2]]; ; CHECK: st.u32 [[[R1]]], [[R2]];
; CHECK: ld.param.u32 [[R3:%r[0-9]+]], [foo_St4x8_param_0+4]; ; CHECK: ld.param.u32 [[R3:%r[0-9]+]], [foo_St4x8_param_0+4];
; CHECK: st.u32 [[[R1]]+4], [[R3]]; ; CHECK: st.u32 [[[R1]]+4], [[R3]];
; CHECK: ld.param.u32 [[R4:%r[0-9]+]], [foo_St4x8_param_0+8]; ; CHECK: ld.param.u32 [[R4:%r[0-9]+]], [foo_St4x8_param_0+8];
; CHECK: st.u32 [[[R1]]+8], [[R4]]; ; CHECK: st.u32 [[[R1]]+8], [[R4]];
; CHECK: ld.param.u32 [[R5:%r[0-9]+]], [foo_St4x8_param_0+12]; ; CHECK: ld.param.u32 [[R5:%r[0-9]+]], [foo_St4x8_param_0+12];
; CHECK: st.u32 [[[R1]]+12], [[R5]]; ; CHECK: st.u32 [[[R1]]+12], [[R5]];
Show All 37 Linesdefine dso_local void @foo_St4x8(ptr nocapture noundef readonly byval(%struct.St4x8) align 4 %in, ptr nocapture noundef writeonly %ret) {
%arrayidx3.7 = getelementptr inbounds [8 x i32], ptr %ret, i64 0, i64 7 %arrayidx3.7 = getelementptr inbounds [8 x i32], ptr %ret, i64 0, i64 7
store i32 %8, ptr %arrayidx3.7, align 4 store i32 %8, ptr %arrayidx3.7, align 4
ret void ret void
} }
define dso_local void @foo_St8x1(ptr nocapture noundef readonly byval(%struct.St8x1) align 8 %in, ptr nocapture noundef writeonly %ret) { define dso_local void @foo_St8x1(ptr nocapture noundef readonly byval(%struct.St8x1) align 8 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func foo_St8x1( ; CHECK-LABEL: .visible .func foo_St8x1(
; CHECK: .param .align 8 .b8 foo_St8x1_param_0[8], ; CHECK: .param .align 8 .b8 foo_St8x1_param_0[8],
; CHECK: .param .b32 foo_St8x1_param_1 ; CHECK: .param .b64 foo_St8x1_param_1
; CHECK: ) ; CHECK: )
; CHECK: ld.param.u32 [[R1:%r[0-9]+]], [foo_St8x1_param_1]; ; CHECK: ld.param.u64 [[R1:%rd[0-9]+]], [foo_St8x1_param_1];
; CHECK: ld.param.u64 [[RD1:%rd[0-9]+]], [foo_St8x1_param_0]; ; CHECK: ld.param.u64 [[RD1:%rd[0-9]+]], [foo_St8x1_param_0];
; CHECK: st.u64 [[[R1]]], [[RD1]]; ; CHECK: st.u64 [[[R1]]], [[RD1]];
; CHECK: ret; ; CHECK: ret;
%1 = load i64, ptr %in, align 8 %1 = load i64, ptr %in, align 8
store i64 %1, ptr %ret, align 8 store i64 %1, ptr %ret, align 8
ret void ret void
} }
define dso_local void @foo_St8x2(ptr nocapture noundef readonly byval(%struct.St8x2) align 8 %in, ptr nocapture noundef writeonly %ret) { define dso_local void @foo_St8x2(ptr nocapture noundef readonly byval(%struct.St8x2) align 8 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func foo_St8x2( ; CHECK-LABEL: .visible .func foo_St8x2(
; CHECK: .param .align 8 .b8 foo_St8x2_param_0[16], ; CHECK: .param .align 8 .b8 foo_St8x2_param_0[16],
; CHECK: .param .b32 foo_St8x2_param_1 ; CHECK: .param .b64 foo_St8x2_param_1
; CHECK: ) ; CHECK: )
; CHECK: ld.param.u32 [[R1:%r[0-9]+]], [foo_St8x2_param_1]; ; CHECK: ld.param.u64 [[R1:%rd[0-9]+]], [foo_St8x2_param_1];
; CHECK: ld.param.u64 [[RD1:%rd[0-9]+]], [foo_St8x2_param_0]; ; CHECK: ld.param.u64 [[RD1:%rd[0-9]+]], [foo_St8x2_param_0];
; CHECK: st.u64 [[[R1]]], [[RD1]]; ; CHECK: st.u64 [[[R1]]], [[RD1]];
; CHECK: ld.param.u64 [[RD2:%rd[0-9]+]], [foo_St8x2_param_0+8]; ; CHECK: ld.param.u64 [[RD2:%rd[0-9]+]], [foo_St8x2_param_0+8];
; CHECK: st.u64 [[[R1]]+8], [[RD2]]; ; CHECK: st.u64 [[[R1]]+8], [[RD2]];
; CHECK: ret; ; CHECK: ret;
%1 = load i64, ptr %in, align 8 %1 = load i64, ptr %in, align 8
store i64 %1, ptr %ret, align 8 store i64 %1, ptr %ret, align 8
%arrayidx.1 = getelementptr inbounds [2 x i64], ptr %in, i64 0, i64 1 %arrayidx.1 = getelementptr inbounds [2 x i64], ptr %in, i64 0, i64 1
%2 = load i64, ptr %arrayidx.1, align 8 %2 = load i64, ptr %arrayidx.1, align 8
%arrayidx3.1 = getelementptr inbounds [2 x i64], ptr %ret, i64 0, i64 1 %arrayidx3.1 = getelementptr inbounds [2 x i64], ptr %ret, i64 0, i64 1
store i64 %2, ptr %arrayidx3.1, align 8 store i64 %2, ptr %arrayidx3.1, align 8
ret void ret void
} }
define dso_local void @foo_St8x3(ptr nocapture noundef readonly byval(%struct.St8x3) align 8 %in, ptr nocapture noundef writeonly %ret) { define dso_local void @foo_St8x3(ptr nocapture noundef readonly byval(%struct.St8x3) align 8 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func foo_St8x3( ; CHECK-LABEL: .visible .func foo_St8x3(
; CHECK: .param .align 8 .b8 foo_St8x3_param_0[24], ; CHECK: .param .align 8 .b8 foo_St8x3_param_0[24],
; CHECK: .param .b32 foo_St8x3_param_1 ; CHECK: .param .b64 foo_St8x3_param_1
; CHECK: ) ; CHECK: )
; CHECK: ld.param.u32 [[R1:%r[0-9]+]], [foo_St8x3_param_1]; ; CHECK: ld.param.u64 [[R1:%rd[0-9]+]], [foo_St8x3_param_1];
; CHECK: ld.param.u64 [[RD1:%rd[0-9]+]], [foo_St8x3_param_0]; ; CHECK: ld.param.u64 [[RD1:%rd[0-9]+]], [foo_St8x3_param_0];
; CHECK: st.u64 [[[R1]]], [[RD1]]; ; CHECK: st.u64 [[[R1]]], [[RD1]];
; CHECK: ld.param.u64 [[RD2:%rd[0-9]+]], [foo_St8x3_param_0+8]; ; CHECK: ld.param.u64 [[RD2:%rd[0-9]+]], [foo_St8x3_param_0+8];
; CHECK: st.u64 [[[R1]]+8], [[RD2]]; ; CHECK: st.u64 [[[R1]]+8], [[RD2]];
; CHECK: ld.param.u64 [[RD3:%rd[0-9]+]], [foo_St8x3_param_0+16]; ; CHECK: ld.param.u64 [[RD3:%rd[0-9]+]], [foo_St8x3_param_0+16];
; CHECK: st.u64 [[[R1]]+16], [[RD3]]; ; CHECK: st.u64 [[[R1]]+16], [[RD3]];
; CHECK: ret; ; CHECK: ret;
%1 = load i64, ptr %in, align 8 %1 = load i64, ptr %in, align 8
store i64 %1, ptr %ret, align 8 store i64 %1, ptr %ret, align 8
%arrayidx.1 = getelementptr inbounds [3 x i64], ptr %in, i64 0, i64 1 %arrayidx.1 = getelementptr inbounds [3 x i64], ptr %in, i64 0, i64 1
%2 = load i64, ptr %arrayidx.1, align 8 %2 = load i64, ptr %arrayidx.1, align 8
%arrayidx3.1 = getelementptr inbounds [3 x i64], ptr %ret, i64 0, i64 1 %arrayidx3.1 = getelementptr inbounds [3 x i64], ptr %ret, i64 0, i64 1
store i64 %2, ptr %arrayidx3.1, align 8 store i64 %2, ptr %arrayidx3.1, align 8
%arrayidx.2 = getelementptr inbounds [3 x i64], ptr %in, i64 0, i64 2 %arrayidx.2 = getelementptr inbounds [3 x i64], ptr %in, i64 0, i64 2
%3 = load i64, ptr %arrayidx.2, align 8 %3 = load i64, ptr %arrayidx.2, align 8
%arrayidx3.2 = getelementptr inbounds [3 x i64], ptr %ret, i64 0, i64 2 %arrayidx3.2 = getelementptr inbounds [3 x i64], ptr %ret, i64 0, i64 2
store i64 %3, ptr %arrayidx3.2, align 8 store i64 %3, ptr %arrayidx3.2, align 8
ret void ret void
} }
define dso_local void @foo_St8x4(ptr nocapture noundef readonly byval(%struct.St8x4) align 8 %in, ptr nocapture noundef writeonly %ret) { define dso_local void @foo_St8x4(ptr nocapture noundef readonly byval(%struct.St8x4) align 8 %in, ptr nocapture noundef writeonly %ret) {
; CHECK-LABEL: .visible .func foo_St8x4( ; CHECK-LABEL: .visible .func foo_St8x4(
; CHECK: .param .align 8 .b8 foo_St8x4_param_0[32], ; CHECK: .param .align 8 .b8 foo_St8x4_param_0[32],
; CHECK: .param .b32 foo_St8x4_param_1 ; CHECK: .param .b64 foo_St8x4_param_1
; CHECK: ) ; CHECK: )
; CHECK: ld.param.u32 [[R1:%r[0-9]+]], [foo_St8x4_param_1]; ; CHECK: ld.param.u64 [[R1:%rd[0-9]+]], [foo_St8x4_param_1];
; CHECK: ld.param.u64 [[RD1:%rd[0-9]+]], [foo_St8x4_param_0]; ; CHECK: ld.param.u64 [[RD1:%rd[0-9]+]], [foo_St8x4_param_0];
; CHECK: st.u64 [[[R1]]], [[RD1]]; ; CHECK: st.u64 [[[R1]]], [[RD1]];
; CHECK: ld.param.u64 [[RD2:%rd[0-9]+]], [foo_St8x4_param_0+8]; ; CHECK: ld.param.u64 [[RD2:%rd[0-9]+]], [foo_St8x4_param_0+8];
; CHECK: st.u64 [[[R1]]+8], [[RD2]]; ; CHECK: st.u64 [[[R1]]+8], [[RD2]];
; CHECK: ld.param.u64 [[RD3:%rd[0-9]+]], [foo_St8x4_param_0+16]; ; CHECK: ld.param.u64 [[RD3:%rd[0-9]+]], [foo_St8x4_param_0+16];
; CHECK: st.u64 [[[R1]]+16], [[RD3]]; ; CHECK: st.u64 [[[R1]]+16], [[RD3]];
; CHECK: ld.param.u64 [[RD4:%rd[0-9]+]], [foo_St8x4_param_0+24]; ; CHECK: ld.param.u64 [[RD4:%rd[0-9]+]], [foo_St8x4_param_0+24];
; CHECK: st.u64 [[[R1]]+24], [[RD4]]; ; CHECK: st.u64 [[[R1]]+24], [[RD4]];
Show All 17 Lines

llvm/test/CodeGen/NVPTX/pow2_mask_cmp.ll

; RUN: llc -march=nvptx -verify-machineinstrs < %s | FileCheck %s ; RUN: llc -march=nvptx64 -verify-machineinstrs < %s | FileCheck %s
; RUN: %if ptxas %{ llc -march=nvptx -verify-machineinstrs < %s | %ptxas-verify %} ; RUN: %if ptxas %{ llc -march=nvptx64 -verify-machineinstrs < %s | %ptxas-verify %}
; Tests the following pattern: ; Tests the following pattern:
; (X & 8) != 0 --> (X & 8) >> 3 ; (X & 8) != 0 --> (X & 8) >> 3
; This produces incorrect code in general when boolean false is ; This produces incorrect code in general when boolean false is
; represented as a negative one. There is however a special ; represented as a negative one. There is however a special
; case when the type has a bitsize of 1, for which the false ; case when the type has a bitsize of 1, for which the false
; value will be identical regardless of the boolean representation. ; value will be identical regardless of the boolean representation.
Show All 10 Lines

llvm/test/CodeGen/NVPTX/pr13291-i1-store.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix=PTX32 ; 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 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefix=PTX64
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
define ptx_kernel void @t1(ptr %a) { define ptx_kernel void @t1(ptr %a) {
; PTX32: mov.u16 %rs{{[0-9]+}}, 0; ; PTX32: mov.u16 %rs{{[0-9]+}}, 0;
; PTX32-NEXT: st.global.u8 [%r{{[0-9]+}}], %rs{{[0-9]+}}; ; PTX32-NEXT: st.global.u8 [%r{{[0-9]+}}], %rs{{[0-9]+}};
; PTX64: mov.u16 %rs{{[0-9]+}}, 0; ; PTX64: mov.u16 %rs{{[0-9]+}}, 0;
; PTX64-NEXT: st.global.u8 [%rd{{[0-9]+}}], %rs{{[0-9]+}}; ; PTX64-NEXT: st.global.u8 [%rd{{[0-9]+}}], %rs{{[0-9]+}};
store i1 false, ptr %a store i1 false, ptr %a
Show All 17 Lines

llvm/test/CodeGen/NVPTX/pr16278.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
@one_f = addrspace(4) global float 1.000000e+00, align 4 @one_f = addrspace(4) global float 1.000000e+00, align 4
define float @foo() { define float @foo() {
; CHECK: ld.const.f32 ; CHECK: ld.const.f32
%val = load float, ptr addrspace(4) @one_f %val = load float, ptr addrspace(4) @one_f
ret float %val ret float %val
} }

llvm/test/CodeGen/NVPTX/pr17529.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
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 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-nvidia-cuda" target triple = "nvptx64-nvidia-cuda"
; Function Attrs: nounwind ; Function Attrs: nounwind
; CHECK: .func kernelgen_memcpy ; CHECK: .func kernelgen_memcpy
define ptx_device void @kernelgen_memcpy(ptr nocapture %dst) #0 { define ptx_device void @kernelgen_memcpy(ptr nocapture %dst) #0 {
entry: entry:
Show All 27 Lines

llvm/test/CodeGen/NVPTX/read-global-variable-constant.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_35 -verify-machineinstrs | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 -verify-machineinstrs | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_35 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_35 | %ptxas-verify %}
; Check load from constant global variables. These loads should be ; Check load from constant global variables. These loads should be
; ld.global.nc (aka ldg). ; ld.global.nc (aka ldg).
@gv_float = external constant float @gv_float = external constant float
@gv_float2 = external constant <2 x float> @gv_float2 = external constant <2 x float>
@gv_float4 = external constant <4 x float> @gv_float4 = external constant <4 x float>
Show All 20 Lines

llvm/test/CodeGen/NVPTX/refl1.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
target triple = "nvptx-nvidia-cuda" target triple = "nvptx-nvidia-cuda"
; Function Attrs: nounwind ; Function Attrs: nounwind
; CHECK: .entry foo ; CHECK: .entry foo
define void @foo(ptr nocapture %a) #0 { define void @foo(ptr nocapture %a) #0 {
%val = load float, ptr %a %val = load float, ptr %a
%tan = tail call fastcc float @__nv_fast_tanf(float %val) %tan = tail call fastcc float @__nv_fast_tanf(float %val)
Show All 30 Lines

llvm/test/CodeGen/NVPTX/reg-types.ll

; Verify register types we generate in PTX. ; Verify register types we generate in PTX.
; RUN: llc -O0 < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc -O0 < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
; RUN: llc -O0 < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s ; RUN: llc -O0 < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: llc -O0 < %s -march=nvptx -mcpu=sm_20 | FileCheck %s -check-prefixes=NO8BIT ; RUN: llc -O0 < %s -march=nvptx -mcpu=sm_20 | FileCheck %s -check-prefixes=NO8BIT
; RUN: llc -O0 < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s -check-prefixes=NO8BIT ; RUN: llc -O0 < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s -check-prefixes=NO8BIT
; RUN: %if ptxas %{ llc -O0 < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc -O0 < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{ llc -O0 < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc -O0 < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; CHECK-LABEL: .visible .func func() ; CHECK-LABEL: .visible .func func()
; NO8BIT-LABEL: .visible .func func() ; NO8BIT-LABEL: .visible .func func()
define void @func() { define void @func() {
entry: entry:
%s8 = alloca i8, align 1 %s8 = alloca i8, align 1
%u8 = alloca i8, align 1 %u8 = alloca i8, align 1
▲ Show 20 Lines • Show All 57 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/rotate.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck --check-prefix=SM20 %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck --check-prefix=SM20 %s
; RUN: llc < %s -march=nvptx -mcpu=sm_35 | FileCheck --check-prefix=SM35 %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 | FileCheck --check-prefix=SM35 %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_35 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_35 | %ptxas-verify %}
declare i32 @llvm.nvvm.rotate.b32(i32, i32) declare i32 @llvm.nvvm.rotate.b32(i32, i32)
declare i64 @llvm.nvvm.rotate.b64(i64, i32) declare i64 @llvm.nvvm.rotate.b64(i64, i32)
declare i64 @llvm.nvvm.rotate.right.b64(i64, i32) declare i64 @llvm.nvvm.rotate.right.b64(i64, i32)
; SM20: rotate32 ; SM20: rotate32
; SM35: rotate32 ; SM35: rotate32
▲ Show 20 Lines • Show All 48 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/rotate_64.ll

; RUN: llc < %s -march=nvptx | FileCheck %s ; RUN: llc < %s -march=nvptx64 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 | %ptxas-verify %}
declare i64 @llvm.nvvm.rotate.b64(i64, i32) declare i64 @llvm.nvvm.rotate.b64(i64, i32)
declare i64 @llvm.nvvm.rotate.right.b64(i64, i32) declare i64 @llvm.nvvm.rotate.right.b64(i64, i32)
; CHECK: rotate64 ; CHECK: rotate64
define i64 @rotate64(i64 %a, i32 %b) { define i64 @rotate64(i64 %a, i32 %b) {
; CHECK: shl.b64 [[LHS:%.*]], [[RD1:%.*]], 3; ; CHECK: shl.b64 [[LHS:%.*]], [[RD1:%.*]], 3;
; CHECK: shr.b64 [[RHS:%.*]], [[RD1]], 61; ; CHECK: shr.b64 [[RHS:%.*]], [[RD1]], 61;
Show All 15 Lines

llvm/test/CodeGen/NVPTX/sched1.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; Ensure source scheduling is working ; Ensure source scheduling is working
define void @foo(ptr %a) { define void @foo(ptr %a) {
; CHECK: .func foo ; CHECK: .func foo
; CHECK: ld.u32 ; CHECK: ld.u32
; CHECK-NEXT: ld.u32 ; CHECK-NEXT: ld.u32
; CHECK-NEXT: ld.u32 ; CHECK-NEXT: ld.u32
Show All 21 Lines

llvm/test/CodeGen/NVPTX/sched2.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
define void @foo(ptr %a) { define void @foo(ptr %a) {
; CHECK: .func foo ; CHECK: .func foo
; CHECK: ld.v2.u32 ; CHECK: ld.v2.u32
; CHECK-NEXT: ld.v2.u32 ; CHECK-NEXT: ld.v2.u32
; CHECK-NEXT: ld.v2.u32 ; CHECK-NEXT: ld.v2.u32
; CHECK-NEXT: ld.v2.u32 ; CHECK-NEXT: ld.v2.u32
; CHECK-NEXT: add.s32 ; CHECK-NEXT: add.s32
Show All 22 Lines

llvm/test/CodeGen/NVPTX/sext-in-reg.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
target datalayout = "e-p:32:32:32-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 datalayout = "e-p:32:32:32-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"
define void @one(i64 %a, i64 %b, ptr %p1, ptr %p2) { define void @one(i64 %a, i64 %b, ptr %p1, ptr %p2) {
; CHECK: cvt.s64.s8 ; CHECK: cvt.s64.s8
; CHECK: cvt.s64.s8 ; CHECK: cvt.s64.s8
entry: entry:
▲ Show 20 Lines • Show All 102 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/sext-params.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
target datalayout = "e-p:32:32:32-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 datalayout = "e-p:32:32:32-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"
define i8 @foo(i8 signext %a) { define i8 @foo(i8 signext %a) {
; CHECK: ld.param.s8 ; CHECK: ld.param.s8
%ret = add i8 %a, 3 %ret = add i8 %a, 3
ret i8 %ret ret i8 %ret
} }
define i8 @bar(i8 zeroext %a) { define i8 @bar(i8 zeroext %a) {
; CHECK: ld.param.u8 ; CHECK: ld.param.u8
%ret = add i8 %a, 3 %ret = add i8 %a, 3
ret i8 %ret ret i8 %ret
} }

llvm/test/CodeGen/NVPTX/shift-parts.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; CHECK: shift_parts_left_128 ; CHECK: shift_parts_left_128
define void @shift_parts_left_128(ptr %val, ptr %amtptr) { define void @shift_parts_left_128(ptr %val, ptr %amtptr) {
; CHECK: shl.b64 ; CHECK: shl.b64
; CHECK: mov.u32 ; CHECK: mov.u32
; CHECK: sub.s32 ; CHECK: sub.s32
; CHECK: shr.u64 ; CHECK: shr.u64
; CHECK: or.b64 ; CHECK: or.b64
Show All 29 Lines

llvm/test/CodeGen/NVPTX/short-ptr.ll

; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefix CHECK-DEFAULT ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefix CHECK-DEFAULT
; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix CHECK-DEFAULT-32 ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix CHECK-DEFAULT-32
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -nvptx-short-ptr | FileCheck %s --check-prefixes CHECK-SHORT-SHARED,CHECK-SHORT-CONST,CHECK-SHORT-LOCAL ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -nvptx-short-ptr | FileCheck %s --check-prefixes CHECK-SHORT-SHARED,CHECK-SHORT-CONST,CHECK-SHORT-LOCAL
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -nvptx-short-ptr | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -nvptx-short-ptr | %ptxas-verify %}
; CHECK-DEFAULT: .visible .shared .align 8 .u64 s ; CHECK-DEFAULT: .visible .shared .align 8 .u64 s
; CHECK-DEFAULT-32: .visible .shared .align 8 .u32 s ; CHECK-DEFAULT-32: .visible .shared .align 8 .u32 s
; CHECK-SHORT-SHARED: .visible .shared .align 8 .u32 s ; CHECK-SHORT-SHARED: .visible .shared .align 8 .u32 s
@s = local_unnamed_addr addrspace(3) global i32 addrspace(3)* null, align 8 @s = local_unnamed_addr addrspace(3) global i32 addrspace(3)* null, align 8
Show All 31 Lines

llvm/test/CodeGen/NVPTX/simple-call.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s ; RUN: llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
; CHECK: .func ({{.*}}) device_func ; CHECK: .func ({{.*}}) device_func
define float @device_func(float %a) noinline { define float @device_func(float %a) noinline {
%ret = fmul float %a, %a %ret = fmul float %a, %a
ret float %ret ret float %ret
} }
; CHECK: .entry kernel_func ; CHECK: .entry kernel_func
define void @kernel_func(ptr %a) { define void @kernel_func(ptr %a) {
%val = load float, ptr %a %val = load float, ptr %a
; CHECK: call.uni (retval0), ; CHECK: call.uni (retval0),
; CHECK: device_func, ; CHECK: device_func,
%mul = call float @device_func(float %val) %mul = call float @device_func(float %val)
store float %mul, ptr %a store float %mul, ptr %a
ret void ret void
} }
!nvvm.annotations = !{!1} !nvvm.annotations = !{!1}
!1 = !{ptr @kernel_func, !"kernel", i32 1} !1 = !{ptr @kernel_func, !"kernel", i32 1}

llvm/test/CodeGen/NVPTX/sqrt-approx.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 -nvptx-prec-divf32=0 -nvptx-prec-sqrtf32=0 \ ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -nvptx-prec-divf32=0 -nvptx-prec-sqrtf32=0 \
; RUN: | FileCheck %s ; RUN: | FileCheck %s
; RUN: %if ptxas %{ \ ; RUN: %if ptxas %{ \
; RUN: llc < %s -march=nvptx -mcpu=sm_20 -nvptx-prec-divf32=0 -nvptx-prec-sqrtf32=0 \ ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -nvptx-prec-divf32=0 -nvptx-prec-sqrtf32=0 \
; RUN: | %ptxas-verify \ ; RUN: | %ptxas-verify \
; RUN: %} ; RUN: %}
target datalayout = "e-p:32:32:32-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 datalayout = "e-p:32:32:32-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"
declare float @llvm.sqrt.f32(float) declare float @llvm.sqrt.f32(float)
declare double @llvm.sqrt.f64(double) declare double @llvm.sqrt.f64(double)
▲ Show 20 Lines • Show All 198 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/st-addrspace.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefixes=ALL,G32,LS32 ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefixes=ALL,G32,LS32
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefixes=ALL,G64,LS64 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefixes=ALL,G64,LS64
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 --nvptx-short-ptr | FileCheck %s --check-prefixes=G64,LS32 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 --nvptx-short-ptr | FileCheck %s --check-prefixes=G64,LS32
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 --nvptx-short-ptr | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 --nvptx-short-ptr | %ptxas-verify %}
;; i8 ;; i8
; ALL-LABEL: st_global_i8 ; ALL-LABEL: st_global_i8
define void @st_global_i8(i8 addrspace(1)* %ptr, i8 %a) { define void @st_global_i8(i8 addrspace(1)* %ptr, i8 %a) {
; G32: st.global.u8 [%r{{[0-9]+}}], %rs{{[0-9]+}} ; G32: st.global.u8 [%r{{[0-9]+}}], %rs{{[0-9]+}}
; G64: st.global.u8 [%rd{{[0-9]+}}], %rs{{[0-9]+}} ; G64: st.global.u8 [%rd{{[0-9]+}}], %rs{{[0-9]+}}
▲ Show 20 Lines • Show All 150 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/st-generic.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix=PTX32 ; 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 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefix=PTX64
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
;; i8 ;; i8
define void @st_global_i8(ptr addrspace(0) %ptr, i8 %a) { define void @st_global_i8(ptr addrspace(0) %ptr, i8 %a) {
; PTX32: st.u8 [%r{{[0-9]+}}], %rs{{[0-9]+}} ; PTX32: st.u8 [%r{{[0-9]+}}], %rs{{[0-9]+}}
; PTX32: ret ; PTX32: ret
; PTX64: st.u8 [%rd{{[0-9]+}}], %rs{{[0-9]+}} ; PTX64: st.u8 [%rd{{[0-9]+}}], %rs{{[0-9]+}}
▲ Show 20 Lines • Show All 59 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/store-retval.ll

; RUN: llc < %s --mtriple=nvptx-unknown-unknown | FileCheck %s ; RUN: llc < %s --mtriple=nvptx64-unknown-unknown | FileCheck %s
; RUN: %if ptxas %{ llc < %s --mtriple=nvptx-unknown-unknown | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s --mtriple=nvptx64-unknown-unknown | %ptxas-verify %}
; ;
; This is IR generated with clang using -O3 optimization level ; This is IR generated with clang using -O3 optimization level
; and nvptx-unknown-unknown target from the following C code. ; and nvptx-unknown-unknown target from the following C code.
; ;
; struct StNoalign { unsigned int field[5]; }; ; struct StNoalign { unsigned int field[5]; };
; struct StAlign8 { _Alignas(8) unsigned int field[5]; }; ; struct StAlign8 { _Alignas(8) unsigned int field[5]; };
; struct StAlign16 { _Alignas(16) unsigned int field[5]; }; ; struct StAlign16 { _Alignas(16) unsigned int field[5]; };
; ;
▲ Show 20 Lines • Show All 62 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/surf-read-cuda.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s --check-prefix=SM20 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s --check-prefix=SM20
; RUN: llc < %s -march=nvptx -mcpu=sm_30 -verify-machineinstrs | FileCheck %s --check-prefix=SM30 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_30 -verify-machineinstrs | FileCheck %s --check-prefix=SM30
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_30 -verify-machineinstrs | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_30 -verify-machineinstrs | %ptxas-verify %}
target triple = "nvptx-unknown-cuda" target triple = "nvptx-unknown-cuda"
declare i32 @llvm.nvvm.suld.1d.i32.trap(i64, i32) declare i32 @llvm.nvvm.suld.1d.i32.trap(i64, i32)
declare i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1)) declare i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1))
; SM20-LABEL: .entry foo ; SM20-LABEL: .entry foo
; SM30-LABEL: .entry foo ; SM30-LABEL: .entry foo
define void @foo(i64 %img, ptr %red, i32 %idx) { define void @foo(i64 %img, ptr %red, i32 %idx) {
; SM20: ld.param.u64 %rd[[SURFREG:[0-9]+]], [foo_param_0]; ; SM20: ld.param.u64 %rd[[SURFREG:[0-9]+]], [foo_param_0];
; SM20: suld.b.1d.b32.trap {%r[[RED:[0-9]+]]}, [%rd[[SURFREG]], {%r{{[0-9]+}}}] ; SM20: suld.b.1d.b32.trap {%r[[RED:[0-9]+]]}, [%rd[[SURFREG]], {%r{{[0-9]+}}}]
; SM30: ld.param.u64 %rd[[SURFREG:[0-9]+]], [foo_param_0]; ; SM30: ld.param.u64 %rd[[SURFREG:[0-9]+]], [foo_param_0];
; SM30: suld.b.1d.b32.trap {%r[[RED:[0-9]+]]}, [%rd[[SURFREG]], {%r{{[0-9]+}}}] ; SM30: suld.b.1d.b32.trap {%r[[RED:[0-9]+]]}, [%rd[[SURFREG]], {%r{{[0-9]+}}}]
%val = tail call i32 @llvm.nvvm.suld.1d.i32.trap(i64 %img, i32 %idx) %val = tail call i32 @llvm.nvvm.suld.1d.i32.trap(i64 %img, i32 %idx)
; SM20: cvt.rn.f32.s32 %f[[REDF:[0-9]+]], %r[[RED]] ; SM20: cvt.rn.f32.s32 %f[[REDF:[0-9]+]], %r[[RED]]
; SM30: cvt.rn.f32.s32 %f[[REDF:[0-9]+]], %r[[RED]] ; SM30: cvt.rn.f32.s32 %f[[REDF:[0-9]+]], %r[[RED]]
%ret = sitofp i32 %val to float %ret = sitofp i32 %val to float
; SM20: st.global.f32 [%r{{[0-9]+}}], %f[[REDF]] ; SM20: st.global.f32 [%rd{{[0-9]+}}], %f[[REDF]]
; SM30: st.global.f32 [%r{{[0-9]+}}], %f[[REDF]] ; SM30: st.global.f32 [%rd{{[0-9]+}}], %f[[REDF]]
store float %ret, ptr %red store float %ret, ptr %red
ret void ret void
} }
@surf0 = internal addrspace(1) global i64 0, align 8 @surf0 = internal addrspace(1) global i64 0, align 8
; SM20-LABEL: .entry bar ; SM20-LABEL: .entry bar
; SM30-LABEL: .entry bar ; SM30-LABEL: .entry bar
define void @bar(ptr %red, i32 %idx) { define void @bar(ptr %red, i32 %idx) {
; SM30: mov.u64 %rd[[SURFHANDLE:[0-9]+]], surf0 ; SM30: mov.u64 %rd[[SURFHANDLE:[0-9]+]], surf0
%surfHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1) @surf0) %surfHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1) @surf0)
; SM20: suld.b.1d.b32.trap {%r[[RED:[0-9]+]]}, [surf0, {%r{{[0-9]+}}}] ; SM20: suld.b.1d.b32.trap {%r[[RED:[0-9]+]]}, [surf0, {%r{{[0-9]+}}}]
; SM30: suld.b.1d.b32.trap {%r[[RED:[0-9]+]]}, [%rd[[SURFHANDLE]], {%r{{[0-9]+}}}] ; SM30: suld.b.1d.b32.trap {%r[[RED:[0-9]+]]}, [%rd[[SURFHANDLE]], {%r{{[0-9]+}}}]
%val = tail call i32 @llvm.nvvm.suld.1d.i32.trap(i64 %surfHandle, i32 %idx) %val = tail call i32 @llvm.nvvm.suld.1d.i32.trap(i64 %surfHandle, i32 %idx)
; SM20: cvt.rn.f32.s32 %f[[REDF:[0-9]+]], %r[[RED]] ; SM20: cvt.rn.f32.s32 %f[[REDF:[0-9]+]], %r[[RED]]
; SM30: cvt.rn.f32.s32 %f[[REDF:[0-9]+]], %r[[RED]] ; SM30: cvt.rn.f32.s32 %f[[REDF:[0-9]+]], %r[[RED]]
%ret = sitofp i32 %val to float %ret = sitofp i32 %val to float
; SM20: st.global.f32 [%r{{[0-9]+}}], %f[[REDF]] ; SM20: st.global.f32 [%rd{{[0-9]+}}], %f[[REDF]]
; SM30: st.global.f32 [%r{{[0-9]+}}], %f[[REDF]] ; SM30: st.global.f32 [%rd{{[0-9]+}}], %f[[REDF]]
store float %ret, ptr %red store float %ret, ptr %red
ret void ret void
} }
!nvvm.annotations = !{!1, !2, !3} !nvvm.annotations = !{!1, !2, !3}
!1 = !{ptr @foo, !"kernel", i32 1} !1 = !{ptr @foo, !"kernel", i32 1}
!2 = !{ptr @bar, !"kernel", i32 1} !2 = !{ptr @bar, !"kernel", i32 1}
!3 = !{ptr addrspace(1) @surf0, !"surface", i32 1} !3 = !{ptr addrspace(1) @surf0, !"surface", i32 1}

llvm/test/CodeGen/NVPTX/surf-read.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
target triple = "nvptx-unknown-nvcl" target triple = "nvptx-unknown-nvcl"
declare i32 @llvm.nvvm.suld.1d.i32.trap(i64, i32) declare i32 @llvm.nvvm.suld.1d.i32.trap(i64, i32)
; CHECK: .entry foo ; CHECK: .entry foo
define void @foo(i64 %img, ptr %red, i32 %idx) { define void @foo(i64 %img, ptr %red, i32 %idx) {
; CHECK: suld.b.1d.b32.trap {%r[[RED:[0-9]+]]}, [foo_param_0, {%r{{[0-9]+}}}] ; CHECK: suld.b.1d.b32.trap {%r[[RED:[0-9]+]]}, [foo_param_0, {%r{{[0-9]+}}}]
%val = tail call i32 @llvm.nvvm.suld.1d.i32.trap(i64 %img, i32 %idx) %val = tail call i32 @llvm.nvvm.suld.1d.i32.trap(i64 %img, i32 %idx)
; CHECK: cvt.rn.f32.s32 %f[[REDF:[0-9]+]], %r[[RED]] ; CHECK: cvt.rn.f32.s32 %f[[REDF:[0-9]+]], %r[[RED]]
%ret = sitofp i32 %val to float %ret = sitofp i32 %val to float
; CHECK: st.f32 [%r{{[0-9]+}}], %f[[REDF]] ; CHECK: st.f32 [%rd{{[0-9]+}}], %f[[REDF]]
store float %ret, ptr %red store float %ret, ptr %red
ret void ret void
} }
!nvvm.annotations = !{!1, !2} !nvvm.annotations = !{!1, !2}
!1 = !{ptr @foo, !"kernel", i32 1} !1 = !{ptr @foo, !"kernel", i32 1}
!2 = !{ptr @foo, !"rdwrimage", i32 0} !2 = !{ptr @foo, !"rdwrimage", i32 0}

llvm/test/CodeGen/NVPTX/surf-write-cuda.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s --check-prefix=SM20 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s --check-prefix=SM20
; RUN: llc < %s -march=nvptx -mcpu=sm_30 -verify-machineinstrs | FileCheck %s --check-prefix=SM30 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_30 -verify-machineinstrs | FileCheck %s --check-prefix=SM30
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_30 -verify-machineinstrs | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_30 -verify-machineinstrs | %ptxas-verify %}
target triple = "nvptx-unknown-cuda" target triple = "nvptx-unknown-cuda"
declare void @llvm.nvvm.sust.b.1d.i32.trap(i64, i32, i32) declare void @llvm.nvvm.sust.b.1d.i32.trap(i64, i32, i32)
declare i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1)) declare i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1))
; SM20-LABEL: .entry foo ; SM20-LABEL: .entry foo
Show All 32 Lines

llvm/test/CodeGen/NVPTX/surf-write.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
target triple = "nvptx-unknown-nvcl" target triple = "nvptx-unknown-nvcl"
declare void @llvm.nvvm.sust.b.1d.i32.trap(i64, i32, i32) declare void @llvm.nvvm.sust.b.1d.i32.trap(i64, i32, i32)
; CHECK: .entry foo ; CHECK: .entry foo
define void @foo(i64 %img, i32 %val, i32 %idx) { define void @foo(i64 %img, i32 %val, i32 %idx) {
; CHECK: sust.b.1d.b32.trap [foo_param_0, {%r{{[0-9]+}}}], {%r{{[0-9]+}}} ; CHECK: sust.b.1d.b32.trap [foo_param_0, {%r{{[0-9]+}}}], {%r{{[0-9]+}}}
tail call void @llvm.nvvm.sust.b.1d.i32.trap(i64 %img, i32 %idx, i32 %val) tail call void @llvm.nvvm.sust.b.1d.i32.trap(i64 %img, i32 %idx, i32 %val)
ret void ret void
} }
!nvvm.annotations = !{!1, !2} !nvvm.annotations = !{!1, !2}
!1 = !{ptr @foo, !"kernel", i32 1} !1 = !{ptr @foo, !"kernel", i32 1}
!2 = !{ptr @foo, !"wroimage", i32 0} !2 = !{ptr @foo, !"wroimage", i32 0}

llvm/test/CodeGen/NVPTX/symbol-naming.ll

; RUN: llc < %s -march=nvptx -mattr=+ptx60 -mcpu=sm_30 | FileCheck %s ; RUN: llc < %s -march=nvptx -mattr=+ptx60 -mcpu=sm_30 | FileCheck %s
; RUN: llc < %s -march=nvptx64 -mattr=+ptx60 -mcpu=sm_30 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mattr=+ptx60 -mcpu=sm_30 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mattr=+ptx60 -mcpu=sm_30 | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mattr=+ptx60 -mcpu=sm_30 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mattr=+ptx60 -mcpu=sm_30 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mattr=+ptx60 -mcpu=sm_30 | %ptxas-verify %}
; Verify that the NVPTX target removes invalid symbol names prior to emitting ; Verify that the NVPTX target removes invalid symbol names prior to emitting
; PTX. ; PTX.
; CHECK-NOT: .str ; CHECK-NOT: .str
; CHECK-NOT: .function. ; CHECK-NOT: .function.
Show All 37 Lines

llvm/test/CodeGen/NVPTX/tex-read-cuda.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s --check-prefix=SM20 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s --check-prefix=SM20
; RUN: llc < %s -march=nvptx -mcpu=sm_30 -verify-machineinstrs | FileCheck %s --check-prefix=SM30 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_30 -verify-machineinstrs | FileCheck %s --check-prefix=SM30
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_30 -verify-machineinstrs | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_30 -verify-machineinstrs | %ptxas-verify %}
target triple = "nvptx-unknown-cuda" target triple = "nvptx-unknown-cuda"
declare { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64, i32) declare { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64, i32)
declare i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1)) declare i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1))
; SM20-LABEL: .entry foo ; SM20-LABEL: .entry foo
; SM30-LABEL: .entry foo ; SM30-LABEL: .entry foo
define void @foo(i64 %img, ptr %red, i32 %idx) { define void @foo(i64 %img, ptr %red, i32 %idx) {
; SM20: ld.param.u64 %rd[[TEXREG:[0-9]+]], [foo_param_0]; ; SM20: ld.param.u64 %rd[[TEXREG:[0-9]+]], [foo_param_0];
; SM20: tex.1d.v4.f32.s32 {%f[[RED:[0-9]+]], %f[[GREEN:[0-9]+]], %f[[BLUE:[0-9]+]], %f[[ALPHA:[0-9]+]]}, [%rd[[TEXREG]], {%r{{[0-9]+}}}] ; SM20: tex.1d.v4.f32.s32 {%f[[RED:[0-9]+]], %f[[GREEN:[0-9]+]], %f[[BLUE:[0-9]+]], %f[[ALPHA:[0-9]+]]}, [%rd[[TEXREG]], {%r{{[0-9]+}}}]
; SM30: ld.param.u64 %rd[[TEXREG:[0-9]+]], [foo_param_0]; ; SM30: ld.param.u64 %rd[[TEXREG:[0-9]+]], [foo_param_0];
; SM30: tex.1d.v4.f32.s32 {%f[[RED:[0-9]+]], %f[[GREEN:[0-9]+]], %f[[BLUE:[0-9]+]], %f[[ALPHA:[0-9]+]]}, [%rd[[TEXREG]], {%r{{[0-9]+}}}] ; SM30: tex.1d.v4.f32.s32 {%f[[RED:[0-9]+]], %f[[GREEN:[0-9]+]], %f[[BLUE:[0-9]+]], %f[[ALPHA:[0-9]+]]}, [%rd[[TEXREG]], {%r{{[0-9]+}}}]
%val = tail call { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64 %img, i32 %idx) %val = tail call { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64 %img, i32 %idx)
%ret = extractvalue { float, float, float, float } %val, 0 %ret = extractvalue { float, float, float, float } %val, 0
; SM20: st.global.f32 [%r{{[0-9]+}}], %f[[RED]] ; SM20: st.global.f32 [%rd{{[0-9]+}}], %f[[RED]]
; SM30: st.global.f32 [%r{{[0-9]+}}], %f[[RED]] ; SM30: st.global.f32 [%rd{{[0-9]+}}], %f[[RED]]
store float %ret, ptr %red store float %ret, ptr %red
ret void ret void
} }
@tex0 = internal addrspace(1) global i64 0, align 8 @tex0 = internal addrspace(1) global i64 0, align 8
; SM20-LABEL: .entry bar ; SM20-LABEL: .entry bar
; SM30-LABEL: .entry bar ; SM30-LABEL: .entry bar
define void @bar(ptr %red, i32 %idx) { define void @bar(ptr %red, i32 %idx) {
; SM30: mov.u64 %rd[[TEXHANDLE:[0-9]+]], tex0 ; SM30: mov.u64 %rd[[TEXHANDLE:[0-9]+]], tex0
%texHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1) @tex0) %texHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1) @tex0)
; SM20: tex.1d.v4.f32.s32 {%f[[RED:[0-9]+]], %f[[GREEN:[0-9]+]], %f[[BLUE:[0-9]+]], %f[[ALPHA:[0-9]+]]}, [tex0, {%r{{[0-9]+}}}] ; SM20: tex.1d.v4.f32.s32 {%f[[RED:[0-9]+]], %f[[GREEN:[0-9]+]], %f[[BLUE:[0-9]+]], %f[[ALPHA:[0-9]+]]}, [tex0, {%r{{[0-9]+}}}]
; SM30: tex.1d.v4.f32.s32 {%f[[RED:[0-9]+]], %f[[GREEN:[0-9]+]], %f[[BLUE:[0-9]+]], %f[[ALPHA:[0-9]+]]}, [%rd[[TEXHANDLE]], {%r{{[0-9]+}}}] ; SM30: tex.1d.v4.f32.s32 {%f[[RED:[0-9]+]], %f[[GREEN:[0-9]+]], %f[[BLUE:[0-9]+]], %f[[ALPHA:[0-9]+]]}, [%rd[[TEXHANDLE]], {%r{{[0-9]+}}}]
%val = tail call { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64 %texHandle, i32 %idx) %val = tail call { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64 %texHandle, i32 %idx)
%ret = extractvalue { float, float, float, float } %val, 0 %ret = extractvalue { float, float, float, float } %val, 0
; SM20: st.global.f32 [%r{{[0-9]+}}], %f[[RED]] ; SM20: st.global.f32 [%rd{{[0-9]+}}], %f[[RED]]
; SM30: st.global.f32 [%r{{[0-9]+}}], %f[[RED]] ; SM30: st.global.f32 [%rd{{[0-9]+}}], %f[[RED]]
store float %ret, ptr %red store float %ret, ptr %red
ret void ret void
} }
declare float @texfunc(i64) declare float @texfunc(i64)
; SM20-LABEL: .entry baz ; SM20-LABEL: .entry baz
; SM30-LABEL: .entry baz ; SM30-LABEL: .entry baz
Show All 9 Lines
; SM20: texfunc, ; SM20: texfunc,
; SM30: texfunc, ; SM30: texfunc,
%texcall = tail call float @texfunc(i64 %texHandle) %texcall = tail call float @texfunc(i64 %texHandle)
; SM20: ld.param.f32 %f[[TEXCALL:[0-9]+]], [[[RETVAL]]+0] ; SM20: ld.param.f32 %f[[TEXCALL:[0-9]+]], [[[RETVAL]]+0]
; SM30: ld.param.f32 %f[[TEXCALL:[0-9]+]], [[[RETVAL]]+0] ; SM30: ld.param.f32 %f[[TEXCALL:[0-9]+]], [[[RETVAL]]+0]
; SM20: add.rn.f32 %f[[RET2:[0-9]+]], %f[[RED]], %f[[TEXCALL]] ; SM20: add.rn.f32 %f[[RET2:[0-9]+]], %f[[RED]], %f[[TEXCALL]]
; SM30: add.rn.f32 %f[[RET2:[0-9]+]], %f[[RED]], %f[[TEXCALL]] ; SM30: add.rn.f32 %f[[RET2:[0-9]+]], %f[[RED]], %f[[TEXCALL]]
%ret2 = fadd float %ret, %texcall %ret2 = fadd float %ret, %texcall
; SM20: st.global.f32 [%r{{[0-9]+}}], %f[[RET2]] ; SM20: st.global.f32 [%rd{{[0-9]+}}], %f[[RET2]]
; SM30: st.global.f32 [%r{{[0-9]+}}], %f[[RET2]] ; SM30: st.global.f32 [%rd{{[0-9]+}}], %f[[RET2]]
store float %ret2, ptr %red store float %ret2, ptr %red
ret void ret void
} }
!nvvm.annotations = !{!1, !2, !3, !4} !nvvm.annotations = !{!1, !2, !3, !4}
!1 = !{ptr @foo, !"kernel", i32 1} !1 = !{ptr @foo, !"kernel", i32 1}
!2 = !{ptr @bar, !"kernel", i32 1} !2 = !{ptr @bar, !"kernel", i32 1}
!3 = !{ptr addrspace(1) @tex0, !"texture", i32 1} !3 = !{ptr addrspace(1) @tex0, !"texture", i32 1}
!4 = !{ptr @baz, !"kernel", i32 1} !4 = !{ptr @baz, !"kernel", i32 1}

llvm/test/CodeGen/NVPTX/tex-read.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
target triple = "nvptx-unknown-nvcl" target triple = "nvptx-unknown-nvcl"
declare { float, float, float, float } @llvm.nvvm.tex.1d.v4f32.s32(i64, i64, i32) declare { float, float, float, float } @llvm.nvvm.tex.1d.v4f32.s32(i64, i64, i32)
; CHECK: .entry foo ; CHECK: .entry foo
define void @foo(i64 %img, i64 %sampler, ptr %red, i32 %idx) { define void @foo(i64 %img, i64 %sampler, ptr %red, i32 %idx) {
; CHECK: tex.1d.v4.f32.s32 {%f[[RED:[0-9]+]], %f[[GREEN:[0-9]+]], %f[[BLUE:[0-9]+]], %f[[ALPHA:[0-9]+]]}, [foo_param_0, foo_param_1, {%r{{[0-9]+}}}] ; CHECK: tex.1d.v4.f32.s32 {%f[[RED:[0-9]+]], %f[[GREEN:[0-9]+]], %f[[BLUE:[0-9]+]], %f[[ALPHA:[0-9]+]]}, [foo_param_0, foo_param_1, {%r{{[0-9]+}}}]
%val = tail call { float, float, float, float } @llvm.nvvm.tex.1d.v4f32.s32(i64 %img, i64 %sampler, i32 %idx) %val = tail call { float, float, float, float } @llvm.nvvm.tex.1d.v4f32.s32(i64 %img, i64 %sampler, i32 %idx)
%ret = extractvalue { float, float, float, float } %val, 0 %ret = extractvalue { float, float, float, float } %val, 0
; CHECK: st.f32 [%r{{[0-9]+}}], %f[[RED]] ; CHECK: st.f32 [%rd{{[0-9]+}}], %f[[RED]]
store float %ret, ptr %red store float %ret, ptr %red
ret void ret void
} }
!nvvm.annotations = !{!1, !2, !3} !nvvm.annotations = !{!1, !2, !3}
!1 = !{ptr @foo, !"kernel", i32 1} !1 = !{ptr @foo, !"kernel", i32 1}
!2 = !{ptr @foo, !"rdoimage", i32 0} !2 = !{ptr @foo, !"rdoimage", i32 0}
!3 = !{ptr @foo, !"sampler", i32 1} !3 = !{ptr @foo, !"sampler", i32 1}

llvm/test/CodeGen/NVPTX/texsurf-queries.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s --check-prefix=SM20 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s --check-prefix=SM20
; RUN: llc < %s -march=nvptx -mcpu=sm_30 -verify-machineinstrs | FileCheck %s --check-prefix=SM30 ; RUN: llc < %s -march=nvptx64 -mcpu=sm_30 -verify-machineinstrs | FileCheck %s --check-prefix=SM30
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_30 -verify-machineinstrs | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_30 -verify-machineinstrs | %ptxas-verify %}
target triple = "nvptx-unknown-cuda" target triple = "nvptx-unknown-cuda"
@tex0 = internal addrspace(1) global i64 0, align 8 @tex0 = internal addrspace(1) global i64 0, align 8
@surf0 = internal addrspace(1) global i64 0, align 8 @surf0 = internal addrspace(1) global i64 0, align 8
declare i32 @llvm.nvvm.txq.width(i64) declare i32 @llvm.nvvm.txq.width(i64)
declare i32 @llvm.nvvm.txq.height(i64) declare i32 @llvm.nvvm.txq.height(i64)
▲ Show 20 Lines • Show All 93 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/tuple-literal.ll

  • This file was deleted.
; RUN: llc < %s -march=nvptx -mcpu=sm_20 %if ptxas %{ | %ptxas-verify %}
define ptx_device void @test_function(ptr) {
ret void
}

llvm/test/CodeGen/NVPTX/vaargs.ll

; RUN: llc < %s -O0 -march=nvptx -mattr=+ptx60 -mcpu=sm_30 | FileCheck %s --check-prefixes=CHECK,CHECK32 ; RUN: llc < %s -O0 -march=nvptx -mattr=+ptx60 -mcpu=sm_30 | FileCheck %s --check-prefixes=CHECK,CHECK32
; RUN: llc < %s -O0 -march=nvptx64 -mattr=+ptx60 -mcpu=sm_30 | FileCheck %s --check-prefixes=CHECK,CHECK64 ; RUN: llc < %s -O0 -march=nvptx64 -mattr=+ptx60 -mcpu=sm_30 | FileCheck %s --check-prefixes=CHECK,CHECK64
; RUN: %if ptxas %{ llc < %s -O0 -march=nvptx -mattr=+ptx60 -mcpu=sm_30 | %ptxas-verify %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -O0 -march=nvptx -mattr=+ptx60 -mcpu=sm_30 | %ptxas-verify %}
; RUN: %if ptxas %{ llc < %s -O0 -march=nvptx64 -mattr=+ptx60 -mcpu=sm_30 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -O0 -march=nvptx64 -mattr=+ptx60 -mcpu=sm_30 | %ptxas-verify %}
; CHECK: .address_size [[BITS:32|64]] ; CHECK: .address_size [[BITS:32|64]]
%struct.__va_list_tag = type { ptr, ptr, i32, i32 } %struct.__va_list_tag = type { ptr, ptr, i32, i32 }
@foo_ptr = internal addrspace(1) global ptr @foo, align 8 @foo_ptr = internal addrspace(1) global ptr @foo, align 8
▲ Show 20 Lines • Show All 99 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/vec-param-load.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
target datalayout = "e-p:32:32:32-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 datalayout = "e-p:32:32:32-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"
define <16 x float> @test_v16f32(<16 x float> %a) { define <16 x float> @test_v16f32(<16 x float> %a) {
; CHECK-LABEL: test_v16f32( ; CHECK-LABEL: test_v16f32(
; CHECK-DAG: ld.param.v4.f32 {[[V_12_15:(%f[0-9]+[, ]*){4}]]}, [test_v16f32_param_0+48]; ; CHECK-DAG: ld.param.v4.f32 {[[V_12_15:(%f[0-9]+[, ]*){4}]]}, [test_v16f32_param_0+48];
; CHECK-DAG: ld.param.v4.f32 {[[V_8_11:(%f[0-9]+[, ]*){4}]]}, [test_v16f32_param_0+32]; ; CHECK-DAG: ld.param.v4.f32 {[[V_8_11:(%f[0-9]+[, ]*){4}]]}, [test_v16f32_param_0+32];
; CHECK-DAG: ld.param.v4.f32 {[[V_4_7:(%f[0-9]+[, ]*){4}]]}, [test_v16f32_param_0+16]; ; CHECK-DAG: ld.param.v4.f32 {[[V_4_7:(%f[0-9]+[, ]*){4}]]}, [test_v16f32_param_0+16];
▲ Show 20 Lines • Show All 73 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/vec8.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
target triple = "nvptx-unknown-cuda" target triple = "nvptx-unknown-cuda"
; CHECK: .visible .func foo ; CHECK: .visible .func foo
define void @foo(<8 x i8> %a, ptr %b) { define void @foo(<8 x i8> %a, ptr %b) {
; CHECK-DAG: ld.param.v4.u8 {[[E0:%rs[0-9]+]], [[E1:%rs[0-9]+]], [[E2:%rs[0-9]+]], [[E3:%rs[0-9]+]]}, [foo_param_0] ; CHECK-DAG: ld.param.v4.u8 {[[E0:%rs[0-9]+]], [[E1:%rs[0-9]+]], [[E2:%rs[0-9]+]], [[E3:%rs[0-9]+]]}, [foo_param_0]
; CHECK-DAG: ld.param.v4.u8 {[[E4:%rs[0-9]+]], [[E5:%rs[0-9]+]], [[E6:%rs[0-9]+]], [[E7:%rs[0-9]+]]}, [foo_param_0+4] ; CHECK-DAG: ld.param.v4.u8 {[[E4:%rs[0-9]+]], [[E5:%rs[0-9]+]], [[E6:%rs[0-9]+]], [[E7:%rs[0-9]+]]}, [foo_param_0+4]
; CHECK-DAG: ld.param.u32 %[[B:r[0-9+]]], [foo_param_1] ; CHECK-DAG: ld.param.u64 %[[B:rd[0-9+]]], [foo_param_1]
; CHECK: add.s16 [[T:%rs[0-9+]]], [[E1]], [[E6]]; ; CHECK: add.s16 [[T:%rs[0-9+]]], [[E1]], [[E6]];
; CHECK: st.u8 [%[[B]]], [[T]]; ; CHECK: st.u8 [%[[B]]], [[T]];
%t0 = extractelement <8 x i8> %a, i32 1 %t0 = extractelement <8 x i8> %a, i32 1
%t1 = extractelement <8 x i8> %a, i32 6 %t1 = extractelement <8 x i8> %a, i32 6
%t = add i8 %t0, %t1 %t = add i8 %t0, %t1
store i8 %t, ptr %b store i8 %t, ptr %b
ret void ret void
} }

llvm/test/CodeGen/NVPTX/vector-args.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
define float @foo(<2 x float> %a) { define float @foo(<2 x float> %a) {
; CHECK: .func (.param .b32 func_retval0) foo ; CHECK: .func (.param .b32 func_retval0) foo
; CHECK: .param .align 8 .b8 foo_param_0[8] ; CHECK: .param .align 8 .b8 foo_param_0[8]
; CHECK: ld.param.v2.f32 {%f{{[0-9]+}}, %f{{[0-9]+}}} ; CHECK: ld.param.v2.f32 {%f{{[0-9]+}}, %f{{[0-9]+}}}
%t1 = fmul <2 x float> %a, %a %t1 = fmul <2 x float> %a, %a
%t2 = extractelement <2 x float> %t1, i32 0 %t2 = extractelement <2 x float> %t1, i32 0
%t3 = extractelement <2 x float> %t1, i32 1 %t3 = extractelement <2 x float> %t1, i32 1
Show All 25 Lines

llvm/test/CodeGen/NVPTX/vector-call.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | %ptxas-verify %}
target triple = "nvptx-unknown-cuda" target triple = "nvptx-unknown-cuda"
declare void @bar(<4 x i32>) declare void @bar(<4 x i32>)
; CHECK-LABEL: .func foo( ; CHECK-LABEL: .func foo(
; CHECK-DAG: ld.param.v4.u32 {[[E0:%r[0-9]+]], [[E1:%r[0-9]+]], [[E2:%r[0-9]+]], [[E3:%r[0-9]+]]}, [foo_param_0]; ; CHECK-DAG: ld.param.v4.u32 {[[E0:%r[0-9]+]], [[E1:%r[0-9]+]], [[E2:%r[0-9]+]], [[E3:%r[0-9]+]]}, [foo_param_0];
; CHECK: .param .align 16 .b8 param0[16]; ; CHECK: .param .align 16 .b8 param0[16];
; CHECK-DAG: st.param.v4.b32 [param0+0], {[[E0]], [[E1]], [[E2]], [[E3]]}; ; CHECK-DAG: st.param.v4.b32 [param0+0], {[[E0]], [[E1]], [[E2]], [[E3]]};
; CHECK: call.uni ; CHECK: call.uni
; CHECK: ret; ; CHECK: ret;
define void @foo(<4 x i32> %a) { define void @foo(<4 x i32> %a) {
tail call void @bar(<4 x i32> %a) tail call void @bar(<4 x i32> %a)
ret void ret void
} }
; CHECK-LABEL: .func foo3( ; CHECK-LABEL: .func foo3(
; CHECK-DAG: ld.param.v2.u32 {[[E0:%r[0-9]+]], [[E1:%r[0-9]+]]}, [foo3_param_0]; ; CHECK-DAG: ld.param.v2.u32 {[[E0:%r[0-9]+]], [[E1:%r[0-9]+]]}, [foo3_param_0];
; CHECK-DAG: ld.param.u32 [[E2:%r[0-9]+]], [foo3_param_0+8]; ; CHECK-DAG: ld.param.u32 [[E2:%r[0-9]+]], [foo3_param_0+8];
; CHECK: .param .align 16 .b8 param0[16]; ; CHECK: .param .align 16 .b8 param0[16];
; CHECK-DAG: st.param.v2.b32 [param0+0], {[[E0]], [[E1]]}; ; CHECK-DAG: st.param.v2.b32 [param0+0], {[[E0]], [[E1]]};
; CHECK-DAG: st.param.b32 [param0+8], [[E2]]; ; CHECK-DAG: st.param.b32 [param0+8], [[E2]];
; CHECK: call.uni ; CHECK: call.uni
; CHECK: ret; ; CHECK: ret;
declare void @bar3(<3 x i32>) declare void @bar3(<3 x i32>)
define void @foo3(<3 x i32> %a) { define void @foo3(<3 x i32> %a) {
tail call void @bar3(<3 x i32> %a) tail call void @bar3(<3 x i32> %a)
ret void ret void
} }

llvm/test/CodeGen/NVPTX/vector-compare.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify -m32 %} ; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify -m32 %}
; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; This test makes sure that the result of vector compares are properly ; This test makes sure that the result of vector compares are properly
; scalarized. If codegen fails, then the type legalizer incorrectly ; scalarized. If codegen fails, then the type legalizer incorrectly
; tried to promote <2 x i1> to <2 x i8> and instruction selection failed. ; tried to promote <2 x i1> to <2 x i8> and instruction selection failed.
; CHECK-LABEL: .visible .func foo( ; CHECK-LABEL: .visible .func foo(
define void @foo(ptr %a, ptr %b, ptr %r1, ptr %r2) { define void @foo(ptr %a, ptr %b, ptr %r1, ptr %r2) {
Show All 19 Lines

llvm/test/CodeGen/NVPTX/vector-loads.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; Even though general vector types are not supported in PTX, we can still ; Even though general vector types are not supported in PTX, we can still
; optimize loads/stores with pseudo-vector instructions of the form: ; optimize loads/stores with pseudo-vector instructions of the form:
; ;
; ld.v2.f32 {%f0, %f1}, [%r0] ; ld.v2.f32 {%f0, %f1}, [%r0]
; ;
; which will load two floats at once into scalar registers. ; which will load two floats at once into scalar registers.
▲ Show 20 Lines • Show All 91 Lines • Show Last 20 Lines

llvm/test/CodeGen/NVPTX/vector-select.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 %if ptxas %{ | %ptxas-verify %} ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 %if ptxas %{ | %ptxas-verify %} ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas && !ptxas-12.0 %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %}
; RUN: %if ptxas %{llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; This test makes sure that vector selects are scalarized by the type legalizer. ; This test makes sure that vector selects are scalarized by the type legalizer.
; If not, type legalization will fail. ; If not, type legalization will fail.
; CHECK-LABEL: .visible .func foo(
define void @foo(ptr addrspace(1) %def_a, ptr addrspace(1) %def_b, ptr addrspace(1) %def_c) { define void @foo(ptr addrspace(1) %def_a, ptr addrspace(1) %def_b, ptr addrspace(1) %def_c) {
entry: entry:
; CHECK: ld.global.v2.u32
; CHECK: ld.global.v2.u32
; CHECK: ld.global.v2.u32
%tmp4 = load <2 x i32>, ptr addrspace(1) %def_a %tmp4 = load <2 x i32>, ptr addrspace(1) %def_a
%tmp6 = load <2 x i32>, ptr addrspace(1) %def_c %tmp6 = load <2 x i32>, ptr addrspace(1) %def_c
%tmp8 = load <2 x i32>, ptr addrspace(1) %def_b %tmp8 = load <2 x i32>, ptr addrspace(1) %def_b
; CHECK: setp.gt.s32
; CHECK: setp.gt.s32
%0 = icmp sge <2 x i32> %tmp4, zeroinitializer %0 = icmp sge <2 x i32> %tmp4, zeroinitializer
; CHECK: selp.b32
; CHECK: selp.b32
%cond = select <2 x i1> %0, <2 x i32> %tmp6, <2 x i32> %tmp8 %cond = select <2 x i1> %0, <2 x i32> %tmp6, <2 x i32> %tmp8
; CHECK: st.global.v2.u32
store <2 x i32> %cond, ptr addrspace(1) %def_c store <2 x i32> %cond, ptr addrspace(1) %def_c
ret void ret void
} }

llvm/test/CodeGen/NVPTX/vector-stores.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; CHECK: .visible .func foo1 ; CHECK: .visible .func foo1
; CHECK: st.v2.f32 ; CHECK: st.v2.f32
define void @foo1(<2 x float> %val, ptr %ptr) { define void @foo1(<2 x float> %val, ptr %ptr) {
store <2 x float> %val, ptr %ptr store <2 x float> %val, ptr %ptr
ret void ret void
} }
Show All 21 Lines

llvm/test/CodeGen/NVPTX/weak-global.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; CHECK: .weak .global .align 4 .u32 g ; CHECK: .weak .global .align 4 .u32 g
@g = common addrspace(1) global i32 zeroinitializer @g = common addrspace(1) global i32 zeroinitializer
define i32 @func0() { define i32 @func0() {
%val = load i32, ptr addrspace(1) @g %val = load i32, ptr addrspace(1) @g
ret i32 %val ret i32 %val
} }

llvm/test/CodeGen/NVPTX/weak-linkage.ll

; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s ; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_20 | %ptxas-verify %} ; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
; CHECK: // .weak foo ; CHECK: // .weak foo
; CHECK: .weak .func foo ; CHECK: .weak .func foo
define weak void @foo() { define weak void @foo() {
ret void ret void
} }
; CHECK: // .weak baz ; CHECK: // .weak baz
Show All 9 Lines

llvm/test/lit.cfg.py

Show First 20 Lines • Show All 299 Lines • ▼ Show 20 Lines if version:
(11, 2), (11, 2),
(11, 3), (11, 3),
(11, 4), (11, 4),
(11, 5), (11, 5),
(11, 6), (11, 6),
(11, 7), (11, 7),
(11, 8), (11, 8),
(12, 0), (12, 0),
(12, 1),
] ]
def version_int(ver): def version_int(ver):
return ver[0] * 100 + ver[1] return ver[0] * 100 + ver[1]
# ignore ptxas if its version is below the minimum supported # ignore ptxas if its version is below the minimum supported
# version # version
min_version = ptxas_known_versions[0] min_version = ptxas_known_versions[0]
▲ Show 20 Lines • Show All 331 Lines • Show Last 20 Lines
llvm/test/CodeGen/NVPTX/tex-read.ll