ComputeNonbondedMICKernel.h

Go to the documentation of this file.

#ifndef COMPUTE_NONBONDED_MIC_KERNEL_H
#define COMPUTE_NONBONDED_MIC_KERNEL_H


// If this is a MIC build for NAMD, then include the contents of this file
#ifdef NAMD_MIC


// General Code Description
//
//    The files beginning with "ComputeNonbondedMIC" are MIC-specific files that
// are involved in supporting MIC within the NAMD application.  Below is a
// general description of that support.
//
// Startup:
//
//    During the startup phase of the application, there are several functions
// that are called, once per device.  First, the function mic_check is used to
// detect if the MIC devices that are to be used available to the application
// (by default, all devices are used, but the "+devices" command line parameter
// can be used to select specific devices).  Then, the mic_init_device function
// is then used to initialize each of the MIC devices.
//
//    Also during startup, some constant tables/variables that will be used
// during the steady state of the application are transfered to the device.
// These functions, again called once per device, include...
//  - mic_bind_table_four : Nonbonded table (includes several sub-tables)
//  - mic_bind_lj_table   : Lennard-Jones (LJ) table
//  - mic_bind_exclusions : Table of exclusion flags
//  - mic_bind_constants  : Several required constant scalar values
//
// Steady State (timesteps executing):
//
//    During the steady state of the application, when timesteps are executing
// one after the next, several functions are used to transfer data to the MIC
// device, including...
//  - mic_bind_patch_pairs : Updates buffers that only need to be updated on
//      the device(s) periodically, usually because of some major event such
//      as atom migration or load balancing (computes shifting between PEs).
//  - mic_nonbonded_forces : Used to initiate computation on the MIC device.
//      Various buffers that are required to be transfered each timestep,
//      such as atoms (input: atom position info) or forces (output: atom
//      force info), along with buffers that are only periodically transfered,
//      such as atom_params (input: misc. atom info).  The body of this function
//      includes an offload pragma that is broken up into two stages.  The first
//      is force computation where the computes (patch_pairs) are executed in
//      parallel with one another.  The second is force reduction, where the
//      forces are reduced in parallel (per patch), accumulating output from
//      each of the computes that contributes to each of the patches.  This
//      offload pragma is asynchronous so that the host can continue to work.
//  - mic_check_[remote|local]_kernel_complete : A function that polls for the
//      completion of the main computation offload pragma.
//
// Force Computation (during steady state):
// 
//    As part of the force computation (first stage of computation on device),
// several functions are used to "process" (calculate forces for) each compute
// object.  These functions are basically the same (at least conceptually, but
// have slight variations from one another based on the type of compute being
// "processed" or the conditions for the timestep.  These functions include...
//  - Base Names:
//    - calc_pair : compute with two different input patches
//    - calc_self : compute with a single input patch (i.e. patch1 == patch2)
//  - Qualifiers:
//    - "_energy" : calculate additional energy info during energy timesteps
//    - "_fullelect" : calculate additional slow_forces during PME timesteps
//  - Examples:
//    - calc_pair : basic caclulations for a pair compute
//    - calc_self_fullelect : basic + slow_force calculations for a self compute
//
//    These force computation functions are generated by including the header
// file "ComputeNonbondedMICKernelBase.h" into "ComputeNonbondedMICKernel.C"
// multiple times, each time with different "macro flags" setup.  The general
// include structure is as follows (indentation indicates inclusion)...
//  - ComputeNonbondedMICKernelBase.h
//    - declares compute function body, along with function setup/cleanup
//    - periodically does pairlist generation
//    - includes ComputeNonbondedMICKernelBase2.h (one per type of interation,
//      such as NORMAL, EXCLUDED, and MODIFIED).
//      - refines pairlist to only include interactions within cutoff for the
//        given timestep (if pairlist refinement is enabled)
//      - includes ComputeNonbondedMICKernelBase2_XXXXX.h, where XXXXX is
//        either "handcode" or "scalar", depending on the build's configuration
//        - force computation loop, that processes each entry in the pairlist
//          for the given type of interaction
//
//    For the most part, the MIC code is based on the host code.  The main
// difference is the restructuring of the pairlists, and thus the force
// computation loops.  For the MIC code, each pairlist entry is a combination
// of both an "i" (outer-loop atom list) index and "j" (inner-loop atom list)
// index that represents a pair of atoms that need to be checked this timestep
// for their exact distance, and if close enough, interact with one another.
// Correspondingly, the force computation loop is only a single loop, rather
// than two nested loops (inner + outer), that iterates over pairlist entries.
// As such, each compute has only one pairlist (per interaction type), rather
// than many smaller pairlists for each outer-loop atom (one per interaction
// type, per outer-loop atom).  There is still one force computation loop,
// setup in ComputeNonbondedMICKernelBase2.h, per interaction type.
//
// Cleanup:
//
//    The cleanup function mic_free_device should be called, once per device, to
// cleanup the various data structures used by the device (and the associated
// data structures on the host) at the end of the overall simulation.
//


#include <stdlib.h>
#include <stdio.h>
#include <offload.h>
#include <string.h>
#include <stdint.h>
#ifndef __MIC__
  #include "charm++.h"
#endif


// NAMD_MIC Configuration Macros

// There are three main ways in which the force computation loop can be compiled,
//   including those listed in the table below.
//
//                VERSION   MIC_HANDCODE_FORCE   MIC_PAD_PLGEN   ALSO RECOMMENDED ("set" these, others "unset")
//    -------------------   ------------------   -------------   ------------------------------------------------------------------------------
//                 scalar                   0               0    
//    compiler-vectorized                   0               1    MIC_HANDCODE_PLGEN(1) (until auto vectorization of pairlist generation loop is working)
//        hand-vectorized                   1             n/a    MIC_HANDCODE_FORCE_CALCR2TABLE(1), MIC_HANDCODE_FORCE_USEGATHER(1), MIC_HANDCODE_FORCE_SOA_VS_AOS(1)
//              --- for all versions of the code ---             MULTIPLE_THREADS(1), MIC_SORT_ATOMS(1), MIC_HANDCODE_FORCE_SINGLE(1)

// Flags to enable/disable the use of handcoding
#define MIC_HANDCODE_FORCE                     ( 1 ) // 1 // 0 - disable, !0 - enable hand-coded force generation loop
  #define MIC_HANDCODE_FORCE_PFDIST            ( 0 )  // (currently broken!) 0 - disable, >=8 - enable software prefetch (value = prefetch distance)
  #define MIC_HANDCODE_FORCE_USEGATHER         ( 1 ) // 1 //
  #define MIC_HANDCODE_FORCE_USEGATHER_NBTBL   ( 0 )
  #define MIC_HANDCODE_FORCE_LOAD_VDW_UPFRONT  ( 0 )  // 0 - disable, !0 - enable loading of vdw type information prior to the cutoff check within the force loops
  #define MIC_HANDCODE_FORCE_COMBINE_FORCES    ( 1 ) // 1 // 0 - disable, !0 - enable a single set of force updates (short-fast and full), rather than two separate updates in the force loop body (allowing indexes, etc. to be reused for both)

// Flags that were originally part of handcoding, but are now more general
#define MIC_HANDCODE_FORCE_SINGLE            ( 1 ) // 1 // 0 - disable, !0 - enable use of single precision to calculate force magnitude in force loops (forces still accumulated in double precision)
#define MIC_HANDCODE_FORCE_CALCR2TABLE       ( 1 ) // 1 // 0 - disable, !0 - enable r2_table value inline calculation (vs r2_table value loads)
#define MIC_HANDCODE_FORCE_SOA_VS_AOS        ( 1 ) // 1 // 0 - SoA, !0 - AoS

#define MIC_SORT_ATOMS                    ( 1 ) // 1 // 0 - disable, !0 - enable atom sorting within each patch
#define MIC_SORT_LISTS                    ( 0 )  // (currently not working, leave at 0)

// Pairlist generation controls
#define MIC_HANDCODE_PLGEN                ( 1 ) // 1 // 0 - disable, !0 - enable hand-coded pairlist generation
#define MIC_TILE_PLGEN                    ( 0 )  // 0 - disable, !0 - enable pairlist tiling (value = number of tiles to create for each atom list)
#define MIC_CONDITION_NORMAL              ( 0 )  // 0 - disable, !0 - enable normal pairlist "conditioning" (enabling this flag will cause the normal pairlist to be separated into two sets, such that pairs that are likely to be within cutoff are grouped together and pairs that are less likely to be within cutoff... so vector lanes tend to all active or all not-active)
#define MIC_PAD_PLGEN                     ( 1 ) // 1 // 0 - disable, !0 - enable pairlist padding

// Flag to enable/disable the use of multiple threads within the force computation
#define MULTIPLE_THREADS                  ( 1 ) // 1 // 0 - disable, !0 - use multiple threads when processing compute objects (a.k.a. patch pairs)

#define MIC_SORT_COMPUTES                 ( 1 )

#define MIC_SYNC_INPUT                    ( 1 )
#define MIC_SYNC_OUTPUT                   ( 1 )

#define MIC_SUBMIT_ATOMS_ON_ARRIVAL       ( 0 )

// Flags used to enable/disable and control pairlist refinement
#define REFINE_PAIRLISTS                  ( 0 )  // 0 - disable, !0 - enable pairlist refinement
  #define REFINE_PAIRLIST_HANDCODE        ( 0 )  // 0 - disable, !0 - use hand-coded version of the pairlist refinement code
  #define REFINE_PAIRLISTS_XYZ            ( 0 )  // 0 - disable, !0 - save dX (delta X), dY, and dZ values along with r2 values as part of the pairlist refinement output process

#define MIC_PREFETCH_DISTANCE             ( 0 )
  #define MIC_PREFETCH_HINT               ( 3 )

// Flag used to enable/disable various checks within and related to the MIC computation code
#define MIC_FULL_CHECK                    ( 1 )  // 0 - disable, !0 - enable various checks use to help verify correct operation, but that are not required to perform the calculations themselves
#define MIC_EXCL_CHECKSUM_FULL            ( 1 )  // NOTE: Mutually exclusive with MIC_EXCL_CHECKSUM
#define MIC_DATA_STRUCT_VERIFY            ( 0 )
  #define MIC_DATA_STRUCT_VERIFY_KERNELS  ( 0 )
#define MIC_DEVICE_FPRINTF                ( 0 )  // Print some debug info from the MICs directly to a device file on the MICs (for debug purposes)
  #define MIC_DEVICE_FPRINTF_REOPEN_FREQ  ( 0 )  // If >0, will cause the code to periodically reopen (and thus truncate) the device file, helping to limit the amount of memory it takes up
#define MIC_DUMP_COMPUTE_MAPPING          ( 0 )  // 0: do nothing, !0: dump directToDevice info from ComputeMap
#define MIC_TRACK_DEVICE_MEM_USAGE        ( 0 )  // 0: do nothing, !0: print memusage on the device every 100 timesteps
#define MIC_HEARTBEAT                     ( 0 )

// Alignment used for various buffers, values, etc., in bytes
#define MIC_ALIGN                         ( 64 ) //4 * 1024 )  // NOTE : Must be > 0 and must be a multiple of the cacheline size (in bytes)

#define MIC_MAX_DEVICES_PER_NODE          ( 16 )

#define MIC_PRINT_CONFIG                  ( 0 )  // 0 - disable, !0 - enable the printing of configuration information (macros defined in this file) at the start of a simulation (TODO | NOTE: This is meant as a reference for repeating performance experiments... not required for production use... disable at some point)

#define MIC_VERBOSE_HVD_LDB               ( 0 )  // 0 - disable, !0 - print extra info about HvD LDB

#define MIC_DEBUG                         ( 0 )  // 0 - disable, !0 - enable extra debugging output to be printed to files (1 file per PE) via MICP statements

#define NUM_PAIRLIST_TYPES                ( 4 )  // Number of pairlist types (i.e. multiplier used for mallocs, etc.)
  #define PL_NORM_INDEX                   ( 0 )  //   Index for normal pairlist
  #define PL_EXCL_INDEX                   ( 1 )  //   Index for excluded pairlist
  #define PL_MOD_INDEX                    ( 2 )  //   Index for modified pairlist
  #define PL_NORMHI_INDEX                 ( 3 )  //   Index for the "normal high" pairlist (temporary pairlist used for entries that are far apart, used to separate the sets of entries, close and far, used in normal pairlist conditioning... see MIC_CONDITION_NORMAL)

// Projections tracing
#define MIC_TRACING                       ( 1 )  // 0 - disable, !0 - enable tracing information to be collected on the host
  #define MIC_DEVICE_TRACING              ( 1 )  // 0 - disable, !0 - enable the collection of extra tracing information on the device itself (only high-level events, like the total time for all computes, for all patches, etc.) (requires MIC_TRACING to be set)
  #define MIC_DEVICE_TRACING_DETAILED     ( 0 )  // 0 - disable, !0 - enable the collection of more tracing information on the device, including start and stop times for each task (requires MIC_DEVICE_TRACING to be sent)
  #define MIC_TRACING_EVENT_BASE          ( 10000 )
  #define MIC_EVENT_OFFLOAD_REMOTE        ( MIC_TRACING_EVENT_BASE +  0 )
  #define MIC_EVENT_OFFLOAD_LOCAL         ( MIC_TRACING_EVENT_BASE +  1 )
  #define MIC_EVENT_OFFLOAD_PRAGMA        ( MIC_TRACING_EVENT_BASE +  2 )
  #define MIC_EVENT_OFFLOAD_POLLSET       ( MIC_TRACING_EVENT_BASE +  3 )
  #define MIC_EVENT_FUNC_DOWORK           ( MIC_TRACING_EVENT_BASE +  4 )
  #define MIC_EVENT_FUNC_FINISHWORK       ( MIC_TRACING_EVENT_BASE +  5 )
  #define MIC_EVENT_ATOMS_SUBMIT          ( MIC_TRACING_EVENT_BASE +  6 )
  #define MIC_EVENT_ATOMS_TRANSFER        ( MIC_TRACING_EVENT_BASE +  7 )
  #define MIC_EVENT_ATOMS_WAIT            ( MIC_TRACING_EVENT_BASE +  8 )
  #define MIC_EVENT_SYNC_INPUT_PRAGMA     ( MIC_TRACING_EVENT_BASE +  9 )
  #define MIC_EVENT_SYNC_OUTPUT_REMOTE_PRAGMA  ( MIC_TRACING_EVENT_BASE + 10 )
  #define MIC_EVENT_SYNC_OUTPUT_LOCAL_PRAGMA   ( MIC_TRACING_EVENT_BASE + 11 )
  #define MIC_EVENT_DEVICE_COMPUTES       ( MIC_TRACING_EVENT_BASE + 12 )
  #define MIC_EVENT_DEVICE_VIRIALS        ( MIC_TRACING_EVENT_BASE + 13 )
  #define MIC_EVENT_DEVICE_PATCHES        ( MIC_TRACING_EVENT_BASE + 14 )
  #define MIC_EVENT_DEVICE_PATCH          ( MIC_TRACING_EVENT_BASE + 15 )
  #define MIC_EVENT_DEVICE_COMPUTE        ( MIC_TRACING_EVENT_BASE + 16 )
  #define MIC_TRACING_NUM_EVENTS          ( 24 )

// DMK - DEBUG
#define MIC_KERNEL_DATA_TRANSFER_STATS    (  0 )

__declspec(target(mic)) void* _mm_malloc_withPrint(size_t s, int a, char * l);
__declspec(target(mic)) void _mm_free_withPrint(void * prt);

// DMK - NOTE : Leaving these macros in as hooks for now
//#ifdef __MIC__
//  #define _MM_MALLOC_WRAPPER(s, a, l) _mm_malloc_withPrint((s), (a), (l))
//  #define _MM_FREE_WRAPPER(p) _mm_free_withPrint(p)
//  #define _MM_MALLOC_VERIFY
//#else
  #define _MM_MALLOC_WRAPPER(s, a, l) _mm_malloc((s), (a))
  #define _MM_FREE_WRAPPER(p) _mm_free(p)
  #define _MM_MALLOC_VERIFY
//#endif


// Configuration Checks/Overrides

// If not compiling for the MIC device, disable all MIC specific code
#include <stdint.h>
#if defined(__MIC__) || defined(__MIC2__)
  #include <immintrin.h>
#else
  #undef MIC_HANDCODE_FORCE
  #define MIC_HANDCODE_FORCE            (0)
  #undef MIC_HANDCODE_PLGEN
  #define MIC_HANDCODE_PLGEN            (0)
#endif

// If the MIC_FULL_CHECK flag is set, enable checks for both __ASSUME and
//   __ASSUME_ALIGNED to ensure the cases for these macros are actually true.
//   Also ensure that MIC_EXCL_CHECKSUM_FULL is enabled.
#if MIC_FULL_CHECK != 0
  #define __FULL_CHECK(X) X
  #define __ASSERT(c) assert(c)
  #define CHECK_ASSUME          (1)
  #define CHECK_ASSUME_ALIGNED  (1)
  #if (MIC_EXCL_CHECKSUM_FULL == 0)
    #warning "MIC_EXCL_CHECKSUM_FULL enabled (because MIC_FULL_CHECK was enabled)"
  #endif
  #undef MIC_EXCL_CHECKSUM_FULL
  #define MIC_EXCL_CHECKSUM_FULL  (1)
#else
  #define __FULL_CHECK(X)
  #define __ASSERT(c)
  #define CHECK_ASSUME          (0)
  #define CHECK_ASSUME_ALIGNED  (0)
#endif

// Create a macro that can be used within the body of offload pragmas that asserts the
//   offload section was executed on a MIC device
#if defined(__MIC__)
  #define __MUST_BE_MIC
#else
  #define __MUST_BE_MIC  __ASSERT(0)
#endif

// For now, if padding the pairlist with zeros, disable pairlist tiling
#if (MIC_PAD_PLGEN != 0) && (MIC_HANDCODE_FORCE != 0)
  #if REFINE_PAIRLISTS != 0
    #warning "REFINE_PAIRLISTS disabled (because MIC_PAD_PLGEN && MIC_HANDCODE_FORCE were enabled; combination not supported yet)"
  #endif
  #undef REFINE_PAIRLISTS
  #define REFINE_PAIRLISTS  (0)
#endif

// If using the handcoded force loop with mixed precision, disable the use of pairlist refinement (not supported yet)
#if MIC_HANDCODE_FORCE_SINGLE != 0
  #if REFINE_PAIRLISTS != 0
    #warning "REFINE_PAIRLISTS disabled (because MIC_HANDCODE_FORCE_SINGLE was enabled)"
  #endif
  #undef REFINE_PAIRLISTS
  #define REFINE_PAIRLISTS  (0)
#endif

// If AOS has been enabled, ensure that loading VDW upfront has been disabled (not supported, doesn't
//   make much sense as x,y,z,q are in a different array than vdwType)
#if MIC_HANDCODE_FORCE_SOA_VS_AOS != 0
  #if MIC_HANDCODE_FORCE_LOAD_VDW_UPFRONT != 0
    #warning "MIC_HANDCODE_FORCE_LOAD_VDW_UPFRONT disabled (because MIC_HANDCODE_FORCE_SOA_VS_AOS was enabled; combination not supported)"
  #endif
  #undef MIC_HANDCODE_FORCE_LOAD_VDW_UPFRONT
  #define MIC_HANDCODE_FORCE_LOAD_VDW_UPFRONT (0)
#endif

// Check for a negative value specified as the prefetch distance (macro's value is used, so require positive)
#if MIC_PREFETCH_DISTANCE < 0
  #warning "INVALID MIC_PREFETCH_DISTANCE value (disabling)"
  #undef MIC_PREFETCH_DISTANCE
  #define MIC_PREFETCH_DISTANCE (0)
#endif

// If tracing has been enabled, ensure that the other required tracing macros are properly set
#if MIC_DEVICE_TRACING_DETAILED != 0
  #if MIC_DEVICE_TRACING == 0
    #warning "MIC_DEVICE_TRACING enabled (because MIC_DEVICE_TRACING_DETAILED was enabled)"
  #endif
  #undef MIC_DEVICE_TRACING
  #define MIC_DEVICE_TRACING  (1)
#endif
#if MIC_DEVICE_TRACING != 0
  #if MIC_TRACING == 0
    #warning "MIC_TRACING enabled (because MIC_DEVICE_TRACING was enabled)"
  #endif
  #undef MIC_TRACING
  #define MIC_TRACING  (1)
#endif


// Tracing Macros (Projections)

#if MIC_TRACING != 0

  #define MIC_TRACING_REGISTER_EVENT(str, id) traceRegisterUserEvent(str, id)

  #define MIC_TRACING_REGISTER_EVENTS \
    MIC_TRACING_REGISTER_EVENT("MIC device offload (remote)", MIC_EVENT_OFFLOAD_REMOTE); \
    MIC_TRACING_REGISTER_EVENT("MIC device offload (local)", MIC_EVENT_OFFLOAD_LOCAL); \
    MIC_TRACING_REGISTER_EVENT("MIC offload pragma", MIC_EVENT_OFFLOAD_PRAGMA); \
    MIC_TRACING_REGISTER_EVENT("MIC polling set", MIC_EVENT_OFFLOAD_POLLSET); \
    MIC_TRACING_REGISTER_EVENT("MIC ::doWork", MIC_EVENT_FUNC_DOWORK); \
    MIC_TRACING_REGISTER_EVENT("MIC ::finishWork", MIC_EVENT_FUNC_FINISHWORK); \
    MIC_TRACING_REGISTER_EVENT("MIC atom submit", MIC_EVENT_ATOMS_SUBMIT); \
    MIC_TRACING_REGISTER_EVENT("MIC atom transfer", MIC_EVENT_ATOMS_TRANSFER); \
    MIC_TRACING_REGISTER_EVENT("MIC atom wait", MIC_EVENT_ATOMS_WAIT); \
    MIC_TRACING_REGISTER_EVENT("MIC sync input pragma", MIC_EVENT_SYNC_INPUT_PRAGMA); \
    MIC_TRACING_REGISTER_EVENT("MIC sync output remote pragma", MIC_EVENT_SYNC_OUTPUT_REMOTE_PRAGMA); \
    MIC_TRACING_REGISTER_EVENT("MIC sync output local pragma", MIC_EVENT_SYNC_OUTPUT_LOCAL_PRAGMA); \
    MIC_TRACING_REGISTER_EVENT("[MIC] computes", MIC_EVENT_DEVICE_COMPUTES); \
    MIC_TRACING_REGISTER_EVENT("[MIC] virials", MIC_EVENT_DEVICE_VIRIALS); \
    MIC_TRACING_REGISTER_EVENT("[MIC] patches", MIC_EVENT_DEVICE_PATCHES); \
    MIC_TRACING_REGISTER_EVENT("[MIC] patch", MIC_EVENT_DEVICE_PATCH); \
    MIC_TRACING_REGISTER_EVENT("[MIC] compute (d0)", MIC_EVENT_DEVICE_COMPUTE + 0); \
    MIC_TRACING_REGISTER_EVENT("[MIC] compute (d1)", MIC_EVENT_DEVICE_COMPUTE + 1); \
    MIC_TRACING_REGISTER_EVENT("[MIC] compute (d2)", MIC_EVENT_DEVICE_COMPUTE + 2); \
    MIC_TRACING_REGISTER_EVENT("[MIC] compute (d3)", MIC_EVENT_DEVICE_COMPUTE + 3); \
    MIC_TRACING_REGISTER_EVENT("[MIC] compute (d4)", MIC_EVENT_DEVICE_COMPUTE + 4); \
    MIC_TRACING_REGISTER_EVENT("[MIC] compute (d5)", MIC_EVENT_DEVICE_COMPUTE + 5); \
    MIC_TRACING_REGISTER_EVENT("[MIC] compute (d6)", MIC_EVENT_DEVICE_COMPUTE + 6); \
    MIC_TRACING_REGISTER_EVENT("[MIC] compute (d7+)", MIC_EVENT_DEVICE_COMPUTE + 7); \
    \
    MIC_TRACING_REGISTER_EVENT("Pair", 11000); \
    MIC_TRACING_REGISTER_EVENT("Self", 11001); \
    MIC_TRACING_REGISTER_EVENT("PME", 11002); \
    MIC_TRACING_REGISTER_EVENT("Tuple", 11003); \
    MIC_TRACING_REGISTER_EVENT("Box Closing", 11010); \
    MIC_TRACING_REGISTER_EVENT("Positions Ready", 11020); \
    MIC_TRACING_REGISTER_EVENT("Post Run Computes", 11021); \
    MIC_TRACING_REGISTER_EVENT("Pre Run Computes", 11022); \
    MIC_TRACING_REGISTER_EVENT("Reduction Submit", 11023); \
    MIC_TRACING_REGISTER_EVENT("MIC Atom Work", 11030); \
    MIC_TRACING_REGISTER_EVENT("Patch::positionsReady", 11031); \
    MIC_TRACING_REGISTER_EVENT("Atom send", 11032); \
    MIC_TRACING_REGISTER_EVENT("DEBUG 0", 11040); \
    MIC_TRACING_REGISTER_EVENT("DEBUG 1", 11041); \
    MIC_TRACING_REGISTER_EVENT("DEBUG 2", 11042); \
    MIC_TRACING_REGISTER_EVENT("DEBUG 3", 11043); \
    MIC_TRACING_REGISTER_EVENT("DEBUG 4", 11044); \
    MIC_TRACING_REGISTER_EVENT("recvImmediateProxyData", 11046); \
    MIC_TRACING_REGISTER_EVENT("msg copy send", 11047); \
    MIC_TRACING_REGISTER_EVENT("ProxyMgr::sendProxyData", 11048); \
    MIC_TRACING_REGISTER_EVENT("MIC box close", 11050); \
    MIC_TRACING_REGISTER_EVENT("MIC reduction submit", 11051); \
    MIC_TRACING_REGISTER_EVENT("WorkDistrib compute", 11070); \
    MIC_TRACING_REGISTER_EVENT("doWork::work", 11080); \
    MIC_TRACING_REGISTER_EVENT("doWork::progress", 11081); 

  #define MIC_TRACING_RECORD(id, start, stop) traceUserBracketEvent(id, start, stop)

  extern void traceUserBracketEvent(int, double, double);
  extern double CmiWallTimer();

#else

  #define MIC_TRACING_REGISTER_EVENT(str, id)
  #define MIC_TRACING_REGISTER_EVENTS
  #define MIC_TRACING_RECORD(id, start, stop)

#endif
  

// Debug functions


extern void debugInit(FILE* fout);
extern void debugClose();
#if MIC_DEBUG != 0
  extern __thread FILE* mic_output;
  extern double CmiWallTimer();
  extern int CmiMyPe();
  #define MICP(fmt, ...)  fprintf(((mic_output == NULL) ? (stdout) : (mic_output)), "[%013.6lf, %05d] " fmt, CmiWallTimer(), CmiMyPe(), ##__VA_ARGS__)
  #define MICPF           fflush((mic_output != NULL) ? (mic_output) : (stdout))
#else
  #define MICP(fmt, ...)
  #define MICPF
#endif
// NOTE: The bodies are defined in ComputeNonbondMIC.C and these should only
//   be used in host code.

#if MIC_DEVICE_FPRINTF != 0
  #define DEVICE_FPRINTF(fmt, ...) if (device__fout != NULL) { fprintf(device__fout, fmt, ##__VA_ARGS__); fflush(device__fout); }
  #define DEVICE_FPRINTF_CLAUSE  nocopy(device__fout)
#else
  #define DEVICE_FPRINTF(fmt, ...)
  #define DEVICE_FPRINTF_CLAUSE
#endif


// Shared Types

#pragma offload_attribute (push, target(mic))

//this type defined in multiple files
typedef float GBReal;

struct  float3 {  float x,y,z  ; };
struct  float4 {  float x,y,z,w; };
struct double3 { double x,y,z  ; };
struct double4 { double x,y,z,w; };
struct    int2 {    int x,y    ; };
struct    int3 {    int x,y,z  ; };
struct    int4 {    int x,y,z,w; };

typedef float mic_position_t;
typedef float3 mic_position3_t;
typedef float4 mic_position4_t;

// Data structure that represents a compute object/task on the MIC device.  An
//   array of these data structures will be passed down to the device, where each
//   is processed in parallel.  Note that for self computes, patch1 and patch2
//   data should match.
struct patch_pair {  // must be multiple of 16!

  mic_position4_t offset;               // Offset of patches (center to center)

  unsigned int patch1_size;             // Number of atoms for patch 1
  unsigned int patch1_force_size;       // Number of non-fixed atoms at start of list
  unsigned int patch1_atom_start;       // Start of patch 1 data within atoms array
  unsigned int patch1_force_start;      // Start of compute-specific force buffer in split-up force array (i.e. sub-buffer per compute-patch)
  unsigned int patch1_force_list_index; // Index for patch 1's force information in force lists array
  unsigned int patch1_force_list_size;  // Number of contributers to patch 1's force buffer

  unsigned int patch2_size;             // Number of atoms for patch 2
  unsigned int patch2_force_size;       // Number of non-fixed atoms at start of list
  unsigned int patch2_atom_start;       // Start of patch 2 data within atoms array
  unsigned int patch2_force_start;      // Start of compute-specific force buffer in split-up force array (i.e. sub-buffer per compute-patch)
  unsigned int patch2_force_list_index; // Index for patch 2's force information in force lists array
  unsigned int patch2_force_list_size;  // Number of contributers to patch 1's force buffer

  double plcutoff;                      // Pairlist cutoff distance

  unsigned int block_flags_start;
  unsigned int virial_start;            // virial output location padded to 16

  int numParts;
  int part;

  #if MIC_SUBMIT_ATOMS_ON_ARRIVAL != 0
    uint64_t patch1_atomDataPtr;
    uint64_t patch2_atomDataPtr;
  #endif

  // DMK - DEBUG
  mic_position_t patch1_center_x;
  mic_position_t patch1_center_y;
  mic_position_t patch1_center_z;
  mic_position_t patch2_center_x;
  mic_position_t patch2_center_y;
  mic_position_t patch2_center_z;

  // DMK - DEBUG
  int p1;
  int p2;
  int cid;

  int __padding0__[1];
};

// Data structure used to associate individual force outputs (per compute object)
//   to the force outputs for patches (per patch object).
struct force_list {  // must be multiple of 16!
  unsigned int force_list_start;     // beginning of compute output
  unsigned int force_list_size;      // number of computes for this patch
  unsigned int patch_size;           // real number of atoms in patch
  unsigned int patch_stride;         // padded number of atoms in patch
  unsigned int force_output_start;   // output array
  unsigned int atom_start;           // atom positions
  unsigned int virial_list_start;    // beginning of compute virial output
  unsigned int virial_output_start;  // virial output location padded to 16
};

// Data structure for the portion of the atom information that is updated on the
//   device each timestep.
// NOTE: This data structure must match the CudaAtom data structure.
// TODO: Modify the code in the patchs to use this data structure rather than CudaAtom.
struct atom {
  //mic_position3_t position; // x, y, z
  mic_position_t x;
  mic_position_t y;
  mic_position_t z;
  mic_position_t charge;    // q
};

// Data structure for the portion of the atom information that is updated
//   periodically on the device (e.g. after atom migration).
struct atom_param {
  int vdw_type; // NOTE: The handcoded version of the MIC code relies on this member ("vdw_type") being the first of four ints
  int index;
  int excl_index;
  int excl_maxdiff;  // NOTE: if maxdiff == 0, then only excluded from self
};

// Data structure containing several constants that are used by the force
//   computation functions.
struct mic_constants {

  double cutoff2;
  double dielectric_1;
  double scaling;
  double scale14;
  double modf_mod;
  double r2_delta;

  int r2_delta_exp;
  int r2_delta_expc;

  int commOnly;
  int singleKernelFlag;
};

// Data structure for each invocation of the offload pragma in mic_nonbonded_force,
//   used to transfer per timestep information related to the kernels each timestep.
struct mic_kernel_data {

  // Inputs
  int isRemote;
  int numLocalAtoms;
  int numLocalComputes;
  int numLocalPatches;
  int doSlow;
  int doEnergy;
  int usePairlists;
  int savePairlists;

  mic_position3_t lata;
  mic_position3_t latb;
  mic_position3_t latc;
  int numAtoms;
  int numPatchPairs;
  int numForceLists;
  size_t forceBuffersReqSize;

  // Outputs
  double virial_xx;
  double virial_xy;
  double virial_xz;
  double virial_yy;
  double virial_yz;
  double virial_zz;
  double fullElectVirial_xx;
  double fullElectVirial_xy;
  double fullElectVirial_xz;
  double fullElectVirial_yy;
  double fullElectVirial_yz;
  double fullElectVirial_zz;
  double vdwEnergy;
  double electEnergy;
  double fullElectEnergy;

  int exclusionSum;

  int __padding_0__;
};

// Data structure that is setup each time a compute object is about to be processed
//   on the device and passed into the force computation function (e.g. calc_pair)
//   that is being called for the given compute on the given timestep.  NOTE: When
//   the computes are being processed on the device, there is a one-to-one relationship
//   between patch_pairs and mic_params (one per iteration of the parallel for
//   processing the computes).
struct mic_params {
  // DMK - TODO | FIXME : Reorder fields to pack data into less memory (if needed)

  patch_pair *pp;            // Pointer to the associated patch_pair structure
  force_list *fl[2];         // Pointers to force list structures (one per input patch)
  mic_constants *constants;  // Pointer to the mic_constants structure used to process the compute
  int **pairlists_ptr;       // Pointer to short array of pointers, one entry per type of pairlist
  int usePairlists;          // Flag if using pairlist (should be the case for MIC)
  int savePairlists;         // Flag if need to save (generate) pairlists this timestep

  // If refining pairlists is enabled, these are pointers to other pointers/scalars so
  //   that they can be maintained across calls/timesteps
  #if REFINE_PAIRLISTS != 0
    int** plArray_ptr;     // Pointer to the scratch buffer pointer for pairlist entries (so can be maintained across calls)
    int* plSize_ptr;       // Pointer to the length of the plArray scratch buffer
    double** r2Array_ptr;  // Pointer to the scratch buffer pointer for r2 (radius squared) values
  #endif

  void *table_four_base_ptr;     // Base pointer for nonbonded force table
  void *lj_table_base_ptr;       // Base pointer for LJ table data
  int table_four_n_16;           // Sub-array length for nonbonded force table
  int lj_table_dim;              // Dimention of the LJ table
  unsigned int *exclusion_bits;  // Pointer to the exclusion table

  mic_position4_t offset;  // Offset between the patches for pair computes
  int numAtoms[2];         // Atom counts for each patch
  int numAtoms_16[2];      // Atom counts for each patch, rounded up to nearest multiple of 16

  // Atom information arrays (input), formatted as SOA, with one pointer per input patch (1 for self, 2 for pair)
  #if MIC_HANDCODE_FORCE_SOA_VS_AOS != 0

    atom* p[2];
    atom_param* pExt[2];
    double4* ff[2];
    double4* fullf[2];

  #else

    mic_position_t *p_x[2];    // Atom's position.x information
    mic_position_t *p_y[2];    // Atom's position.y information
    mic_position_t *p_z[2];    // Atom's position.z information
    mic_position_t *p_q[2];    // Atom's charge information
    int *pExt_vdwType[2];      // Atom's Van Der Waal information
    int *pExt_index[2];        // Atom's global index (i.e. vs the index in this array)
    int *pExt_exclIndex[2];    // Atom's exclusion table offset
    int *pExt_exclMaxDiff[2];  // Atom's exclusion table entry size (# entries = 2 * maxDiff + 1)
  
    // Atom force arrays (output), formatted as SOA, with one pointer per input patch (1 for self, 2 for pair).
    //   The "ff" arrays are forces for each timestep, while "fullf" arrays are slow_forces calculated during
    //   PME timesteps.  Otherwise, just x, y, z, and w (exclusion accumulators) fields.
    double *ff_x[2];
    double *ff_y[2];
    double *ff_z[2];
    double *ff_w[2];
    double *fullf_x[2];
    double *fullf_y[2];
    double *fullf_z[2];
    double *fullf_w[2];

  #endif
  
  // Virial and energy output for a given compute function ("virial" = every timestep,
  //   "fullElectVirial" = PME timesteps, "Energy" = energy timesteps) that is added
  //   to a global reduction.
  double virial_xx;
  double virial_xy;
  double virial_xz;
  double virial_yy;
  double virial_yz;
  double virial_zz;
  double fullElectVirial_xx;
  double fullElectVirial_xy;
  double fullElectVirial_xz;
  double fullElectVirial_yy;
  double fullElectVirial_yz;
  double fullElectVirial_zz;
  double vdwEnergy;
  double electEnergy;
  double fullElectEnergy;

  int doSlow;  // Flag that indicates if slow_forces should be calculated this timestep or not

  // DMK - DEBUG - Record the center of each patch
  mic_position_t patch1_center_x;
  mic_position_t patch1_center_y;
  mic_position_t patch1_center_z;
  mic_position_t patch2_center_x;
  mic_position_t patch2_center_y;
  mic_position_t patch2_center_z;

  // DMK - DEBUG
  int ppI;
  int p1;
  int p2;
  int pe;
  int timestep;

  int exclusionSum;

  int __padding0__[1];
};

#pragma offload_attribute (pop)


// Functions

// NOTE: Also see comments for these functions in ComputeNonbondedMICKernel.C.

// A function to initialize variables, etc. related to the device's operation.
void mic_init_device(const int pe, const int node, const int deviceNum);

// A functions to free variables, etc. related to the device's operation.
void mic_free_device(const int deviceNum);

// A function to test if the function's execution is actually offloaded to
//   a device (i.e. test if offloading is working correctly).
//   Returns 1 if offloading is working for device 'dev', and 0 if not.
int mic_check(int deviceNum);

// A function that transfers the nonbonded force table to the device
void mic_bind_table_four(const int deviceNum,
                         const double * table_four,
                         const int table_four_n,
                         const int table_four_n_16
                        );

// A function that transfers the Lennard-Jones table to the device
void mic_bind_lj_table(const int deviceNum,
                       const char * lj_table,
                       const int lj_table_dim,
                       const int lj_table_size
                      );

// A function that transfers the exclusion table to the device
void mic_bind_exclusions(const int deviceNum,
                         unsigned int *exclusion_bits,
                         const long int exclusion_bits_size
                        );

// A function that transfers various constants to the device
void mic_bind_constants(const int deviceNum,
                        const double cutoff2,
                        const double dielectric_1,
                        const double scaling,
                        const double scale14,
                        const double r2_delta,
                        const int r2_delta_exp,
                        const int commOnly
                       );



void mic_bind_patch_pairs_only(const int deviceNum,
                               const patch_pair *patch_pairs,
                               const int patch_pairs_size,
                               const int patch_pairs_bufSize
                              );

void mic_bind_force_lists_only(const int deviceNum,
                               force_list *force_lists,
                               const int force_lists_size,
                               const int force_lists_bufSize
                              );

void mic_bind_atoms_only(const int deviceNum,
                         atom *atoms,
                         atom_param *atom_params,
                         double4 *forces,
                         double4 *slow_forces,
                         const int atoms_size,
                         const int atoms_bufSize
                        );

void mic_bind_force_buffers_only(const int deviceNum, const size_t force_buffers_size);


#if MIC_SUBMIT_ATOMS_ON_ARRIVAL != 0
void mic_submit_patch_data(const int deviceNum,
                           void * const hostPtr,
                           void* &hostPtr_prev,
                           const int transferBytes,
                           const int allocBytes_host,
                           int &allocBytes_device,
                           uint64_t &devicePtr,
                           void* &signal
                          );
#endif  // MIC_SUBMIT_ATOMS_ON_ARRIVAL


// A function that initiates the force computation that is done on the device
//   each timestep, resulting in forces being transferred from the device
//   to the host.
void mic_nonbonded_forces(const int deviceNum,
                          const int isRemote,
                          const int numLocalAtoms,
                          const int numLocalComputes,
                          const int numLocalPatches,
                          const mic_position3_t lata,
                          const mic_position3_t latb,
                          const mic_position3_t latc,
                          const int doSlow,
                          const int doEnergy,
                          const int usePairlists,
                          const int savePairlists,
                          const int atomsChanged
                         );

void mic_transfer_output(const int deviceNum,
                         const int isRemote,
                         const int numLocalAtoms,
                         const int doSlow
                        );

// Functions that are used to poll the for the completion of the asynchronous
//   calculation initiated by mic_nonboned_forces.
int mic_check_remote_kernel_complete(const int deviceNum);
int mic_check_local_kernel_complete(const int deviceNum);

// DMK - DEBUG
void _mic_print_stats(const int deviceNum);
void mic_dumpHostDeviceComputeMap();


// DMK - DEBUG
#if MIC_TRACK_DEVICE_MEM_USAGE != 0

__declspec(target(mic))
typedef struct _mem_info {
  size_t memTotal;
  size_t memFree;
  size_t cached;
  size_t active;
  size_t inactive;
  size_t vmSize;
  size_t vmPeak;
} MemInfo;

__declspec(target(mic))
void printMemInfo(int device__pe, int device__timestep, MemInfo * mi);

__declspec(target(mic))
void readMemInfo_processLine(MemInfo * memInfo, char * n, char * v, char * u);

__declspec(target(mic))
bool readMemInfo(MemInfo * memInfo);

#endif  // MIC_TRACK_DEVICE_MEM_USAGE != 0


#endif  // NAMD_MIC
#endif  // COMPUTE_NONBONDED_MIC_KERNEL_H