[offload][OpenMP] Add strict flag for blocks and threads in kernel arguments (llvm#199483)

kevinsala · web-flow · commit f15904ec71a8 · 2026-05-30T02:28:35.000-07:00
Until now, strict behavior in the number of threads and blocks has been
applied only when the kernel is in bare mode. When this mode is enabled,
the values passed in UserNumBlocks and UserThreadLimit are not adjusted
and are the definitive values used to launch the kernel. This commit
detaches the strictness from the kernel mode.

This is going to be used by the kernel replay tool. Additionally, it
starts clearing the path for the upcoming OpenMP dims modifier, used to
configure multidimensional teams and leagues, which will include
strictness choices for teams and threads.

All the bare kernels must indicate strict behavior. Asserts are added to
check this condition.
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -10982,7 +10982,8 @@ static void emitTargetCallKernelLaunch(
 
     llvm::OpenMPIRBuilder::TargetKernelArgs Args(
         NumTargetItems, RTArgs, NumIterations, NumTeams, NumThreads,
-        DynCGroupMem, HasNoWait, DynCGroupMemFallback);
+        DynCGroupMem, HasNoWait, /*StrictBlocksAndThreads=*/IsBare,
+        DynCGroupMemFallback);
 
     llvm::OpenMPIRBuilder::InsertPointTy AfterIP =
         cantFail(OMPRuntime->getOMPBuilder().emitKernelLaunch(
diff --git a/clang/test/OpenMP/target_teams_codegen.cpp b/clang/test/OpenMP/target_teams_codegen.cpp
diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
@@ -2771,6 +2771,9 @@ class OpenMPIRBuilder {
     Value *DynCGroupMem = nullptr;
     /// True if the kernel has 'no wait' clause.
     bool HasNoWait = false;
+    /// True if the kernel strictly requires the number of blocks and threads
+    /// above to run.
+    bool StrictBlocksAndThreads = false;
     /// The fallback mechanism for the shared memory.
     omp::OMPDynGroupprivateFallbackType DynCGroupMemFallback =
         omp::OMPDynGroupprivateFallbackType::Abort;
@@ -2780,12 +2783,13 @@ class OpenMPIRBuilder {
     TargetKernelArgs(unsigned NumTargetItems, TargetDataRTArgs RTArgs,
                      Value *NumIterations, ArrayRef<Value *> NumTeams,
                      ArrayRef<Value *> NumThreads, Value *DynCGroupMem,
-                     bool HasNoWait,
+                     bool HasNoWait, bool StrictBlocksAndThreads,
                      omp::OMPDynGroupprivateFallbackType DynCGroupMemFallback)
         : NumTargetItems(NumTargetItems), RTArgs(RTArgs),
           NumIterations(NumIterations), NumTeams(NumTeams),
           NumThreads(NumThreads), DynCGroupMem(DynCGroupMem),
-          HasNoWait(HasNoWait), DynCGroupMemFallback(DynCGroupMemFallback) {}
+          HasNoWait(HasNoWait), StrictBlocksAndThreads(StrictBlocksAndThreads),
+          DynCGroupMemFallback(DynCGroupMemFallback) {}
   };
 
   /// Create the kernel args vector used by emitTargetKernel. This function
diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -647,7 +647,12 @@ void OpenMPIRBuilder::getKernelArgsVector(TargetKernelArgs &KernelArgs,
   Value *DynCGroupMemFallbackFlag =
       Builder.getInt64(static_cast<uint64_t>(KernelArgs.DynCGroupMemFallback));
   DynCGroupMemFallbackFlag = Builder.CreateShl(DynCGroupMemFallbackFlag, 2);
+
+  Value *StrictFlag = Builder.getInt64(KernelArgs.StrictBlocksAndThreads);
+  StrictFlag = Builder.CreateShl(StrictFlag, 6);
+
   Value *Flags = Builder.CreateOr(HasNoWaitFlag, DynCGroupMemFallbackFlag);
+  Flags = Builder.CreateOr(Flags, StrictFlag);
 
   assert(!KernelArgs.NumTeams.empty() && !KernelArgs.NumThreads.empty());
 
@@ -9783,7 +9788,7 @@ static void emitTargetCall(
 
     KArgs = OpenMPIRBuilder::TargetKernelArgs(
         NumTargetItems, RTArgs, TripCount, NumTeamsC, NumThreadsC, DynCGroupMem,
-        HasNoWait, DynCGroupMemFallback);
+        HasNoWait, /*StrictBlocksAndThreads=*/false, DynCGroupMemFallback);
 
     // Assume no error was returned because TaskBodyCB and
     // EmitTargetCallFallbackCB don't produce any.
diff --git a/offload/include/Shared/APITypes.h b/offload/include/Shared/APITypes.h
@@ -106,8 +106,10 @@ struct KernelArgsTy {
     uint64_t DynCGroupMemFallback : 2; // The fallback for dynamic cgroup mem.
     uint64_t Cooperative : 1; // Was this kernel spawned as cooperative.
     uint64_t IsPtrArgs : 1;   // Arguments are laid out as an array of pointers.
-    uint64_t Unused : 58;
-  } Flags = {0, 0, 0, 0, 0, 0};
+    uint64_t StrictBlocksAndThreads
+        : 1; // The user-requested number of blocks and threads are strict.
+    uint64_t Unused : 57;
+  } Flags = {0, 0, 0, 0, 0, 0, 0};
   // User-requested number of blocks (for x,y,z dimension).
   uint32_t UserNumBlocks[3] = {0, 0, 0};
   // User-requested number of threads (for x,y,z dimension).
diff --git a/offload/liboffload/src/OffloadImpl.cpp b/offload/liboffload/src/OffloadImpl.cpp
@@ -1124,6 +1124,7 @@ Error olLaunchKernel_impl(ol_queue_handle_t Queue, ol_device_handle_t Device,
   LaunchArgs.UserThreadLimit[1] = LaunchSizeArgs->GroupSize.y;
   LaunchArgs.UserThreadLimit[2] = LaunchSizeArgs->GroupSize.z;
   LaunchArgs.DynCGroupMem = LaunchSizeArgs->DynSharedMemory;
+  LaunchArgs.Flags.StrictBlocksAndThreads = true;
 
   while (Properties && Properties->type != OL_KERNEL_LAUNCH_PROP_TYPE_NONE) {
     switch (Properties->type) {
diff --git a/offload/libomptarget/KernelLanguage/API.cpp b/offload/libomptarget/KernelLanguage/API.cpp
@@ -68,6 +68,7 @@ unsigned llvmLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim,
   Args.UserThreadLimit[2] = blockDim.z;
   Args.ArgPtrs = reinterpret_cast<void **>(args);
   Args.Flags.IsCUDA = true;
+  Args.Flags.StrictBlocksAndThreads = true;
   return __tgt_target_kernel(nullptr, 0, gridDim.x, blockDim.x, func, &Args);
 }
 }
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
@@ -2481,6 +2481,7 @@ int target_replay(ident_t *Loc, DeviceTy &Device, void *HostPtr,
   KernelArgs.UserThreadLimit[1] = 1;
   KernelArgs.UserThreadLimit[2] = 1;
   KernelArgs.DynCGroupMem = SharedMemorySize;
+  KernelArgs.Flags.StrictBlocksAndThreads = true;
 
   KernelExtraArgsTy KernelExtraArgs{};
   KernelExtraArgs.ReplayOutcome = ReplayOutcome;
diff --git a/offload/plugins-nextgen/common/include/PluginInterface.h b/offload/plugins-nextgen/common/include/PluginInterface.h
@@ -529,16 +529,15 @@ struct GenericKernelTy {
   /// Get the effective number of threads for the kernel based on the
   /// user-defined number of threads.
   uint32_t getEffectiveNumThreads(GenericDeviceTy &GenericDevice,
-                                  uint32_t UserThreadLimit[3]) const;
+                                  uint32_t UserThreadLimit) const;
 
   /// Get the effective number of blocks for the kernel based on the
   /// user-defined number of blocks and the loop trip count.
   /// The number of threads \p NumThreads can be adjusted by this method.
   /// \p IsNumThreadsFromUser is true is \p NumThreads is defined by user via
   /// thread_limit clause.
   uint32_t getEffectiveNumBlocks(GenericDeviceTy &GenericDevice,
-                                 uint32_t UserNumBlocks[3],
-                                 uint64_t LoopTripCount,
+                                 uint32_t UserNumBlocks, uint64_t LoopTripCount,
                                  uint32_t &EffectiveNumThreads,
                                  bool IsNumThreadsFromUser) const;
 
diff --git a/offload/plugins-nextgen/common/src/PluginInterface.cpp b/offload/plugins-nextgen/common/src/PluginInterface.cpp
@@ -249,15 +249,27 @@ Error GenericKernelTy::launch(GenericDeviceTy &GenericDevice, void **ArgPtrs,
   uint32_t EffectiveNumBlocks[3] = {KernelArgs.UserNumBlocks[0],
                                     KernelArgs.UserNumBlocks[1],
                                     KernelArgs.UserNumBlocks[2]};
-  if (!isBareMode()) {
-    assert(
-        EffectiveNumThreads[1] == 1 && EffectiveNumThreads[2] == 1 &&
-        EffectiveNumBlocks[1] == 1 && EffectiveNumBlocks[2] == 1 &&
-        "Non-bare mode should only use the first thread and block dimensions");
+
+  // Multidimensional is only supported with bare mode for now.
+  assert(isBareMode() ||
+         EffectiveNumThreads[1] == 1 && EffectiveNumThreads[2] == 1 &&
+             EffectiveNumBlocks[1] == 1 && EffectiveNumBlocks[2] == 1 &&
+             "Non-bare mode should only use the first thread and block "
+             "dimensions");
+
+  assert(!KernelArgs.Flags.StrictBlocksAndThreads ||
+         EffectiveNumThreads[0] > 0 && EffectiveNumThreads[1] > 0 &&
+             EffectiveNumThreads[2] > 0 && EffectiveNumBlocks[0] > 0 &&
+             EffectiveNumBlocks[1] > 0 && EffectiveNumBlocks[2] > 0 &&
+             "Strict requires number of blocks and threads greater than zero");
+
+  // Calculate or adjust the effective number of threads and blocks if needed.
+  if (!KernelArgs.Flags.StrictBlocksAndThreads) {
     EffectiveNumThreads[0] =
-        getEffectiveNumThreads(GenericDevice, EffectiveNumThreads);
+        getEffectiveNumThreads(GenericDevice, EffectiveNumThreads[0]);
+
     EffectiveNumBlocks[0] = getEffectiveNumBlocks(
-        GenericDevice, EffectiveNumBlocks, KernelArgs.Tripcount,
+        GenericDevice, EffectiveNumBlocks[0], KernelArgs.Tripcount,
         EffectiveNumThreads[0], KernelArgs.UserThreadLimit[0] > 0);
   }
 
@@ -364,34 +376,27 @@ GenericKernelTy::prepareArgs(GenericDeviceTy &GenericDevice, void **ArgPtrs,
 
 uint32_t
 GenericKernelTy::getEffectiveNumThreads(GenericDeviceTy &GenericDevice,
-                                        uint32_t UserThreadLimit[3]) const {
+                                        uint32_t UserThreadLimit) const {
   assert(!isBareMode() && "bare kernel should not call this function");
 
-  assert(UserThreadLimit[1] == 1 && UserThreadLimit[2] == 1 &&
-         "Multi dimensional launch not supported yet.");
+  if (UserThreadLimit > 0 && isGenericMode())
+    UserThreadLimit += GenericDevice.getWarpSize();
 
-  if (UserThreadLimit[0] > 0 && isGenericMode())
-    UserThreadLimit[0] += GenericDevice.getWarpSize();
-
-  return std::min(MaxNumThreads, (UserThreadLimit[0] > 0)
-                                     ? UserThreadLimit[0]
-                                     : PreferredNumThreads);
+  return std::min(MaxNumThreads, (UserThreadLimit > 0) ? UserThreadLimit
+                                                       : PreferredNumThreads);
 }
 
 uint32_t GenericKernelTy::getEffectiveNumBlocks(
-    GenericDeviceTy &GenericDevice, uint32_t UserNumBlocks[3],
+    GenericDeviceTy &GenericDevice, uint32_t UserNumBlocks,
     uint64_t LoopTripCount, uint32_t &EffectiveNumThreads,
     bool IsNumThreadsFromUser) const {
   assert(!isBareMode() && "bare kernel should not call this function");
 
-  assert(UserNumBlocks[1] == 1 && UserNumBlocks[2] == 1 &&
-         "Multi dimensional launch not supported yet.");
-
-  if (UserNumBlocks[0] > 0) {
+  if (UserNumBlocks > 0) {
     // TODO: We need to honor any value and consequently allow more than the
     // block limit. For this we might need to start multiple kernels or let the
     // blocks start again until the requested number has been started.
-    return std::min(UserNumBlocks[0], GenericDevice.getBlockLimit());
+    return std::min(UserNumBlocks, GenericDevice.getBlockLimit());
   }
 
   // Return the number of blocks required to cover the loop iterations.
diff --git a/offload/test/offloading/ompx_bare_gridsize.c b/offload/test/offloading/ompx_bare_gridsize.c
@@ -0,0 +1,153 @@
+// RUN: %libomptarget-compile-generic
+// RUN: env LIBOMPTARGET_INFO=63 %libomptarget-run-generic 2>&1 | \
+// RUN:   %fcheck-generic
+//
+// REQUIRES: gpu
+
+#include <ompx.h>
+#include <stdio.h>
+
+void get_gridsizes(int *nblocks, int *nthreads) {
+  if (ompx_block_id_x() == 0 && ompx_thread_id_x() == 0 &&
+      ompx_block_id_y() == 0 && ompx_thread_id_y() == 0 &&
+      ompx_block_id_z() == 0 && ompx_thread_id_z() == 0) {
+    nblocks[0] = ompx_grid_dim_x();
+    nblocks[1] = ompx_grid_dim_y();
+    nblocks[2] = ompx_grid_dim_z();
+    nthreads[0] = ompx_block_dim_x();
+    nthreads[1] = ompx_block_dim_y();
+    nthreads[2] = ompx_block_dim_z();
+  }
+}
+
+int main(int argc, char *argv[]) {
+  int nblocks[3], nthreads[3];
+
+  // CHECK: PluginInterface device 0 info: Launching kernel
+  // CHECK-SAME: __omp_offloading_{{.*}} with [64,1,1] blocks and [32,1,1]
+  // CHECK-SAME: threads in BARE mode
+  nblocks[0] = nblocks[1] = nblocks[2] = nthreads[0] = nthreads[1] =
+      nthreads[2] = 0;
+#pragma omp target teams ompx_bare num_teams(64) thread_limit(32)              \
+    map(tofrom : nblocks, nthreads)
+  {
+    get_gridsizes(nblocks, nthreads);
+  }
+  // CHECK: nblocks: 64 1 1, nthreads: 32 1 1
+  fprintf(stderr, "nblocks: %d %d %d, nthreads: %d %d %d\n", nblocks[0],
+          nblocks[1], nblocks[2], nthreads[0], nthreads[1], nthreads[2]);
+
+  // CHECK: PluginInterface device 0 info: Launching kernel
+  // CHECK-SAME: __omp_offloading_{{.*}} with [64,1,1] blocks and [32,4,1]
+  // CHECK-SAME: threads in BARE mode
+  nblocks[0] = nblocks[1] = nblocks[2] = nthreads[0] = nthreads[1] =
+      nthreads[2] = 0;
+#pragma omp target teams ompx_bare num_teams(64) thread_limit(32, 4)           \
+    map(tofrom : nblocks, nthreads)
+  {
+    get_gridsizes(nblocks, nthreads);
+  }
+  // CHECK: nblocks: 64 1 1, nthreads: 32 4 1
+  fprintf(stderr, "nblocks: %d %d %d, nthreads: %d %d %d\n", nblocks[0],
+          nblocks[1], nblocks[2], nthreads[0], nthreads[1], nthreads[2]);
+
+  // CHECK: PluginInterface device 0 info: Launching kernel
+  // CHECK-SAME: __omp_offloading_{{.*}} with [64,1,1] blocks and [32,4,2]
+  // CHECK-SAME: threads in BARE mode
+  nblocks[0] = nblocks[1] = nblocks[2] = nthreads[0] = nthreads[1] =
+      nthreads[2] = 0;
+#pragma omp target teams ompx_bare num_teams(64) thread_limit(32, 4, 2)        \
+    map(tofrom : nblocks, nthreads)
+  {
+    get_gridsizes(nblocks, nthreads);
+  }
+  // CHECK: nblocks: 64 1 1, nthreads: 32 4 2
+  fprintf(stderr, "nblocks: %d %d %d, nthreads: %d %d %d\n", nblocks[0],
+          nblocks[1], nblocks[2], nthreads[0], nthreads[1], nthreads[2]);
+
+  // CHECK: PluginInterface device 0 info: Launching kernel
+  // CHECK-SAME: __omp_offloading_{{.*}} with [64,16,1] blocks and [32,1,1]
+  // CHECK-SAME: threads in BARE mode
+  nblocks[0] = nblocks[1] = nblocks[2] = nthreads[0] = nthreads[1] =
+      nthreads[2] = 0;
+#pragma omp target teams ompx_bare num_teams(64, 16, 1) thread_limit(32)       \
+    map(tofrom : nblocks, nthreads)
+  {
+    get_gridsizes(nblocks, nthreads);
+  }
+  // CHECK: nblocks: 64 16 1, nthreads: 32 1 1
+  fprintf(stderr, "nblocks: %d %d %d, nthreads: %d %d %d\n", nblocks[0],
+          nblocks[1], nblocks[2], nthreads[0], nthreads[1], nthreads[2]);
+
+  // CHECK: PluginInterface device 0 info: Launching kernel
+  // CHECK-SAME: __omp_offloading_{{.*}} with [64,16,1] blocks and [32,4,1]
+  // CHECK-SAME: threads in BARE mode
+  nblocks[0] = nblocks[1] = nblocks[2] = nthreads[0] = nthreads[1] =
+      nthreads[2] = 0;
+#pragma omp target teams ompx_bare num_teams(64, 16, 1) thread_limit(32, 4)    \
+    map(tofrom : nblocks, nthreads)
+  {
+    get_gridsizes(nblocks, nthreads);
+  }
+  // CHECK: nblocks: 64 16 1, nthreads: 32 4 1
+  fprintf(stderr, "nblocks: %d %d %d, nthreads: %d %d %d\n", nblocks[0],
+          nblocks[1], nblocks[2], nthreads[0], nthreads[1], nthreads[2]);
+
+  // CHECK: PluginInterface device 0 info: Launching kernel
+  // CHECK-SAME: __omp_offloading_{{.*}} with [64,16,1] blocks and [32,4,2]
+  // CHECK-SAME: threads in BARE mode
+  nblocks[0] = nblocks[1] = nblocks[2] = nthreads[0] = nthreads[1] =
+      nthreads[2] = 0;
+#pragma omp target teams ompx_bare num_teams(64, 16, 1) thread_limit(32, 4, 2) \
+    map(tofrom : nblocks, nthreads)
+  {
+    get_gridsizes(nblocks, nthreads);
+  }
+  // CHECK: nblocks: 64 16 1, nthreads: 32 4 2
+  fprintf(stderr, "nblocks: %d %d %d, nthreads: %d %d %d\n", nblocks[0],
+          nblocks[1], nblocks[2], nthreads[0], nthreads[1], nthreads[2]);
+
+  // CHECK: PluginInterface device 0 info: Launching kernel
+  // CHECK-SAME: __omp_offloading_{{.*}} with [64,16,8] blocks and [32,1,1]
+  // CHECK-SAME: threads in BARE mode
+  nblocks[0] = nblocks[1] = nblocks[2] = nthreads[0] = nthreads[1] =
+      nthreads[2] = 0;
+#pragma omp target teams ompx_bare num_teams(64, 16, 8) thread_limit(32)       \
+    map(tofrom : nblocks, nthreads)
+  {
+    get_gridsizes(nblocks, nthreads);
+  }
+  // CHECK: nblocks: 64 16 8, nthreads: 32 1 1
+  fprintf(stderr, "nblocks: %d %d %d, nthreads: %d %d %d\n", nblocks[0],
+          nblocks[1], nblocks[2], nthreads[0], nthreads[1], nthreads[2]);
+
+  // CHECK: PluginInterface device 0 info: Launching kernel
+  // CHECK-SAME: __omp_offloading_{{.*}} with [64,16,8] blocks and [32,4,1]
+  // CHECK-SAME: threads in BARE mode
+  nblocks[0] = nblocks[1] = nblocks[2] = nthreads[0] = nthreads[1] =
+      nthreads[2] = 0;
+#pragma omp target teams ompx_bare num_teams(64, 16, 8) thread_limit(32, 4)    \
+    map(tofrom : nblocks, nthreads)
+  {
+    get_gridsizes(nblocks, nthreads);
+  }
+  // CHECK: nblocks: 64 16 8, nthreads: 32 4 1
+  fprintf(stderr, "nblocks: %d %d %d, nthreads: %d %d %d\n", nblocks[0],
+          nblocks[1], nblocks[2], nthreads[0], nthreads[1], nthreads[2]);
+
+  // CHECK: PluginInterface device 0 info: Launching kernel
+  // CHECK-SAME: __omp_offloading_{{.*}} with [64,16,8] blocks and [32,4,2]
+  // CHECK-SAME: threads in BARE mode
+  nblocks[0] = nblocks[1] = nblocks[2] = nthreads[0] = nthreads[1] =
+      nthreads[2] = 0;
+#pragma omp target teams ompx_bare num_teams(64, 16, 8) thread_limit(32, 4, 2) \
+    map(tofrom : nblocks, nthreads)
+  {
+    get_gridsizes(nblocks, nthreads);
+  }
+  // CHECK: nblocks: 64 16 8, nthreads: 32 4 2
+  fprintf(stderr, "nblocks: %d %d %d, nthreads: %d %d %d\n", nblocks[0],
+          nblocks[1], nblocks[2], nthreads[0], nthreads[1], nthreads[2]);
+
+  return 0;
+}
diff --git a/offload/unittests/OffloadAPI/device_code/CMakeLists.txt b/offload/unittests/OffloadAPI/device_code/CMakeLists.txt
@@ -13,6 +13,7 @@ add_offload_test_device_code(global.cpp global)
 add_offload_test_device_code(global_ctor.cpp global_ctor)
 add_offload_test_device_code(global_dtor.cpp global_dtor)
 add_offload_test_device_code(sequence.cpp sequence)
+add_offload_test_device_code(gridsize.cpp gridsize)
 
 add_custom_target(offload_device_binaries DEPENDS
     foo.bin
@@ -29,5 +30,6 @@ add_custom_target(offload_device_binaries DEPENDS
     global_ctor.bin
     global_dtor.bin
     sequence.bin
+    gridsize.bin
 )
 set(OFFLOAD_TEST_DEVICE_CODE_PATH ${CMAKE_CURRENT_BINARY_DIR} PARENT_SCOPE)
diff --git a/offload/unittests/OffloadAPI/device_code/gridsize.cpp b/offload/unittests/OffloadAPI/device_code/gridsize.cpp
@@ -0,0 +1,13 @@
+#include <gpuintrin.h>
+#include <stdint.h>
+
+extern "C" __gpu_kernel void gridsize(uint32_t *nblocks, uint32_t *nthreads) {
+  if (__gpu_block_id(0) == 0 && __gpu_block_id(1) == 0 &&
+      __gpu_block_id(2) == 0 && __gpu_thread_id(0) == 0 &&
+      __gpu_thread_id(1) == 0 && __gpu_thread_id(2) == 0) {
+    for (int i = 0; i < 3; ++i) {
+      nblocks[i] = __gpu_num_blocks(i);
+      nthreads[i] = __gpu_num_threads(i);
+    }
+  }
+}
diff --git a/offload/unittests/OffloadAPI/kernel/olLaunchKernel.cpp b/offload/unittests/OffloadAPI/kernel/olLaunchKernel.cpp

Original file line number	Diff line number	Diff line change
`@@ -68,6 +68,7 @@ unsigned llvmLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim,`
`68`	`68`	`Args.UserThreadLimit[2] = blockDim.z;`
`69`	`69`	`Args.ArgPtrs = reinterpret_cast<void **>(args);`
`70`	`70`	`Args.Flags.IsCUDA = true;`
	`71`	`+ Args.Flags.StrictBlocksAndThreads = true;`
`71`	`72`	`return __tgt_target_kernel(nullptr, 0, gridDim.x, blockDim.x, func, &Args);`
`72`	`73`	`}`
`73`	`74`	`}`