[PATCH] D45391: [NVPTX] add support for initializing fp16 arrays.

[Artem Belevich via Phabricator via llvm-commits]

This revision was automatically updated to reflect the committed changes.
Closed by commit rL329463: [NVPTX] add support for initializing fp16 arrays. (authored by tra, committed by ).

Changed prior to commit:
  https://reviews.llvm.org/D45391?vs=141434&id=141441#toc

Repository:
  rL LLVM

https://reviews.llvm.org/D45391

Files:
  llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
  llvm/trunk/test/CodeGen/NVPTX/half.ll


Index: llvm/trunk/test/CodeGen/NVPTX/half.ll
===================================================================
--- llvm/trunk/test/CodeGen/NVPTX/half.ll
+++ llvm/trunk/test/CodeGen/NVPTX/half.ll
@@ -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]+}}]
Index: llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
===================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
+++ llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
@@ -1945,11 +1945,17 @@
       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);


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D45391.141441.patch
Type: text/x-patch
Size: 1627 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20180406/34ea2d52/attachment.bin>