[llvm] r329463 - [NVPTX] add support for initializing fp16 arrays.
Artem Belevich via llvm-commits
llvm-commits at lists.llvm.org
Fri Apr 6 15:25:08 PDT 2018
Author: tra
Date: Fri Apr 6 15:25:08 2018
New Revision: 329463
URL: http://llvm.org/viewvc/llvm-project?rev=329463&view=rev
Log:
[NVPTX] add support for initializing fp16 arrays.
Previously HalfTy was not handled which would either trigger an assertion,
or result in array initialized with garbage.
Differential Revision: https://reviews.llvm.org/D45391
Modified:
llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
llvm/trunk/test/CodeGen/NVPTX/half.ll
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp?rev=329463&r1=329462&r2=329463&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp Fri Apr 6 15:25:08 2018
@@ -1945,11 +1945,17 @@ void NVPTXAsmPrinter::bufferLEByte(const
llvm_unreachable("unsupported integer const type");
break;
}
+ case Type::HalfTyID:
case Type::FloatTyID:
case Type::DoubleTyID: {
const ConstantFP *CFP = dyn_cast<ConstantFP>(CPV);
Type *Ty = CFP->getType();
- if (Ty == Type::getFloatTy(CPV->getContext())) {
+ if (Ty == Type::getHalfTy(CPV->getContext())) {
+ APInt API = CFP->getValueAPF().bitcastToAPInt();
+ uint16_t float16 = API.getLoBits(16).getZExtValue();
+ ConvertIntToBytes<>(ptr, float16);
+ aggBuffer->addBytes(ptr, 2, Bytes);
+ } else if (Ty == Type::getFloatTy(CPV->getContext())) {
float float32 = (float) CFP->getValueAPF().convertToFloat();
ConvertFloatToBytes(ptr, float32);
aggBuffer->addBytes(ptr, 4, Bytes);
Modified: llvm/trunk/test/CodeGen/NVPTX/half.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/NVPTX/half.ll?rev=329463&r1=329462&r2=329463&view=diff
==============================================================================
--- llvm/trunk/test/CodeGen/NVPTX/half.ll (original)
+++ llvm/trunk/test/CodeGen/NVPTX/half.ll Fri Apr 6 15:25:08 2018
@@ -1,5 +1,9 @@
; RUN: llc < %s -march=nvptx | FileCheck %s
+; CHECK: .b8 half_array[8] = {1, 2, 3, 4, 5, 6, 7, 8};
+@"half_array" = addrspace(1) constant [4 x half]
+ [half 0xH0201, half 0xH0403, half 0xH0605, half 0xH0807]
+
define void @test_load_store(half addrspace(1)* %in, half addrspace(1)* %out) {
; CHECK-LABEL: @test_load_store
; CHECK: ld.global.b16 [[TMP:%h[0-9]+]], [{{%r[0-9]+}}]
More information about the llvm-commits
mailing list