[polly] r310941 - [Polly] [GPUJIT] Set min size to 1 on CUDA allocation calls. [NFC]

[Siddharth Bhat via llvm-commits]

Author: bollu
Date: Tue Aug 15 11:21:38 2017
New Revision: 310941

URL: http://llvm.org/viewvc/llvm-project?rev=310941&view=rev
Log:
[Polly] [GPUJIT] Set min size to 1 on CUDA allocation calls. [NFC]

Requesting size 0 allocations from `cuMalloc` / `cuMallocManaged` fails.
If there is a size 0 allocation that can be statically proved, the we
fail at PPCGCodeGeneration. This is because if size 0 allocation could
take place, we should not generate code that tries to use this array.

However, there are cases where we cannot statically prove this, and at
runtime we get a request for 0 bytes of memory. We choose to allocate
size 1 to allow the program to continue running.

Differential Revision: https://reviews.llvm.org/D36751

Modified:
    polly/trunk/tools/GPURuntime/GPUJIT.c

Modified: polly/trunk/tools/GPURuntime/GPUJIT.c
URL: http://llvm.org/viewvc/llvm-project/polly/trunk/tools/GPURuntime/GPUJIT.c?rev=310941&r1=310940&r2=310941&view=diff
==============================================================================
--- polly/trunk/tools/GPURuntime/GPUJIT.c (original)
+++ polly/trunk/tools/GPURuntime/GPUJIT.c Tue Aug 15 11:21:38 2017
@@ -35,6 +35,7 @@
 
 static int DebugMode;
 static int CacheMode;
+#define max(x, y) ((x) > (y) ? (x) : (y))
 
 static PollyGPURuntime Runtime = RUNTIME_NONE;
 
@@ -1455,6 +1456,16 @@ void polly_freeManaged(void *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) {
@@ -1474,6 +1485,11 @@ static void freeDeviceMemoryCUDA(PollyGP
 }
 
 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));