[OpenMP][Offloading] Replaced default stream with an actual per-device unblocking stream in NVPTX implementation

[tianshilei1992]

Summary

In this patch, default stream is replaced with an actual per-device stream for better performance as there're plenty of constraints in default stream, according to https://developer.download.nvidia.com/CUDA/training/StreamsAndConcurrencyWebinar.pdf. Later I'll enable multiple streams to improve the concurrency.

[ABataev]

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

260

1. Use real type instead of auto
2. Variables must start with the upper case letter.

[tianshilei1992]

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

260

Sure but just to keep align with existing code. Do I need to update existing code as well?

[ABataev]

Also, the main question, how does it affect the exiting execution model? What if we have target region in a parallel region, will they be executed asynchronously? We need some tests for this if we don't have such tests.

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

260

In a separate patch, please

[tianshilei1992]

In D70010#1738930, @ABataev wrote:

Also, the main question, how does it affect the exiting execution model? What if we have target region in a parallel region, will they be executed asynchronously? We need some tests for this if we don't have such tests.

According to https://developer.download.nvidia.com/CUDA/training/StreamsAndConcurrencyWebinar.pdf, non-default stream can improve performance. This is actually the first step to use multiple streams I'm gonna implement later.

[ABataev]

In D70010#1739049, @tianshilei1992 wrote:

In D70010#1738930, @ABataev wrote:

Also, the main question, how does it affect the exiting execution model? What if we have target region in a parallel region, will they be executed asynchronously? We need some tests for this if we don't have such tests.

According to https://developer.download.nvidia.com/CUDA/training/StreamsAndConcurrencyWebinar.pdf, non-default stream can improve performance. This is actually the first step to use multiple streams I'm gonna implement later.

My question is different. Does it affect execution of the existing code anyhow?

[tianshilei1992]

In D70010#1739061, @ABataev wrote:

In D70010#1739049, @tianshilei1992 wrote:

In D70010#1738930, @ABataev wrote:

Also, the main question, how does it affect the exiting execution model? What if we have target region in a parallel region, will they be executed asynchronously? We need some tests for this if we don't have such tests.

According to https://developer.download.nvidia.com/CUDA/training/StreamsAndConcurrencyWebinar.pdf, non-default stream can improve performance. This is actually the first step to use multiple streams I'm gonna implement later.

My question is different. Does it affect execution of the existing code anyhow?

AFAIK, no. Currently we still only have one stream for each device, but it's just not the default stream. Kernels in a stream are executed in order. The asynchronous execution requires multiple streams. I'll check whether existing cases can cover it, and will write one if no.

[jdoerfert]

Can you please write a commit message explaining the change and the plan?

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

260

You can go either way but I would just keep it as the surrounding code for now.
Making the code blend in is arguably a good think and we have not decided how we handle the plugins in the short term (wrt. coding standard).

@JonChesterfield @grokos What do you think? Should we run once over the plugin and adjust the coding style now or keep it consistent for the time being?

[grokos]

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

260

I would say let's keep it consistent. Later on, we can adjust the code style for the whole library but until then I prefer consistency over mixed styles.

[tianshilei1992]

In D70010#1739316, @jdoerfert wrote:

Can you please write a commit message explaining the change and the plan?

No problem. Will do.

[tianshilei1992]

This one can be abandoned now...

[grokos]

Superseded by D74145. You can abandon this one.

[tianshilei1992]

In D70010#1966048, @grokos wrote:

Superseded by D74145. You can abandon this one.

I didn't see the option?

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

 /// Class containing all the device information.
 class RTLDeviceInfoTy {
   std::vector<std::list<FuncOrGblEntryTy>> FuncGblEntries;
 
 public:
   int NumberOfDevices;
   std::vector<CUmodule> Modules;
   std::vector<CUcontext> Contexts;
+  std::vector<CUstream> Streams;
 
   // Device properties
   std::vector<int> ThreadsPerBlock;
   std::vector<int> BlocksPerGrid;
   std::vector<int> WarpSize;
 
   // OpenMP properties
   std::vector<int> NumTeams;
#endif // OMPTARGET_DEBUG
 
     if (NumberOfDevices == 0) {
       DP("There are no devices supporting CUDA.\n");
       return;
     }
 
     FuncGblEntries.resize(NumberOfDevices);
     Contexts.resize(NumberOfDevices);
+    Streams.resize(NumberOfDevices);
     ThreadsPerBlock.resize(NumberOfDevices);
     BlocksPerGrid.resize(NumberOfDevices);
     WarpSize.resize(NumberOfDevices);
     NumTeams.resize(NumberOfDevices);
     NumThreads.resize(NumberOfDevices);
 
     // Get environment variables regarding teams
     char *envStr = getenv("OMP_TEAM_LIMIT");
~RTLDeviceInfoTy() {
     for (auto &ctx : Contexts)
       if (ctx) {
         CUresult err = cuCtxDestroy(ctx);
         if (err != CUDA_SUCCESS) {
           DP("Error when destroying CUDA context\n");
           CUDA_ERR_STRING(err);
         }
       }
+
+    // Destroy streams
+    for (auto &stream : Streams)
+      if (stream) {
+        CUresult err = cuStreamDestroy(stream);
+        if (err != CUDA_SUCCESS) {
+          DP("Error when destroying CUDA stream\n");
+          CUDA_ERR_STRING(err);
+        }
+      }
   }
 };
 
 static RTLDeviceInfoTy DeviceInfo;
 
 #ifdef __cplusplus
 extern "C" {
 #endif
int32_t __tgt_rtl_init_device(int32_t device_id) {
   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;
   }
 
+  // Set current context for later creating corresponding stream
+  err = cuCtxSetCurrent(DeviceInfo.Contexts[device_id]);
+  if (err != CUDA_SUCCESS) {
+    DP("Error when setting current CUDA context\n");
+    CUDA_ERR_STRING(err);
+    return OFFLOAD_FAIL;
+  }
+
+  //Create a stream for each device
+  err = cuStreamCreate(&DeviceInfo.Streams[device_id], CU_STREAM_NON_BLOCKING);
+  if (err != CUDA_SUCCESS) {
+    DP("Error when creating CUDA stream\n");
+    CUDA_ERR_STRING(err);
+    return OFFLOAD_FAIL;
+  }
+
   // Query attributes to determine number of threads/block and blocks/grid.
   int maxGridDimX;
   err = cuDeviceGetAttribute(&maxGridDimX, CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_X,
                              cuDevice);
   if (err != CUDA_SUCCESS) {
     DP("Error getting max grid dimension, use default\n");
     DeviceInfo.BlocksPerGrid[device_id] = RTLDeviceInfoTy::DefaultNumTeams;
   } else if (maxGridDimX <= RTLDeviceInfoTy::HardTeamLimit) {
if (team_num <= 0) {
     DP("Using requested number of teams %d\n", team_num);
   }
 
   // Run on the device.
   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);
+                       cudaThreadsPerBlock, 1, 1, 0 /*bytes of shared memory*/,
+                       DeviceInfo.Streams[device_id], &args[0], 0);
   if (err != CUDA_SUCCESS) {
     DP("Device kernel launch failed!\n");
     CUDA_ERR_STRING(err);
     return OFFLOAD_FAIL;
   }
 
   DP("Launch of entry point at " DPxMOD " successful!\n",
       DPxPTR(tgt_entry_ptr));