1fdbe69849
Lower `fadd`, `fsub`, `fmul`, and `fma` to f32x2 variants introduced in PTX 8.6 for sm_100+. Adds a new register class for v2f32 as a b64 register in PTX. This causes other vector operations like loads and stores to lower as .b64 instead of .v2.b32 as appropriate. Also update test cases to use the autogenerator.
50 lines
1.5 KiB
LLVM
50 lines
1.5 KiB
LLVM
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
|
|
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_35 -verify-machineinstrs | FileCheck %s
|
|
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_35 | %ptxas-verify %}
|
|
|
|
; Check load from constant global variables. These loads should be
|
|
; ld.global.nc (aka ldg).
|
|
|
|
@gv_float = external constant float
|
|
@gv_float2 = external constant <2 x float>
|
|
@gv_float4 = external constant <4 x float>
|
|
|
|
define float @test_gv_float() {
|
|
; CHECK-LABEL: test_gv_float(
|
|
; CHECK: {
|
|
; CHECK-NEXT: .reg .b32 %r<2>;
|
|
; CHECK-EMPTY:
|
|
; CHECK-NEXT: // %bb.0:
|
|
; CHECK-NEXT: ld.global.nc.b32 %r1, [gv_float];
|
|
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
|
|
; CHECK-NEXT: ret;
|
|
%v = load float, ptr @gv_float
|
|
ret float %v
|
|
}
|
|
|
|
define <2 x float> @test_gv_float2() {
|
|
; CHECK-LABEL: test_gv_float2(
|
|
; CHECK: {
|
|
; CHECK-NEXT: .reg .b64 %rd<2>;
|
|
; CHECK-EMPTY:
|
|
; CHECK-NEXT: // %bb.0:
|
|
; CHECK-NEXT: ld.global.nc.b64 %rd1, [gv_float2];
|
|
; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
|
|
; CHECK-NEXT: ret;
|
|
%v = load <2 x float>, ptr @gv_float2
|
|
ret <2 x float> %v
|
|
}
|
|
|
|
define <4 x float> @test_gv_float4() {
|
|
; CHECK-LABEL: test_gv_float4(
|
|
; CHECK: {
|
|
; CHECK-NEXT: .reg .b64 %rd<3>;
|
|
; CHECK-EMPTY:
|
|
; CHECK-NEXT: // %bb.0:
|
|
; CHECK-NEXT: ld.global.nc.v2.b64 {%rd1, %rd2}, [gv_float4];
|
|
; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd1, %rd2};
|
|
; CHECK-NEXT: ret;
|
|
%v = load <4 x float>, ptr @gv_float4
|
|
ret <4 x float> %v
|
|
}
|