[Openmp-commits] [PATCH] D42686: [libomptarget] Only use CUDA Driver API

[Jonas Hahnfeld via Phabricator via Openmp-commits]

Hahnfeld created this revision.
Hahnfeld added a reviewer: grokos.
Herald added a subscriber: openmp-commits.
Hahnfeld added a dependent revision: D42643: [libomptarget] Check for library with CUDA Driver API.

Use equivalents for the last calls to the Runtime API.


Repository:
  rOMP OpenMP

https://reviews.llvm.org/D42686

Files:
  libomptarget/plugins/cuda/src/rtl.cpp


Index: libomptarget/plugins/cuda/src/rtl.cpp
===================================================================
--- libomptarget/plugins/cuda/src/rtl.cpp
+++ libomptarget/plugins/cuda/src/rtl.cpp
@@ -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 @@
   }
 
   // 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 @@
           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
@@ -685,10 +684,12 @@
   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);
+    assert(sync_err == CUDA_SUCCESS &&
+           "Unable to synchronize target execution!");
     return OFFLOAD_FAIL;
   } else {
     DP("Kernel execution at " DPxMOD " successful!\n", DPxPTR(tgt_entry_ptr));


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D42686.131956.patch
Type: text/x-patch
Size: 1989 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20180130/f4187ddc/attachment.bin>