[Polly] [GPUJIT] Set min size to 1 on allocation calls.

[bollu]

Summary

All cuda allocation functions fail on size 0 allocations. So, make sure that
we allocate a minimum of size 1.

[bollu]

• [NFC] document why we care about size 0 allocs

[bollu]

@singam-sanjay Review please.

[bollu]

• [NFC] remove TODO with actual code

[grosser]

LGTM

polly/trunk/tools/GPURuntime/GPUJIT.c

 #include <dlfcn.h>
 #include <stdarg.h>
 #include <stdio.h>
 #include <string.h>
 #include <unistd.h>
 
 static int DebugMode;
 static int CacheMode;
+#define max(x, y) ((x) > (y) ? (x) : (y))
 
 static PollyGPURuntime Runtime = RUNTIME_NONE;
 
 static void debug_print(const char *format, ...) {
   if (!DebugMode)
     return;
 
   va_list args;
if (isManagedPtr(mem)) {
     }
     return;
   } else {
     free(mem);
   }
 }
 
 void *polly_mallocManaged(size_t size) {
+  // Note: [Size 0 allocations]
+  // Sometimes, some runtime computation of size could create a size of 0
+  // for an allocation. In these cases, we do not wish to fail.
+  // The CUDA API fails on size 0 allocations.
+  // So, we allocate size a minimum of size 1.
+  if (!size && DebugMode)
+    fprintf(stderr, "cudaMallocManaged called with size 0. "
+                    "Promoting to size 1");
+  size = max(size, 1);
+  polly_initContextCUDA();
   dump_function();
   void *a;
   if (cudaMallocManaged(&a, size, cudaMemAttachGlobal) != cudaSuccess) {
     fprintf(stderr, "cudaMallocManaged failed for size: %zu\n", size);
     exit(-1);
   }
   addManagedPtr(a);
   return a;
 }
 
 static void freeDeviceMemoryCUDA(PollyGPUDevicePtr *Allocation) {
   dump_function();
   CUDADevicePtr *DevPtr = (CUDADevicePtr *)Allocation->DevicePtr;
   CuMemFreeFcnPtr((CUdeviceptr)DevPtr->Cuda);
   free(DevPtr);
   free(Allocation);
 }
 
 static PollyGPUDevicePtr *allocateMemoryForDeviceCUDA(long MemSize) {
+  if (!MemSize && DebugMode)
+    fprintf(stderr, "allocateMemoryForDeviceCUDA called with size 0. "
+                    "Promoting to size 1");
+  // see: [Size 0 allocations]
+  MemSize = max(MemSize, 1);
   dump_function();
 
   PollyGPUDevicePtr *DevData = malloc(sizeof(PollyGPUDevicePtr));
   if (DevData == 0) {
     fprintf(stderr,
             "Allocate memory for GPU device memory pointer failed."
             " Line: %d | Size: %ld\n",
             __LINE__, MemSize);