[mlir][sparse][gpu] fixing broken literal names in cuda runner macros

[K-Wu]

more fixes credit to Aart

[K-Wu]

format

mlir/lib/ExecutionEngine/CudaRuntimeWrappers.cpp

 }
 
 ///
 /// Wrapper methods for the cuSparse library.
 ///
 
 // Some macro magic to get float/double alpha and beta on host.
 #define ALPHABETA(dtp, alpha, beta)                                            \
-  __nv_bfloat16(alpha##bf16) = 1.0f;                                           \
-  __nv_bfloat16(beta##bf16) = 1.0f;                                            \
-  __half(alpha##f16) = 1.0f;                                                   \
-  __half(beta##f16) = 1.0f;                                                    \
+  __nv_bfloat16(alpha##16bf) = 1.0f;                                           \
+  __nv_bfloat16(beta##16bf) = 1.0f;                                            \
+  __half(alpha##16f) = 1.0f;                                                   \
+  __half(beta##16f) = 1.0f;                                                    \
   float(alpha##f) = 1.0f;                                                      \
   float(beta##f) = 1.0f;                                                       \
   double(alpha##d) = 1.0;                                                      \
   double(beta##d) = 1.0;                                                       \
   const void *(alpha##p) = nullptr;                                            \
   const void *(beta##p) = nullptr;                                             \
   if (dtp == CUDA_R_16BF || dtp == CUDA_C_16BF) {                              \
     (alpha##p) = reinterpret_cast<void *>(&(alpha##16bf));                     \
     (beta##p) = reinterpret_cast<void *>(&(beta##16bf));                       \
   } else if (dtp == CUDA_R_16F || dtp == CUDA_C_16F) {                         \
     (alpha##p) = reinterpret_cast<void *>(&(alpha##16f));                      \
     (beta##p) = reinterpret_cast<void *>(&(beta##16f));                        \
   } else if (dtp == CUDA_R_32F || dtp == CUDA_C_32F) {                         \
     (alpha##p) = reinterpret_cast<void *>(&(alpha##f));                        \
     (beta##p) = reinterpret_cast<void *>(&(beta##f));                          \
-  } else if (dtp == CUDA_R_64F || dtp == CUDA_C_64F) {                         \
+  } else {                                                                     \
     (alpha##p) = reinterpret_cast<void *>(&(alpha##d));                        \
     (beta##p) = reinterpret_cast<void *>(&(beta##d));                          \
-  } else {                                                                     \
-    llvm_unreachable("Unsupported data type");                                 \
   }
 
 extern "C" MLIR_CUDA_WRAPPERS_EXPORT void *
 mgpuCreateSparseEnv(CUstream /*stream*/) {
   cusparseHandle_t handle = nullptr;
   CUSPARSE_REPORT_IF_ERROR(cusparseCreate(&handle))
   return reinterpret_cast<void *>(handle);
 }

 
 extern "C" MLIR_CUDA_WRAPPERS_EXPORT void *
 mgpuCreateCsr(intptr_t rows, intptr_t cols, intptr_t nnz, void *rowPos,
               void *colIdxs, void *values, int32_t ptp, int32_t itp,
               int32_t dtp, CUstream /*stream*/) {
   cusparseSpMatDescr_t mat = nullptr;
   auto pTp = static_cast<cusparseIndexType_t>(ptp);
   auto iTp = static_cast<cusparseIndexType_t>(itp);
+  auto dTp = static_cast<cudaDataType_t>(dtp);
   CUSPARSE_REPORT_IF_ERROR(cusparseCreateCsr(&mat, rows, cols, nnz, rowPos,
                                              colIdxs, values, pTp, iTp,
                                              CUSPARSE_INDEX_BASE_ZERO, dtp))
   return reinterpret_cast<void *>(mat);
 }
 
 extern "C" MLIR_CUDA_WRAPPERS_EXPORT void
 mgpuDestroySpMat(void *m, CUstream /*stream*/) {