[PATCH] D45780: [CUDA] added missing __ldg(const signed char *)

[Artem Belevich via Phabricator via cfe-commits]

tra created this revision.
tra added a reviewer: jlebar.
Herald added a subscriber: sanjoy.

Until now we only had variants for 'char' and 'unsigned char'. In C++ 'char' 
'unsigned char' and 'signed char' are three different types and we need 
overloads for all of them.


https://reviews.llvm.org/D45780

Files:
  clang/lib/Headers/__clang_cuda_intrinsics.h


Index: clang/lib/Headers/__clang_cuda_intrinsics.h
===================================================================
--- clang/lib/Headers/__clang_cuda_intrinsics.h
+++ clang/lib/Headers/__clang_cuda_intrinsics.h
@@ -277,6 +277,9 @@
 inline __device__ unsigned char __ldg(const unsigned char *ptr) {
   return __nvvm_ldg_uc(ptr);
 }
+inline __device__ signed char __ldg(const signed char *ptr) {
+  return __nvvm_ldg_uc((const unsigned char *)ptr);
+}
 inline __device__ unsigned short __ldg(const unsigned short *ptr) {
   return __nvvm_ldg_us(ptr);
 }



Index: clang/lib/Headers/__clang_cuda_intrinsics.h
===================================================================
--- clang/lib/Headers/__clang_cuda_intrinsics.h
+++ clang/lib/Headers/__clang_cuda_intrinsics.h
@@ -277,6 +277,9 @@
 inline __device__ unsigned char __ldg(const unsigned char *ptr) {
   return __nvvm_ldg_uc(ptr);
 }
+inline __device__ signed char __ldg(const signed char *ptr) {
+  return __nvvm_ldg_uc((const unsigned char *)ptr);
+}
 inline __device__ unsigned short __ldg(const unsigned short *ptr) {
   return __nvvm_ldg_us(ptr);
 }

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits