[Openmp-commits] [openmp] r323786 - [libomptarget] Only use CUDA Driver API

[Jonas Hahnfeld via Openmp-commits]

Author: hahnfeld
Date: Tue Jan 30 08:49:06 2018
New Revision: 323786

URL: http://llvm.org/viewvc/llvm-project?rev=323786&view=rev
Log:
[libomptarget] Only use CUDA Driver API

Use equivalents for the last calls to the Runtime API. Remove
stray assert in case of an error found during review, we should
only return OFFLOAD_FAIL.

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

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

Modified: openmp/trunk/libomptarget/plugins/cuda/src/rtl.cpp
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/plugins/cuda/src/rtl.cpp?rev=323786&r1=323785&r2=323786&view=diff
==============================================================================
--- openmp/trunk/libomptarget/plugins/cuda/src/rtl.cpp (original)
+++ openmp/trunk/libomptarget/plugins/cuda/src/rtl.cpp Tue Jan 30 08:49:06 2018
@@ -14,7 +14,6 @@
 #include <cassert>
 #include <cstddef>
 #include <cuda.h>
-#include <cuda_runtime_api.h>
 #include <list>
 #include <string>
 #include <vector>
@@ -280,9 +279,9 @@ int32_t __tgt_rtl_init_device(int32_t de
   }
 
   // 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;
@@ -314,8 +313,8 @@ int32_t __tgt_rtl_init_device(int32_t de
           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
@@ -678,17 +677,16 @@ int32_t __tgt_rtl_run_target_team_region
   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));