[libomptarget] Build a minimal deviceRTL for amdgcn

[JonChesterfield]

Summary

[libomptarget] Build a minimal deviceRTL for amdgcn

Repeat of D70414, with an include path fixed. Diff for sanity checking.

The CMakeLists.txt file is functionally identical to the one used in the aomp fork.
Whitespace changes were made based on nvptx/CMakeLists.txt, plus the
copyright notice updated to match (Greg was the original author so would
like his sign off on that here).

This change will build a small subset of the deviceRTL if an appropriate toolchain is
available, e.g. a local install of rocm. Support.h is moved from nvptx as a dependency
of debug.h.

[ABataev]

It builds, but I have questions about this patch

openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt

17–19

We can't build the runtime with regular clang?

openmp/libomptarget/deviceRTLs/amdgcn/src/device_environment.h

18–26 ↗

(On Diff #231946)

Agan this thing. Can we make it common if it is really required? Again, why do we need these new fields? Could you remove it from the current patch?

[JonChesterfield]

Glad it builds. Replies to comments inline

openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt

17–19

Sure, it builds fine with trunk. I'm using a local build of clang 9.0 on one machine and the openmp/fortran fork on a second.

The find package seems to be a list of anywhere that clang might plausibly be found. Copied straight out of our fork, happy to remove any entries you don't like the look of.

The NVPTX_CUDA_COMPILER_DIR doesn't seem especially appropriately named. I use the same compiler for nvptx and for amdgcn.

openmp/libomptarget/deviceRTLs/amdgcn/src/device_environment.h

18–26 ↗

(On Diff #231946)

It's different on nvptx and amdgcn, so not a great candidate for common. The extra fields are used to implement a couple of functions that are stubs in trunk which you weren't sure are necessary for nvptx.

[ABataev]

openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt

17–19

Can we then just use the clang (I think, you proposed this for nvcc) to build the runtime? Why shall we lookup for rocm/aomp?

openmp/libomptarget/deviceRTLs/amdgcn/src/device_environment.h

18–26 ↗

(On Diff #231946)

Maybe, it is better to exclude this from this patch? Maybe it would be good to have these fields for nvptx too or we don't need them to for amdgcn. Could you remind where we discussed this structure?

[JonChesterfield]

openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt

17–19

Which clang? CMAKE_CXX_COMPILER_DIR may resolve to gcc. Or do you mean introduce LIBOMPTARGET_AMDGCN_CUDA_COMPILER_DIR and expect the user to match that up to a clang build?

openmp/libomptarget/deviceRTLs/amdgcn/src/device_environment.h

18–26 ↗

(On Diff #231946)

Discussion was in D69424. I could drop the file from this patch, but it's included by debug.h so that'll break the amdgcn build

[JonChesterfield]

openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt

17–19

rocm/aomp is a likely location to find clang if using the rocm developer tools stack, and I think it gets installed into different places depending on whether it's from apt-get or github.

I'm a little unclear on how in tree compiler plus out of tree library (e.g. libc, opencl runtime, or the cuda toolchain) is supposed to work logistically.

[ABataev]

openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt

17–19

Freshly built clang which is used to build .bc library for nvptx. Plus, if need it in some cases, I think it would be good to have a cmake configuration option for the user so he could point to his own installation of AOMP. Thoughts?

openmp/libomptarget/deviceRTLs/amdgcn/src/device_environment.h

18–26 ↗

(On Diff #231946)

Ah, found it. Again, the same question. These fields are required to add 2 new omp_... functions. The question is the same. Are they required by the standard? If yes, we must have the same fields for nvptx. Otherwise, I don't think we need these fields at all.

[jdoerfert]

openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt

17–19

Take the cmake flag that is used to build the .bc file of the library.

[JonChesterfield]

openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt

17–19

There's some discussion on another diff about building libomp independently of the compiler so there may not be a freshly built clang available.

I've just had a go at tracing the logic for deciding what directory will be used for clang and it appears to be LLVM_INSTALL_PREFIX if set, otherwise LLVM_BUILD_BINARY_DIR, where find_package looks in:

• the AOMP environment variable (I think our tests use this)
• various install places
• wherever the nvptx compiler is (I'm OK with dropping this)
• the CXX compiler dir (one of my local builds uses this)

It's a bit of a mess but looks like canonical cmake to me.

openmp/libomptarget/deviceRTLs/amdgcn/src/device_environment.h

18–26 ↗

(On Diff #231946)

"When called from a target region the effect of this routing is unspecified"

So nvptx returning zero is fine, and andgcn returning a meaningful value is also fine, as far as the standard is concerned.

[ABataev]

openmp/libomptarget/deviceRTLs/amdgcn/src/device_environment.h

18–26 ↗

(On Diff #231946)

Why do we need to store these values? And if we want these functions to return some meaningful values, better to implement the same behavior for both nvptx and amdgcn.

[JonChesterfield]

openmp/libomptarget/deviceRTLs/amdgcn/src/device_environment.h

18–26 ↗

(On Diff #231946)

The plugin knows how many devices are available, so it writes the values to memory on the device, such that the device can answer the query accurately. I want the function to return a meaningful value on amdgcn as a quality of implementation point - it doesn't directly affect me whether nvptx supports the same. It'll cost a few bytes of memory, in difficult to optimise out fashion.

[ABataev]

openmp/libomptarget/deviceRTLs/amdgcn/src/device_environment.h

18–26 ↗

(On Diff #231946)

WE need to support the same behavior on all platforms.

[JonChesterfield]

openmp/libomptarget/deviceRTLs/amdgcn/src/device_environment.h

18–26 ↗

(On Diff #231946)

Who is we in this context, and do you consider unconditionally returning zero everywhere to be better than returning meaningful values on some targets and zero on others?

[ABataev]

openmp/libomptarget/deviceRTLs/amdgcn/src/device_environment.h

18–26 ↗

(On Diff #231946)

The behavior of the runtime library must be the same on all platforms.

[JonChesterfield]

openmp/libomptarget/deviceRTLs/amdgcn/src/device_environment.h

18–26 ↗

(On Diff #231946)

I don't think that's the right requirement. The runtime library is correct on a given target if it meets the openmp standard. If one target fails to meet a part of the standard, or chooses to exceed it, that should neither hobble nor demand a comparable extension from all other targets.

[ABataev]

openmp/libomptarget/deviceRTLs/amdgcn/src/device_environment.h

18–26 ↗

(On Diff #231946)

No, sorry, I can't agree with you here. The same library must behave the same way everywhere, on each target, on each platform. Supporting different behavior of the same library on different platforms is the bad design and leads to a lot of pain with support.

[JonChesterfield]

openmp/libomptarget/deviceRTLs/amdgcn/src/device_environment.h

18–26 ↗

(On Diff #231946)

It's easier for the user if changing platform causes no change in behaviour, sure. But that also means the user can only use a feature that is supported on all systems.

This is the advantage of implementation defined behaviour. People can write portable code, winning easy porting at the cost of performance. Or they can depend on features that are only available on some systems, knowing roughly the price paid for that.

Both are reasonable positions for developers to take. I don't think we should mandate portable application code, especially given the popularity of openmp on HPC where code may only ever run on one machine.

[ABataev]

openmp/libomptarget/deviceRTLs/amdgcn/src/device_environment.h

18–26 ↗

(On Diff #231946)

No, not reasonable at all. I insist that the target dependent code must contain only really target dependent code, the functiinality must be the same on all platforms.

[JonChesterfield]

openmp/libomptarget/deviceRTLs/amdgcn/src/device_environment.h

18–26 ↗

(On Diff #231946)

In which direction would you like this one resolved? Force all targets to return zero or enhance nvptx and future targets beyond the requirements of the standard?

On the other hand, take the openmp locking code. On trunk, as far as I can tell, this is functionally correct on volta and deadlocks on earlier nvidia hardware. Imagine for a moment that we are unable to make it work on non-volta cards, would the fix be to deadlock everywhere or to drop support for non-volta cards? Or to make the locks a nop, a serious QOI degradation on volta.

That one is of particular interest as it is unclear whether amdgcn is able to support unstructured locks.

[ABataev]

openmp/libomptarget/deviceRTLs/amdgcn/src/device_environment.h

18–26 ↗

(On Diff #231946)

Does not matter. I would impelement it to have better performance but you can reuse the original amdgcn solution.

As to locks, they works on all gpus, like pascal, volta, etc. And this code is platform-dependent. If the target does not support it, it is hardware limitation. In your case the code does not have any platform specific dependencies.

[JonChesterfield]

• drop a couple of fields

[JonChesterfield]

openmp/libomptarget/deviceRTLs/amdgcn/src/device_environment.h

18–26 ↗

(On Diff #231946)

I've deleted the extra fields from this patch. Given amdgcn doesn't have a libcall implementation yet the discussion about implementation divergence is arguably premature.

I don't understand your requirements here. Are hardware limitations acceptable as a reason for different behaviour between systems? E.g. floating point arithmetic will inevitably return different values on different architectures when compiled with fast-math.

[ABataev]

I don't understand your requirements here. Are hardware limitations acceptable as a reason for different behaviour between systems? E.g. floating point arithmetic will inevitably return different values on different architectures when compiled with fast-math.

Each particular case must be investigated individually. Precision is platform dependent here and thus it is ok if the results may vary.

openmp/libomptarget/deviceRTLs/amdgcn/src/device_environment.h

18 ↗

(On Diff #232010)

I think this can be moved to the common part of the code since the structure is absolutely the same.

[JonChesterfield]

In D70971#1768180, @ABataev wrote:

I don't understand your requirements here. Are hardware limitations acceptable as a reason for different behaviour between systems? E.g. floating point arithmetic will inevitably return different values on different architectures when compiled with fast-math.

Each particular case must be investigated individually. Precision is platform dependent here and thus it is ok if the results may vary.

Case by case I can deal with. Thanks for clarifying.

openmp/libomptarget/deviceRTLs/amdgcn/src/device_environment.h

18 ↗

(On Diff #232010)

It's only the same because I've deleted the fields you disliked. This interface between plugin and deviceRTL is unlikely to be uniform across architectures and is known to be different for nvptx and amdgcn.

[ABataev]

openmp/libomptarget/deviceRTLs/amdgcn/src/device_environment.h

18 ↗

(On Diff #232010)

Currently they are the same. So, we don't need to copy entities and can use the common one.

[JonChesterfield]

• merge temporarily equivalent headers as requested

[JonChesterfield]

• use device macro

[jdoerfert]

@ABataev Accept when you are ready, don't wait for me.

[ABataev]

LG

openmp/libomptarget/deviceRTLs/CMakeLists.txt

 ##===----------------------------------------------------------------------===##
 #
 # Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 # See https://llvm.org/LICENSE.txt for license information.
 # SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 #
 # ##===----------------------------------------------------------------------===##
 #
-# Build a device RTL for each available machine available.
+# Build a device RTL for each available machine.
 #
 ##===----------------------------------------------------------------------===##
 
+add_subdirectory(amdgcn)
 add_subdirectory(nvptx)

openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt

+##===----------------------------------------------------------------------===##
+#
+# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+# See https://llvm.org/LICENSE.txt for license information.
+# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+#
+##===----------------------------------------------------------------------===##
+#
+# Build the AMDGCN Device RTL if the ROCM tools are available
+#
+##===----------------------------------------------------------------------===##
+
+find_package(LLVM QUIET CONFIG
+  PATHS
+  $ENV{AOMP}
+  $ENV{HOME}/rocm/aomp
+  /opt/rocm/aomp
+  /usr/lib/rocm/aomp
+  ${LIBOMPTARGET_NVPTX_CUDA_COMPILER_DIR}
+  ${LIBOMPTARGET_NVPTX_CUDA_LINKER_DIR}
+  ${CMAKE_CXX_COMPILER_DIR}
+  NO_DEFAULT_PATH)
+
+if (LLVM_DIR)
+  libomptarget_say("Found LLVM ${LLVM_PACKAGE_VERSION}. Configure: ${LLVM_DIR}/LLVMConfig.cmake")
+else()
+  libomptarget_say("Not building AMDGCN device RTL: AOMP not found")
+  return()
+endif()
+
+set(AOMP_INSTALL_PREFIX ${LLVM_INSTALL_PREFIX})
+
+if (AOMP_INSTALL_PREFIX)
+  set(AOMP_BINDIR ${AOMP_INSTALL_PREFIX}/bin)
+else()
+  set(AOMP_BINDIR ${LLVM_BUILD_BINARY_DIR}/bin)
+endif()
+
+libomptarget_say("Building AMDGCN device RTL. LLVM_COMPILER_PATH=${AOMP_BINDIR}")
+
+project(omptarget-amdgcn)
+
+add_custom_target(omptarget-amdgcn ALL)
+
+#optimization level
+set(optimization_level 2)
+
+# Activate RTL message dumps if requested by the user.
+if(LIBOMPTARGET_NVPTX_DEBUG)
+  set(CUDA_DEBUG -DOMPTARGET_NVPTX_DEBUG=-1)
+endif()
+
+get_filename_component(devicertl_base_directory
+  ${CMAKE_CURRENT_SOURCE_DIR}
+  DIRECTORY)
+
+set(cuda_sources
+  ${devicertl_base_directory}/common/src/cancel.cu
+  ${devicertl_base_directory}/common/src/critical.cu)
+
+set(h_files
+  ${CMAKE_CURRENT_SOURCE_DIR}/src/amdgcn_interface.h
+  ${CMAKE_CURRENT_SOURCE_DIR}/src/target_impl.h
+  ${devicertl_base_directory}/common/debug.h
+  ${devicertl_base_directory}/common/device_environment.h
+  ${devicertl_base_directory}/common/state-queue.h
+  ${devicertl_base_directory}/common/state-queuei.h
+  ${devicertl_base_directory}/common/support.h)
+
+# for both in-tree and out-of-tree build
+if (NOT CMAKE_ARCHIVE_OUTPUT_DIRECTORY)
+  set(OUTPUTDIR ${CMAKE_CURRENT_BINARY_DIR})
+else()
+  set(OUTPUTDIR ${CMAKE_ARCHIVE_OUTPUT_DIRECTORY})
+endif()
+
+# create libraries
+set(mcpus gfx700 gfx701 gfx801 gfx803 gfx900)
+if (DEFINED LIBOMPTARGET_AMDGCN_GFXLIST)
+  set(mcpus ${LIBOMPTARGET_AMDGCN_GFXLIST})
+endif()
+
+macro(add_cuda_bc_library)
+  set(cu_cmd ${AOMP_BINDIR}/clang++
+    -std=c++11
+    -fcuda-rdc
+    -fvisibility=default
+    --cuda-device-only
+    -Wno-unused-value
+    -x hip
+    -O${optimization_level}
+    --cuda-gpu-arch=${mcpu}
+    ${CUDA_DEBUG}
+    -I${CMAKE_CURRENT_SOURCE_DIR}/src
+    -I${devicertl_base_directory})
+
+  set(bc1_files)
+
+  foreach(file ${ARGN})
+    get_filename_component(fname ${file} NAME_WE)
+    set(bc1_filename ${fname}.${mcpu}.bc)
+
+    add_custom_command(
+      OUTPUT ${bc1_filename}
+      COMMAND ${cu_cmd} ${file} -o ${bc1_filename}
+      DEPENDS ${file} ${h_files})
+
+    list(APPEND bc1_files ${bc1_filename})
+  endforeach()
+
+  add_custom_command(
+    OUTPUT linkout.cuda.${mcpu}.bc
+    COMMAND ${AOMP_BINDIR}/llvm-link ${bc1_files} -o linkout.cuda.${mcpu}.bc
+    DEPENDS ${bc1_files})
+
+  list(APPEND bc_files linkout.cuda.${mcpu}.bc)
+endmacro()
+
+set(libname "omptarget-amdgcn")
+
+foreach(mcpu ${mcpus})
+  set(bc_files)
+  add_cuda_bc_library(${cuda_sources})
+
+  set(bc_libname lib${libname}-${mcpu}.bc)
+  add_custom_command(
+    OUTPUT ${bc_libname}
+    COMMAND ${AOMP_BINDIR}/llvm-link ${bc_files} | ${AOMP_BINDIR}/opt --always-inline -o ${OUTPUTDIR}/${bc_libname}
+    DEPENDS ${bc_files})
+
+  add_custom_target(lib${libname}-${mcpu} ALL DEPENDS ${bc_libname})
+
+  install(FILES ${OUTPUTDIR}/${bc_libname}
+     DESTINATION "${OPENMP_INSTALL_LIBDIR}/libdevice"
+  )
+endforeach()

openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h

 EXTERN uint64_t __ballot64(int predicate);
 // initialized with a 64-bit mask with bits set in positions less than the
 // thread's lane number in the warp
 EXTERN uint64_t __lanemask_lt();
 // initialized with a 64-bit mask with bits set in positions greater than the
 // thread's lane number in the warp
 EXTERN uint64_t __lanemask_gt();
 
-EXTERN void llvm_amdgcn_s_barrier();
-
 // CU id
 EXTERN unsigned __smid();
 
 INLINE void __kmpc_impl_unpack(uint64_t val, uint32_t &lo, uint32_t &hi) {
   lo = (uint32_t)(val & UINT64_C(0x00000000FFFFFFFF));
   hi = (uint32_t)((val & UINT64_C(0xFFFFFFFF00000000)) >> 32);
 }
 

 INLINE __kmpc_impl_lanemask_t __kmpc_impl_lanemask_gt() {
   return __lanemask_gt();
 }
 
 INLINE uint32_t __kmpc_impl_smid() {
   return __smid();
 }
 
-INLINE uint64_t __kmpc_impl_ffs(uint64_t x) { return __ffsll(x); }
+INLINE uint64_t __kmpc_impl_ffs(uint64_t x) { return __builtin_ffsl(x); }
 
-INLINE uint64_t __kmpc_impl_popc(uint64_t x) { return __popcll(x); }
+INLINE uint64_t __kmpc_impl_popc(uint64_t x) { return __builtin_popcountl(x); }
 
 INLINE __kmpc_impl_lanemask_t __kmpc_impl_activemask() {
   return __ballot64(1);
 }
 
-INLINE int32_t __kmpc_impl_shfl_sync(__kmpc_impl_lanemask_t, int32_t Var,
-                                     int32_t SrcLane) {
-  return __shfl(Var, SrcLane, WARPSIZE);
-}
+EXTERN int32_t __kmpc_impl_shfl_sync(__kmpc_impl_lanemask_t, int32_t Var,
+                                     int32_t SrcLane);
 
-INLINE int32_t __kmpc_impl_shfl_down_sync(__kmpc_impl_lanemask_t, int32_t Var,
-                                          uint32_t Delta, int32_t Width) {
-  return __shfl_down(Var, Delta, Width);
-}
+EXTERN int32_t __kmpc_impl_shfl_down_sync(__kmpc_impl_lanemask_t, int32_t Var,
+                                          uint32_t Delta, int32_t Width);
 
-INLINE void __kmpc_impl_syncthreads() { llvm_amdgcn_s_barrier(); }
+INLINE void __kmpc_impl_syncthreads() { __builtin_amdgcn_s_barrier(); }
 
 INLINE void __kmpc_impl_named_sync(int barrier, uint32_t num_threads) {
   // we have protected the master warp from releasing from its barrier
   // due to a full workgroup barrier in the middle of a work function.
   // So it is ok to issue a full workgroup barrier here.
   __builtin_amdgcn_s_barrier();
 }
 
+// DEVICE versions of part of libc
+extern "C" {
+DEVICE __attribute__((noreturn)) void
+__assertfail(const char *, const char *, unsigned, const char *, size_t);
+INLINE static void __assert_fail(const char *__message, const char *__file,
+                                 unsigned int __line, const char *__function) {
+  __assertfail(__message, __file, __line, __function, sizeof(char));
+}
+DEVICE int printf(const char *, ...);
+}
+
 #endif

openmp/libomptarget/deviceRTLs/common/debug.h

 //   WARNING(flag, str, args)      : if warning flag is on, print the warning
 //   WON(flag)                     : return true if warning flag is on
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef _OMPTARGET_NVPTX_DEBUG_H_
 #define _OMPTARGET_NVPTX_DEBUG_H_
 
-#include "device_environment.h"
+#include "common/device_environment.h"
 
 ////////////////////////////////////////////////////////////////////////////////
 // set desired level of debugging
 ////////////////////////////////////////////////////////////////////////////////
 
 #define LD_SET_NONE 0ULL /* none */
 #define LD_SET_ALL -1ULL /* all */
 

 #endif
 
 ////////////////////////////////////////////////////////////////////////////////
 // implemtation for debug
 ////////////////////////////////////////////////////////////////////////////////
 
 #if OMPTARGET_NVPTX_DEBUG || OMPTARGET_NVPTX_TEST || OMPTARGET_NVPTX_WARNING
 #include <stdio.h>
-#include "support.h"
+#include "common/support.h"
 
 template <typename... Arguments>
 NOINLINE static void log(const char *fmt, Arguments... parameters) {
   printf(fmt, (int)GetBlockIdInKernel(), (int)GetThreadIdInBlock(),
          (int)GetWarpId(), (int)GetLaneId(), parameters...);
 }
 
 #endif

openmp/libomptarget/deviceRTLs/common/device_environment.h

 #define _OMPTARGET_DEVICE_ENVIRONMENT_H_
 
 #include "target_impl.h"
 
 struct omptarget_device_environmentTy {
   int32_t debug_level;
 };
 
-extern __device__ omptarget_device_environmentTy omptarget_device_environment;
+extern DEVICE omptarget_device_environmentTy omptarget_device_environment;
 
 #endif

openmp/libomptarget/deviceRTLs/nvptx/src/omp_data.cu

 //===------------ omp_data.cu - NVPTX OpenMP GPU objects --------- CUDA -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file contains the data objects used on the GPU device.
 //
 //===----------------------------------------------------------------------===//
 
 #include "omptarget-nvptx.h"
-#include "device_environment.h"
+#include "common/device_environment.h"
 
 ////////////////////////////////////////////////////////////////////////////////
 // global device environment
 ////////////////////////////////////////////////////////////////////////////////
 
 __device__ omptarget_device_environmentTy omptarget_device_environment;
 
 ////////////////////////////////////////////////////////////////////////////////

openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h

 #include <math.h>
 #include <stdlib.h>
 
 // local includes
 #include "target_impl.h"
 #include "common/debug.h"     // debug
 #include "interface.h" // interfaces with omp, compiler, and user
 #include "common/state-queue.h"
-#include "support.h"
+#include "common/support.h"
 
 #define OMPTARGET_NVPTX_VERSION 1.1
 
 // used by the library for the interface with the app
 #define DISPATCH_FINISHED 0
 #define DISPATCH_NOTFINISHED 1
 
 // used by dynamic scheduling

openmp/libomptarget/deviceRTLs/nvptx/src/support.cu

 //===--------- support.cu - NVPTX OpenMP support functions ------- CUDA -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // Wrapper implementation to some functions natively supported by the GPU.
 //
 //===----------------------------------------------------------------------===//
 
-#include "support.h"
+#include "common/support.h"
 #include "common/debug.h"
 #include "omptarget-nvptx.h"
 
 ////////////////////////////////////////////////////////////////////////////////
 // Execution Parameters
 ////////////////////////////////////////////////////////////////////////////////
 
 DEVICE void setExecutionParameters(ExecutionMode EMode, RuntimeMode RMode) {