diff --git a/openmp/libomptarget/plugins/cuda/src/rtl.cpp b/openmp/libomptarget/plugins/cuda/src/rtl.cpp --- a/openmp/libomptarget/plugins/cuda/src/rtl.cpp +++ b/openmp/libomptarget/plugins/cuda/src/rtl.cpp @@ -580,11 +580,12 @@ DeviceData[DeviceId].BlocksPerGrid = EnvTeamLimit; } - INFO(DeviceId, - "Device supports up to %d CUDA blocks and %d threads with a " - "warp size of %d\n", - DeviceData[DeviceId].BlocksPerGrid, - DeviceData[DeviceId].ThreadsPerBlock, DeviceData[DeviceId].WarpSize); + if (getDebugLevel() || (getInfoLevel() & OMP_INFOTYPE_PLUGIN_KERNEL)) + INFO(DeviceId, + "Device supports up to %d CUDA blocks and %d threads with a " + "warp size of %d\n", + DeviceData[DeviceId].BlocksPerGrid, + DeviceData[DeviceId].ThreadsPerBlock, DeviceData[DeviceId].WarpSize); // Set default number of teams if (EnvNumTeams > 0) { @@ -1004,14 +1005,15 @@ CudaBlocksPerGrid = TeamNum; } - INFO(DeviceId, - "Launching kernel %s with %d blocks and %d threads in %s " - "mode\n", - (getOffloadEntry(DeviceId, TgtEntryPtr)) - ? getOffloadEntry(DeviceId, TgtEntryPtr)->name - : "(null)", - CudaBlocksPerGrid, CudaThreadsPerBlock, - (KernelInfo->ExecutionMode == SPMD) ? "SPMD" : "Generic"); + if (getDebugLevel() || (getInfoLevel() & OMP_INFOTYPE_PLUGIN_KERNEL)) + INFO(DeviceId, + "Launching kernel %s with %d blocks and %d threads in %s " + "mode\n", + (getOffloadEntry(DeviceId, TgtEntryPtr)) + ? getOffloadEntry(DeviceId, TgtEntryPtr)->name + : "(null)", + CudaBlocksPerGrid, CudaThreadsPerBlock, + (KernelInfo->ExecutionMode == SPMD) ? "SPMD" : "Generic"); CUstream Stream = getStream(DeviceId, AsyncInfo); Err = cuLaunchKernel(KernelInfo->Func, CudaBlocksPerGrid, /* gridDimY */ 1, diff --git a/openmp/libomptarget/src/private.h b/openmp/libomptarget/src/private.h --- a/openmp/libomptarget/src/private.h +++ b/openmp/libomptarget/src/private.h @@ -142,7 +142,7 @@ type = "private"; else if (ArgTypes[i] & OMP_TGT_MAPTYPE_LITERAL) type = "firstprivate"; - else if (ArgTypes[i] & OMP_TGT_MAPTYPE_TARGET_PARAM && ArgSizes[i] != 0) + else if (ArgSizes[i] != 0) type = "alloc"; else type = "use_address";