Index: tools/GPURuntime/GPUJIT.c =================================================================== --- tools/GPURuntime/GPUJIT.c +++ tools/GPURuntime/GPUJIT.c @@ -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,18 @@ } 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"); + // TODO: print warning message if size == 0 if POLLY_DEBUG=1, so that + // the user knows that this behaviour is kicking in. + size = max(size, 1); + polly_initContextCUDA(); dump_function(); void *a; if (cudaMallocManaged(&a, size, cudaMemAttachGlobal) != cudaSuccess) { @@ -1474,6 +1487,11 @@ } 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));