[NVPTX] Expose LDU builtins

[jchlanda]

Summary

Also check if native half types are supported to give more descriptive error message, without it clang only reports incorrect intrinsic return type.

[tra]

Nice. Thank you!

clang/lib/CodeGen/CGBuiltin.cpp

18104

I'd add a comment describing a meaning of the fields in the returned pair.

18116

Can we use the standard StringRef Name = getContext().BuiltinInfo.getName(BuiltinID); to figure out the builtin name?

18261–18263

Nit: this would be a bit more readable:

std::string BuiltinName{HalfSupport.second};
CGM.Error(E->getExprLoc(),  BuiltinName + " requires native half type support.")`

You may also consider changing returned BuiltinName to be std::string, so you would not need an explicit temp var.

clang/test/CodeGen/builtins-nvptx-native-half-type-err.c

3

I think we could've done it all in one run if we were to do both ldg and ldu in one function.

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

33

Should alignment be 2 ?

60

Hmm. I wonder why we end up with u64 here and not b64. Not that it matters in this case, but it is a discrepancy vs. f16.

[jchlanda]

Address PR comments

clang/lib/CodeGen/CGBuiltin.cpp

18116

Ha, had a feeling it would exist, couldn't find it. Thanks.

18261–18263

Done the std::string return, thanks.

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

60

That is copy/paste sloppiness on my part, sorry. I've updated the test to check generated PTX, not just the labels, and fixed the values.

It generates correct kinds of loads, based on the type, the only discrepancy is that it doesn't distinguish between signed and unsigned loads, always choosing the unsigned variant. I think that's by design, at ISel there is a check if the load needs to be extended and a correct CVT instruction will be added.

[tra]

clang/lib/CodeGen/CGBuiltin.cpp

18259–18263

I think we can simplify it all further.

auto HasHalfSupport = [&](unsigned BuiltinID) {
    auto &Context = getContext();
    return Context.getLangOpts().NativeHalfType || !Context.getTargetInfo().useFP16ConversionIntrinsics();
}
...

if (!HasHalfSupport(BuiltinID)) {
      CGM.Error(E->getExprLoc(),   getContext().BuiltinInfo.getName(BuiltinID) + " requires native half type support.");

[jchlanda]

Simplify the check for half tys support.

[jchlanda]

clang/lib/CodeGen/CGBuiltin.cpp

18259–18263

Done, although we need a string append there, StringRef and Twine would not work.

[tra]

clang/lib/CodeGen/CGBuiltin.cpp

18259–18263

We don't really need append().
CGM.Error(E->getExprLoc(), getContext().BuiltinInfo.getName(BuiltinID).str() + " requires native half type support."); works.

[jchlanda]

append -> +

[jchlanda]

clang/lib/CodeGen/CGBuiltin.cpp

18259–18263

Done.

[jchlanda]

@tra is there anything else I should do for this patch? Thank you.

[tra]

LGTM with a test nit.

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

30

Nit. Function names usually checked with CHECK-LABEL.

[jchlanda]

Use CHECK-LABEL.

clang/include/clang/Basic/BuiltinsNVPTX.def

 BUILTIN(__nvvm_atom_cas_gen_ll, "LLiLLiD*LLiLLi", "n")
 TARGET_BUILTIN(__nvvm_atom_cta_cas_gen_ll, "LLiLLiD*LLiLLi", "n", SM_60)
 TARGET_BUILTIN(__nvvm_atom_sys_cas_gen_ll, "LLiLLiD*LLiLLi", "n", SM_60)
 
 // Compiler Error Warn
 BUILTIN(__nvvm_compiler_error, "vcC*4", "n")
 BUILTIN(__nvvm_compiler_warn, "vcC*4", "n")
 
-// __ldg.  This is not implemented as a builtin by nvcc.
+BUILTIN(__nvvm_ldu_c, "ccC*", "")
+BUILTIN(__nvvm_ldu_s, "ssC*", "")
+BUILTIN(__nvvm_ldu_i, "iiC*", "")
+BUILTIN(__nvvm_ldu_l, "LiLiC*", "")
+BUILTIN(__nvvm_ldu_ll, "LLiLLiC*", "")
+
+BUILTIN(__nvvm_ldu_uc, "UcUcC*", "")
+BUILTIN(__nvvm_ldu_us, "UsUsC*", "")
+BUILTIN(__nvvm_ldu_ui, "UiUiC*", "")
+BUILTIN(__nvvm_ldu_ul, "ULiULiC*", "")
+BUILTIN(__nvvm_ldu_ull, "ULLiULLiC*", "")
+
+BUILTIN(__nvvm_ldu_h, "hhC*", "")
+BUILTIN(__nvvm_ldu_f, "ffC*", "")
+BUILTIN(__nvvm_ldu_d, "ddC*", "")
+
+BUILTIN(__nvvm_ldu_c2, "E2cE2cC*", "")
+BUILTIN(__nvvm_ldu_c4, "E4cE4cC*", "")
+BUILTIN(__nvvm_ldu_s2, "E2sE2sC*", "")
+BUILTIN(__nvvm_ldu_s4, "E4sE4sC*", "")
+BUILTIN(__nvvm_ldu_i2, "E2iE2iC*", "")
+BUILTIN(__nvvm_ldu_i4, "E4iE4iC*", "")
+BUILTIN(__nvvm_ldu_ll2, "E2LLiE2LLiC*", "")
+
+BUILTIN(__nvvm_ldu_uc2, "E2UcE2UcC*", "")
+BUILTIN(__nvvm_ldu_uc4, "E4UcE4UcC*", "")
+BUILTIN(__nvvm_ldu_us2, "E2UsE2UsC*", "")
+BUILTIN(__nvvm_ldu_us4, "E4UsE4UsC*", "")
+BUILTIN(__nvvm_ldu_ui2, "E2UiE2UiC*", "")
+BUILTIN(__nvvm_ldu_ui4, "E4UiE4UiC*", "")
+BUILTIN(__nvvm_ldu_ull2, "E2ULLiE2ULLiC*", "")
+
+BUILTIN(__nvvm_ldu_h2, "E2hE2hC*", "")
+BUILTIN(__nvvm_ldu_f2, "E2fE2fC*", "")
+BUILTIN(__nvvm_ldu_f4, "E4fE4fC*", "")
+BUILTIN(__nvvm_ldu_d2, "E2dE2dC*", "")
+
 BUILTIN(__nvvm_ldg_c, "ccC*", "")
 BUILTIN(__nvvm_ldg_s, "ssC*", "")
 BUILTIN(__nvvm_ldg_i, "iiC*", "")
 BUILTIN(__nvvm_ldg_l, "LiLiC*", "")
 BUILTIN(__nvvm_ldg_ll, "LLiLLiC*", "")
 
 BUILTIN(__nvvm_ldg_uc, "UcUcC*", "")
 BUILTIN(__nvvm_ldg_us, "UsUsC*", "")

clang/lib/CodeGen/CGBuiltin.cpp

 #undef MMA_VARIANTS_B1_AND
 #undef MMA_VARIANTS_B1_XOR
 }
 
 } // namespace
 
 Value *
 CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) {
-  auto MakeLdg = [&](unsigned IntrinsicID) {
+  auto HasHalfSupport = [&](unsigned BuiltinID) {
+    auto &Context = getContext();
+    const bool Res = Context.getLangOpts().NativeHalfType ||
+                     !Context.getTargetInfo().useFP16ConversionIntrinsics();
+    if (Res)
+      return std::make_pair(true, "");
+
+    const char *BuiltinName;
+    switch (BuiltinID) {
+    default:
+      llvm_unreachable("Unknown builtin.");
+    case NVPTX::BI__nvvm_ldu_h:
+      BuiltinName = "__nvvm_ldu_h";
+      break;
+    case NVPTX::BI__nvvm_ldu_h2:
+      BuiltinName = "__nvvm_ldu_h2";
+      break;
+    case NVPTX::BI__nvvm_ldg_h:
+      BuiltinName = "__nvvm_ldg_h";
+      break;
+    case NVPTX::BI__nvvm_ldg_h2:
+      BuiltinName = "__nvvm_ldg_h2";
+      break;
+    }
+    return std::make_pair(false, BuiltinName);
+  };
+  auto MakeLdgLdu = [&](unsigned IntrinsicID) {
     Value *Ptr = EmitScalarExpr(E->getArg(0));
     QualType ArgType = E->getArg(0)->getType();
     clang::CharUnits Align = CGM.getNaturalPointeeTypeAlignment(ArgType);
     llvm::Type *ElemTy = ConvertTypeForMem(ArgType->getPointeeType());
     return Builder.CreateCall(
         CGM.getIntrinsic(IntrinsicID, {ElemTy, Ptr->getType()}),
         {Ptr, ConstantInt::get(Builder.getInt32Ty(), Align.getQuantity())});
   };
CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID, const CallExpr *E) {
   case NVPTX::BI__nvvm_ldg_ui2:
   case NVPTX::BI__nvvm_ldg_ui4:
   case NVPTX::BI__nvvm_ldg_ul:
   case NVPTX::BI__nvvm_ldg_ull:
   case NVPTX::BI__nvvm_ldg_ull2:
     // PTX Interoperability section 2.2: "For a vector with an even number of
     // elements, its alignment is set to number of elements times the alignment
     // of its member: n*alignof(t)."
-    return MakeLdg(Intrinsic::nvvm_ldg_global_i);
+    return MakeLdgLdu(Intrinsic::nvvm_ldg_global_i);
   case NVPTX::BI__nvvm_ldg_h:
+  case NVPTX::BI__nvvm_ldg_h2: {
+    auto HalfSupport = HasHalfSupport(BuiltinID);
+    if (!HalfSupport.first) {
+      std::string ErrMsg{HalfSupport.second};
+      CGM.Error(E->getExprLoc(),
+                ErrMsg.append(" requires native half type support.").c_str());
+      return nullptr;
+    }
+  }
+    [[fallthrough]];
   case NVPTX::BI__nvvm_ldg_f:
-  case NVPTX::BI__nvvm_ldg_h2:
   case NVPTX::BI__nvvm_ldg_f2:
   case NVPTX::BI__nvvm_ldg_f4:
   case NVPTX::BI__nvvm_ldg_d:
   case NVPTX::BI__nvvm_ldg_d2:
-    return MakeLdg(Intrinsic::nvvm_ldg_global_f);
+    return MakeLdgLdu(Intrinsic::nvvm_ldg_global_f);
+
+  case NVPTX::BI__nvvm_ldu_c:
+  case NVPTX::BI__nvvm_ldu_c2:
+  case NVPTX::BI__nvvm_ldu_c4:
+  case NVPTX::BI__nvvm_ldu_s:
+  case NVPTX::BI__nvvm_ldu_s2:
+  case NVPTX::BI__nvvm_ldu_s4:
+  case NVPTX::BI__nvvm_ldu_i:
+  case NVPTX::BI__nvvm_ldu_i2:
+  case NVPTX::BI__nvvm_ldu_i4:
+  case NVPTX::BI__nvvm_ldu_l:
+  case NVPTX::BI__nvvm_ldu_ll:
+  case NVPTX::BI__nvvm_ldu_ll2:
+  case NVPTX::BI__nvvm_ldu_uc:
+  case NVPTX::BI__nvvm_ldu_uc2:
+  case NVPTX::BI__nvvm_ldu_uc4:
+  case NVPTX::BI__nvvm_ldu_us:
+  case NVPTX::BI__nvvm_ldu_us2:
+  case NVPTX::BI__nvvm_ldu_us4:
+  case NVPTX::BI__nvvm_ldu_ui:
+  case NVPTX::BI__nvvm_ldu_ui2:
+  case NVPTX::BI__nvvm_ldu_ui4:
+  case NVPTX::BI__nvvm_ldu_ul:
+  case NVPTX::BI__nvvm_ldu_ull:
+  case NVPTX::BI__nvvm_ldu_ull2:
+    return MakeLdgLdu(Intrinsic::nvvm_ldu_global_i);
+  case NVPTX::BI__nvvm_ldu_h:
+  case NVPTX::BI__nvvm_ldu_h2: {
+    auto HalfSupport = HasHalfSupport(BuiltinID);
+    if (!HalfSupport.first) {
+      std::string ErrMsg{HalfSupport.second};
+      CGM.Error(E->getExprLoc(),
+                ErrMsg.append(" requires native half type support.").c_str());
+      return nullptr;
+    }
+  }
+    [[fallthrough]];
+  case NVPTX::BI__nvvm_ldu_f:
+  case NVPTX::BI__nvvm_ldu_f2:
+  case NVPTX::BI__nvvm_ldu_f4:
+  case NVPTX::BI__nvvm_ldu_d:
+  case NVPTX::BI__nvvm_ldu_d2:
+    return MakeLdgLdu(Intrinsic::nvvm_ldu_global_f);
 
   case NVPTX::BI__nvvm_atom_cta_add_gen_i:
   case NVPTX::BI__nvvm_atom_cta_add_gen_l:
   case NVPTX::BI__nvvm_atom_cta_add_gen_ll:
     return MakeScopedAtomic(Intrinsic::nvvm_atomic_add_gen_i_cta);
   case NVPTX::BI__nvvm_atom_sys_add_gen_i:
   case NVPTX::BI__nvvm_atom_sys_add_gen_l:
   case NVPTX::BI__nvvm_atom_sys_add_gen_ll:

clang/test/CodeGen/builtins-nvptx-native-half-type-err.c

+// REQUIRES: nvptx-registered-target
+//
+// RUN: not %clang_cc1 -DLDG -fsyntax-only -ffp-contract=off -triple nvptx-unknown-unknown -target-cpu \
+// RUN:   sm_75 -target-feature +ptx70 -fcuda-is-device -x cuda -emit-llvm -o - %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=CHECK-ERROR-LDG %s
+
+// RUN: not %clang_cc1 -DLDU -fsyntax-only -ffp-contract=off -triple nvptx-unknown-unknown -target-cpu \
+// RUN:   sm_75 -target-feature +ptx70 -fcuda-is-device -x cuda -emit-llvm -o - %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=CHECK-ERROR-LDU %s
+#define __device__ __attribute__((device))
+typedef __fp16 __fp16v2 __attribute__((ext_vector_type(2)));
+
+#ifdef LDG
+__device__ void nvvm_ldg_native_half_types(const void *p) {
+  __nvvm_ldg_h((const __fp16 *)p);
+  __nvvm_ldg_h2((const __fp16v2 *)p);
+}
+#endif // LDG
+
+#ifdef LDU
+__device__ void nvvm_ldu_native_half_types(const void *p) {
+  __nvvm_ldu_h((const __fp16 *)p);
+  __nvvm_ldu_h2((const __fp16v2 *)p);
+}
+#endif // LDU
+
+// CHECK-ERROR-LDG: error: __nvvm_ldg_h requires native half type support.
+// CHECK-ERROR-LDG: error: __nvvm_ldg_h2 requires native half type support.
+// CHECK-ERROR-LDU: error: __nvvm_ldu_h requires native half type support.
+// CHECK-ERROR-LDU: error: __nvvm_ldu_h2 requires native half type support.

clang/test/CodeGen/builtins-nvptx-native-half-type.c

#if __CUDA_ARCH__ >= 860
   // CHECK_PTX72_SM86: call <2 x half> @llvm.nvvm.fmax.nan.xorsign.abs.f16x2
   __nvvm_fmax_nan_xorsign_abs_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});
   // CHECK_PTX72_SM86: call <2 x half> @llvm.nvvm.fmax.ftz.nan.xorsign.abs.f16x2
   __nvvm_fmax_ftz_nan_xorsign_abs_f16x2({0.1f16, 0.7f16}, {0.1f16, 0.7f16});
 #endif
   // CHECK: ret void
 }
 
+typedef __fp16 __fp16v2 __attribute__((ext_vector_type(2)));
+
 // CHECK-LABEL: nvvm_ldg_native_half_types
 __device__ void nvvm_ldg_native_half_types(const void *p) {
   // CHECK: call half @llvm.nvvm.ldg.global.f.f16.p0
   __nvvm_ldg_h((const __fp16 *)p);
-  typedef __fp16 __fp16v2 __attribute__((ext_vector_type(2)));
   // CHECK: call <2 x half> @llvm.nvvm.ldg.global.f.v2f16.p0
   __nvvm_ldg_h2((const __fp16v2 *)p);
 }
+
+// CHECK-LABEL: nvvm_ldu_native_half_types
+__device__ void nvvm_ldu_native_half_types(const void *p) {
+  // CHECK: call half @llvm.nvvm.ldu.global.f.f16.p0
+  __nvvm_ldu_h((const __fp16 *)p);
+  // CHECK: call <2 x half> @llvm.nvvm.ldu.global.f.v2f16.p0
+  __nvvm_ldu_h2((const __fp16v2 *)p);
+}

clang/test/CodeGen/builtins-nvptx.c

__device__ void nvvm_ldg(const void *p) {
   typedef float float4 __attribute__((ext_vector_type(4)));
   __nvvm_ldg_f4((const float4 *)p);
 
   // CHECK: call <2 x double> @llvm.nvvm.ldg.global.f.v2f64.p0(ptr {{%[0-9]+}}, i32 16)
   typedef double double2 __attribute__((ext_vector_type(2)));
   __nvvm_ldg_d2((const double2 *)p);
 }
 
+// CHECK-LABEL: nvvm_ldu
+__device__ void nvvm_ldu(const void *p) {
+  // CHECK: call i8 @llvm.nvvm.ldu.global.i.i8.p0(ptr {{%[0-9]+}}, i32 1)
+  // CHECK: call i8 @llvm.nvvm.ldu.global.i.i8.p0(ptr {{%[0-9]+}}, i32 1)
+  __nvvm_ldu_c((const char *)p);
+  __nvvm_ldu_uc((const unsigned char *)p);
+
+  // CHECK: call i16 @llvm.nvvm.ldu.global.i.i16.p0(ptr {{%[0-9]+}}, i32 2)
+  // CHECK: call i16 @llvm.nvvm.ldu.global.i.i16.p0(ptr {{%[0-9]+}}, i32 2)
+  __nvvm_ldu_s((const short *)p);
+  __nvvm_ldu_us((const unsigned short *)p);
+
+  // CHECK: call i32 @llvm.nvvm.ldu.global.i.i32.p0(ptr {{%[0-9]+}}, i32 4)
+  // CHECK: call i32 @llvm.nvvm.ldu.global.i.i32.p0(ptr {{%[0-9]+}}, i32 4)
+  __nvvm_ldu_i((const int *)p);
+  __nvvm_ldu_ui((const unsigned int *)p);
+
+  // LP32: call i32 @llvm.nvvm.ldu.global.i.i32.p0(ptr {{%[0-9]+}}, i32 4)
+  // LP32: call i32 @llvm.nvvm.ldu.global.i.i32.p0(ptr {{%[0-9]+}}, i32 4)
+  // LP64: call i64 @llvm.nvvm.ldu.global.i.i64.p0(ptr {{%[0-9]+}}, i32 8)
+  // LP64: call i64 @llvm.nvvm.ldu.global.i.i64.p0(ptr {{%[0-9]+}}, i32 8)
+  __nvvm_ldu_l((const long *)p);
+  __nvvm_ldu_ul((const unsigned long *)p);
+
+  // CHECK: call float @llvm.nvvm.ldu.global.f.f32.p0(ptr {{%[0-9]+}}, i32 4)
+  __nvvm_ldu_f((const float *)p);
+  // CHECK: call double @llvm.nvvm.ldu.global.f.f64.p0(ptr {{%[0-9]+}}, i32 8)
+  __nvvm_ldu_d((const double *)p);
+
+  // CHECK: call <2 x i8> @llvm.nvvm.ldu.global.i.v2i8.p0(ptr {{%[0-9]+}}, i32 2)
+  // CHECK: call <2 x i8> @llvm.nvvm.ldu.global.i.v2i8.p0(ptr {{%[0-9]+}}, i32 2)
+  typedef char char2 __attribute__((ext_vector_type(2)));
+  typedef unsigned char uchar2 __attribute__((ext_vector_type(2)));
+  __nvvm_ldu_c2((const char2 *)p);
+  __nvvm_ldu_uc2((const uchar2 *)p);
+
+  // CHECK: call <4 x i8> @llvm.nvvm.ldu.global.i.v4i8.p0(ptr {{%[0-9]+}}, i32 4)
+  // CHECK: call <4 x i8> @llvm.nvvm.ldu.global.i.v4i8.p0(ptr {{%[0-9]+}}, i32 4)
+  typedef char char4 __attribute__((ext_vector_type(4)));
+  typedef unsigned char uchar4 __attribute__((ext_vector_type(4)));
+  __nvvm_ldu_c4((const char4 *)p);
+  __nvvm_ldu_uc4((const uchar4 *)p);
+
+  // CHECK: call <2 x i16> @llvm.nvvm.ldu.global.i.v2i16.p0(ptr {{%[0-9]+}}, i32 4)
+  // CHECK: call <2 x i16> @llvm.nvvm.ldu.global.i.v2i16.p0(ptr {{%[0-9]+}}, i32 4)
+  typedef short short2 __attribute__((ext_vector_type(2)));
+  typedef unsigned short ushort2 __attribute__((ext_vector_type(2)));
+  __nvvm_ldu_s2((const short2 *)p);
+  __nvvm_ldu_us2((const ushort2 *)p);
+
+  // CHECK: call <4 x i16> @llvm.nvvm.ldu.global.i.v4i16.p0(ptr {{%[0-9]+}}, i32 8)
+  // CHECK: call <4 x i16> @llvm.nvvm.ldu.global.i.v4i16.p0(ptr {{%[0-9]+}}, i32 8)
+  typedef short short4 __attribute__((ext_vector_type(4)));
+  typedef unsigned short ushort4 __attribute__((ext_vector_type(4)));
+  __nvvm_ldu_s4((const short4 *)p);
+  __nvvm_ldu_us4((const ushort4 *)p);
+
+  // CHECK: call <2 x i32> @llvm.nvvm.ldu.global.i.v2i32.p0(ptr {{%[0-9]+}}, i32 8)
+  // CHECK: call <2 x i32> @llvm.nvvm.ldu.global.i.v2i32.p0(ptr {{%[0-9]+}}, i32 8)
+  typedef int int2 __attribute__((ext_vector_type(2)));
+  typedef unsigned int uint2 __attribute__((ext_vector_type(2)));
+  __nvvm_ldu_i2((const int2 *)p);
+  __nvvm_ldu_ui2((const uint2 *)p);
+
+  // CHECK: call <4 x i32> @llvm.nvvm.ldu.global.i.v4i32.p0(ptr {{%[0-9]+}}, i32 16)
+  // CHECK: call <4 x i32> @llvm.nvvm.ldu.global.i.v4i32.p0(ptr {{%[0-9]+}}, i32 16)
+  typedef int int4 __attribute__((ext_vector_type(4)));
+  typedef unsigned int uint4 __attribute__((ext_vector_type(4)));
+  __nvvm_ldu_i4((const int4 *)p);
+  __nvvm_ldu_ui4((const uint4 *)p);
+
+  // CHECK: call <2 x i64> @llvm.nvvm.ldu.global.i.v2i64.p0(ptr {{%[0-9]+}}, i32 16)
+  // CHECK: call <2 x i64> @llvm.nvvm.ldu.global.i.v2i64.p0(ptr {{%[0-9]+}}, i32 16)
+  typedef long long longlong2 __attribute__((ext_vector_type(2)));
+  typedef unsigned long long ulonglong2 __attribute__((ext_vector_type(2)));
+  __nvvm_ldu_ll2((const longlong2 *)p);
+  __nvvm_ldu_ull2((const ulonglong2 *)p);
+
+  // CHECK: call <2 x float> @llvm.nvvm.ldu.global.f.v2f32.p0(ptr {{%[0-9]+}}, i32 8)
+  typedef float float2 __attribute__((ext_vector_type(2)));
+  __nvvm_ldu_f2((const float2 *)p);
+
+  // CHECK: call <4 x float> @llvm.nvvm.ldu.global.f.v4f32.p0(ptr {{%[0-9]+}}, i32 16)
+  typedef float float4 __attribute__((ext_vector_type(4)));
+  __nvvm_ldu_f4((const float4 *)p);
+
+  // CHECK: call <2 x double> @llvm.nvvm.ldu.global.f.v2f64.p0(ptr {{%[0-9]+}}, i32 16)
+  typedef double double2 __attribute__((ext_vector_type(2)));
+  __nvvm_ldu_d2((const double2 *)p);
+}
+
 // CHECK-LABEL: nvvm_shfl
 __device__ void nvvm_shfl(int i, float f, int a, int b) {
   // CHECK: call i32 @llvm.nvvm.shfl.down.i32(i32
   __nvvm_shfl_down_i32(i, a, b);
   // CHECK: call float @llvm.nvvm.shfl.down.f32(float
   __nvvm_shfl_down_f32(f, a, b);
   // CHECK: call i32 @llvm.nvvm.shfl.up.i32(i32
   __nvvm_shfl_up_i32(i, a, b);

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

 ; RUN: llc < %s -march=nvptx -mcpu=sm_32 | FileCheck %s
 ; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_32 | %ptxas-verify %}
 
 
 declare i8 @llvm.nvvm.ldu.global.i.i8.p1(ptr addrspace(1) %ptr, i32 %align)
+declare i16 @llvm.nvvm.ldu.global.i.i16.p1(ptr addrspace(1) %ptr, i32 %align)
 declare i32 @llvm.nvvm.ldu.global.i.i32.p1(ptr addrspace(1) %ptr, i32 %align)
+declare i64 @llvm.nvvm.ldu.global.i.i64.p1(ptr addrspace(1) %ptr, i32 %align)
+declare float @llvm.nvvm.ldu.global.f.f32.p1(ptr addrspace(1) %ptr, i32 %align)
+declare double @llvm.nvvm.ldu.global.f.f64.p1(ptr addrspace(1) %ptr, i32 %align)
+declare half @llvm.nvvm.ldu.global.f.f16.p1(ptr addrspace(1) %ptr, i32 %align)
+declare <2 x half> @llvm.nvvm.ldu.global.f.v2f16.p1(ptr addrspace(1) %ptr, i32 %align)
 
 declare i8 @llvm.nvvm.ldg.global.i.i8.p1(ptr addrspace(1) %ptr, i32 %align)
 declare i16 @llvm.nvvm.ldg.global.i.i16.p1(ptr addrspace(1) %ptr, i32 %align)
 declare i32 @llvm.nvvm.ldg.global.i.i32.p1(ptr addrspace(1) %ptr, i32 %align)
 declare i64 @llvm.nvvm.ldg.global.i.i64.p1(ptr addrspace(1) %ptr, i32 %align)
 declare float @llvm.nvvm.ldg.global.f.f32.p1(ptr addrspace(1) %ptr, i32 %align)
 declare double @llvm.nvvm.ldg.global.f.f64.p1(ptr addrspace(1) %ptr, i32 %align)
 declare half @llvm.nvvm.ldg.global.f.f16.p1(ptr addrspace(1) %ptr, i32 %align)
 declare <2 x half> @llvm.nvvm.ldg.global.f.v2f16.p1(ptr addrspace(1) %ptr, i32 %align)
 
 ; CHECK: test_ldu_i8
 define i8 @test_ldu_i8(ptr addrspace(1) %ptr) {
 ; ldu.global.u8
   %val = tail call i8 @llvm.nvvm.ldu.global.i.i8.p1(ptr addrspace(1) %ptr, i32 4)
   ret i8 %val
 }
 
+; CHECK: test_ldu_i16
+define i16 @test_ldu_i16(ptr addrspace(1) %ptr) {
+; ldu.global.u16
+  %val = tail call i16 @llvm.nvvm.ldu.global.i.i16.p1(ptr addrspace(1) %ptr, i32 4)
+  ret i16 %val
+}
+
 ; CHECK: test_ldu_i32
 define i32 @test_ldu_i32(ptr addrspace(1) %ptr) {
 ; ldu.global.u32
   %val = tail call i32 @llvm.nvvm.ldu.global.i.i32.p1(ptr addrspace(1) %ptr, i32 4)
   ret i32 %val
 }
 
+; CHECK: test_ldu_i64
+define i64 @test_ldu_i64(ptr addrspace(1) %ptr) {
+; ldu.global.u64
+  %val = tail call i64 @llvm.nvvm.ldu.global.i.i64.p1(ptr addrspace(1) %ptr, i32 8)
+  ret i64 %val
+}
+
+; CHECK: test_ldu_f32
+define float @test_ldu_f32(ptr addrspace(1) %ptr) {
+; ldu.global.u64
+  %val = tail call float @llvm.nvvm.ldu.global.f.f32.p1(ptr addrspace(1) %ptr, i32 4)
+  ret float %val
+}
+
+; CHECK: test_ldu_f64
+define double @test_ldu_f64(ptr addrspace(1) %ptr) {
+; ldu.global.u64
+  %val = tail call double @llvm.nvvm.ldu.global.f.f64.p1(ptr addrspace(1) %ptr, i32 8)
+  ret double %val
+}
+
+; CHECK: test_ldu_f16
+define half @test_ldu_f16(ptr addrspace(1) %ptr) {
+; ldu.global.b16
+  %val = tail call half @llvm.nvvm.ldu.global.f.f16.p1(ptr addrspace(1) %ptr, i32 4)
+  ret half %val
+}
+
+; CHECK: test_ldu_v2f16
+define <2 x half> @test_ldu_v2f16(ptr addrspace(1) %ptr) {
+; ldu.global.b32
+  %val = tail call <2 x half> @llvm.nvvm.ldu.global.f.v2f16.p1(ptr addrspace(1) %ptr, i32 4)
+  ret <2 x half> %val
+}
+
 ; CHECK: test_ldg_i8
 define i8 @test_ldg_i8(ptr addrspace(1) %ptr) {
 ; ld.global.nc.u8
   %val = tail call i8 @llvm.nvvm.ldg.global.i.i8.p1(ptr addrspace(1) %ptr, i32 4)
   ret i8 %val
 }
 
 ; CHECK: test_ldg_i16