[NVPTX] add support for initializing fp16 arrays.

[tra]

Summary

Previously HalfTy was not handled which would either trigger an assertion,
or result in array initialized with garbage.

[jlebar]

llvm/test/CodeGen/NVPTX/half.ll

6

Should we also test a [N x half] array, in addition to this nested array?

[tra]

Simplified the test case.

llvm/test/CodeGen/NVPTX/half.ll

6

Being nested is incidental -- I've just copy/pasted it from the test case that triggered the problem.
I've simplified the test to use a 1D array.

llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp

if (ETy == Type::getInt8Ty(CPV->getContext())) {
           break;
         }
       }
       llvm_unreachable("unsupported integer const type");
     } else
       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);
     } else if (Ty == Type::getDoubleTy(CPV->getContext())) {
       double float64 = CFP->getValueAPF().convertToDouble();
       ConvertDoubleToBytes(ptr, float64);
       aggBuffer->addBytes(ptr, 8, Bytes);
     } else {

llvm/test/CodeGen/NVPTX/half.ll

 ; 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 [2 x [2 x half]]
+                [[2 x half] [half 0xH0201, half 0xH0403],
+                 [2 x half] [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]+}}]
 ; CHECK: st.global.b16 [{{%r[0-9]+}}], [[TMP]]
   %val = load half, half addrspace(1)* %in
   store half %val, half addrspace(1) * %out
   ret void
 }