[libomptarget] Only use CUDA Driver API

[Hahnfeld]

Summary

Use equivalents for the last calls to the Runtime API.

[grokos]

libomptarget/plugins/cuda/src/rtl.cpp

691–692 ↗

(On Diff #131956)

This assertion is guaranteed to trip - it's inside an if (sync_err != CUDA_SUCCESS) - so it doesn't make sense to have it. Besides, if kernel execution fails, the plugin should just return OFFLOAD_FAIL to the parent runtime, not terminate the program.

[Hahnfeld]

libomptarget/plugins/cuda/src/rtl.cpp

691–692 ↗

(On Diff #131956)

I know, I was just mimicking the same behavior of the previous if statement. If that's not what we want, we should probably remove the first assert as well.

[grokos]

libomptarget/plugins/cuda/src/rtl.cpp

691–692 ↗

(On Diff #131956)

Right, the previous assertion must be removed, it shouldn't be there.

[Hahnfeld]

Remove asserts.

[grokos]

LGTM.

openmp/trunk/libomptarget/plugins/cuda/src/rtl.cpp

 //===----RTLs/cuda/src/rtl.cpp - Target RTLs Implementation ------- C++ -*-===//
 //
 //                     The LLVM Compiler Infrastructure
 //
 // This file is dual licensed under the MIT and the University of Illinois Open
 // Source Licenses. See LICENSE.txt for details.
 //
 //===----------------------------------------------------------------------===//
 //
 // RTL for CUDA machine
 //
 //===----------------------------------------------------------------------===//
 
 #include <cassert>
 #include <cstddef>
 #include <cuda.h>
-#include <cuda_runtime_api.h>
 #include <list>
 #include <string>
 #include <vector>
 
 #include "omptargetplugin.h"
 
 #ifndef TARGET_NAME
 #define TARGET_NAME CUDA
err = cuCtxCreate(&DeviceInfo.Contexts[device_id], CU_CTX_SCHED_BLOCKING_SYNC,
                     cuDevice);
   if (err != CUDA_SUCCESS) {
     DP("Error when creating a CUDA context\n");
     CUDA_ERR_STRING(err);
     return OFFLOAD_FAIL;
   }
 
   // scan properties to determine number of threads/block and blocks/grid.
-  struct cudaDeviceProp Properties;
-  cudaError_t error = cudaGetDeviceProperties(&Properties, device_id);
-  if (error != cudaSuccess) {
+  CUdevprop Properties;
+  err = cuDeviceGetProperties(&Properties, cuDevice);
+  if (err != CUDA_SUCCESS) {
     DP("Error getting device Properties, use defaults\n");
     DeviceInfo.BlocksPerGrid[device_id] = RTLDeviceInfoTy::DefaultNumTeams;
     DeviceInfo.ThreadsPerBlock[device_id] = RTLDeviceInfoTy::DefaultNumThreads;
     DeviceInfo.WarpSize[device_id] = 32;
   } else {
     // Get blocks per grid
     if (Properties.maxGridSize[0] <= RTLDeviceInfoTy::HardTeamLimit) {
       DeviceInfo.BlocksPerGrid[device_id] = Properties.maxGridSize[0];
if (Properties.maxThreadsDim[0] <= RTLDeviceInfoTy::HardThreadLimit) {
       }
     } else {
       DeviceInfo.ThreadsPerBlock[device_id] = RTLDeviceInfoTy::HardThreadLimit;
       DP("Max CUDA threads per block %d exceeds the hard thread limit %d, "
           "capping at the hard limit\n", Properties.maxThreadsDim[0],
           RTLDeviceInfoTy::HardThreadLimit);
     }
 
-    // Get warp size
-    DeviceInfo.WarpSize[device_id] = Properties.warpSize;
+    // According to the documentation, SIMDWidth is "Warp size in threads".
+    DeviceInfo.WarpSize[device_id] = Properties.SIMDWidth;
   }
 
   // Adjust teams to the env variables
   if (DeviceInfo.EnvTeamLimit > 0 &&
       DeviceInfo.BlocksPerGrid[device_id] > DeviceInfo.EnvTeamLimit) {
     DeviceInfo.BlocksPerGrid[device_id] = DeviceInfo.EnvTeamLimit;
     DP("Capping max CUDA blocks per grid to OMP_TEAM_LIMIT=%d\n",
         DeviceInfo.EnvTeamLimit);
int32_t __tgt_rtl_run_target_team_region(int32_t device_id, void *tgt_entry_ptr,
   DP("Launch kernel with %d blocks and %d threads\n", cudaBlocksPerGrid,
      cudaThreadsPerBlock);
 
   err = cuLaunchKernel(KernelInfo->Func, cudaBlocksPerGrid, 1, 1,
       cudaThreadsPerBlock, 1, 1, 0 /*bytes of shared memory*/, 0, &args[0], 0);
   if (err != CUDA_SUCCESS) {
     DP("Device kernel launch failed!\n");
     CUDA_ERR_STRING(err);
-    assert(err == CUDA_SUCCESS && "Unable to launch target execution!");
     return OFFLOAD_FAIL;
   }
 
   DP("Launch of entry point at " DPxMOD " successful!\n",
       DPxPTR(tgt_entry_ptr));
 
-  cudaError_t sync_error = cudaDeviceSynchronize();
-  if (sync_error != cudaSuccess) {
-  DP("Kernel execution error at " DPxMOD ", %s.\n", DPxPTR(tgt_entry_ptr),
-      cudaGetErrorString(sync_error));
+  CUresult sync_err = cuCtxSynchronize();
+  if (sync_err != CUDA_SUCCESS) {
+    DP("Kernel execution error at " DPxMOD "!\n", DPxPTR(tgt_entry_ptr));
+    CUDA_ERR_STRING(sync_err);
     return OFFLOAD_FAIL;
   } else {
     DP("Kernel execution at " DPxMOD " successful!\n", DPxPTR(tgt_entry_ptr));
   }
 
   return OFFLOAD_SUCCESS;
 }