Merge 'master' into 'amd-master'

Change-Id: Id44416aee5dc592d7b6351926abf53d4a49e9b45

.gitignore

@@ -13,5 +13,6 @@ bin/hipInfo
 bin/hipBusBandwidth
 bin/hipDispatchLatency
 bin/hipify-clang
+include/hip/hcc_detail/hip_prof_str.h
 
 samples/1_Utils/hipInfo/hipInfo

CMakeLists.txt

@@ -160,9 +160,21 @@ if(NOT DEFINED COMPILE_HIP_ATP_MARKER)
 endif()
 add_to_config(_buildInfo COMPILE_HIP_ATP_MARKER)
 
-################
-# Detect profiling API
-################
+#############################
+# Profiling API support
+#############################
+# Generate profiling API macros/structures header
+set(PROF_API_STR "${CMAKE_CURRENT_SOURCE_DIR}/include/hip/hcc_detail/hip_prof_str.h")
+set(PROF_API_HDR "${CMAKE_CURRENT_SOURCE_DIR}/include/hip/hcc_detail/hip_runtime_api.h")
+set(PROF_API_SRC "${CMAKE_CURRENT_SOURCE_DIR}/src")
+set(PROF_API_GEN "${CMAKE_CURRENT_SOURCE_DIR}/hip_prof_gen.py")
+set(PROF_API_LOG "${PROJECT_BINARY_DIR}/hip_prof_gen.log.txt")
+set(PROF_API_CMD "${PROF_API_GEN} -v ${OPT_PROF_API} ${PROF_API_HDR} ${PROF_API_SRC} ${PROF_API_STR} >${PROF_API_LOG}")
+MESSAGE(STATUS "Generating profiling promitives: ${PROF_API_STR}")
+execute_process(COMMAND sh -c "rm -f ${PROF_API_STR}; ${PROF_API_CMD}")
+set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS ${PROF_API_GEN} ${PROF_API_HDR} ${PROF_API_STR})
+
+# Enable profiling API
 if(USE_PROF_API EQUAL 1)
     find_path(PROF_API_HEADER_DIR prof_protocol.h
         HINTS

bin/hipcc

@@ -196,8 +196,10 @@ if ($HIP_PLATFORM eq "clang") {
     $HIPLDFLAGS .= " $HIP_DEVLIB_FLAGS -L$HIP_LIB_PATH";
     if (not $isWindows) {
       $HIPLDFLAGS .= " -Wl,--rpath=$HIP_LIB_PATH";
+      $HIPLDFLAGS .= " -lhip_hcc";
+    } else {
+      $HIPLDFLAGS .= " -lamdhip64";
     }
-    $HIPLDFLAGS .= " -lhip_hcc";
     if ($HIP_CLANG_HCC_COMPAT_MODE) {
         ## Allow __fp16 as function parameter and return type.
         $HIPCXXFLAGS .= " -Xclang -fallow-half-arguments-and-returns -D__HIP_HCC_COMPAT_MODE__=1";

docs/markdown/hip_faq.md

@@ -77,7 +77,7 @@ HIP code provides the same performance as native CUDA code, plus the benefits of
 HIP APIs and features do not map to a specific CUDA version. HIP provides a strong subset of functionality provided in CUDA, and the hipify tools can 
 scan code to identify any unsupported CUDA functions - this is useful for identifying the specific features required by a given application.
 
-However, we can provide a rough summary of the features included in each CUDA SDK and the support level in HIP:
+However, we can provide a rough summary of the features included in each CUDA SDK and the support level in HIP. Each bullet below lists the major new language features in each CUDA release and then indicate which are supported/not supported in HIP:
 
 - CUDA 4.0 and earlier :  
     - HIP supports CUDA 4.0 except for the limitations described above.

hip_prof_gen.py

@@ -0,0 +1,478 @@
+#!/usr/bin/python
+import os, sys, re
+
+PROF_HEADER = "hip_prof_str.h"
+OUTPUT = PROF_HEADER
+REC_MAX_LEN = 1024
+
+# Messages and errors controll
+verbose = 0
+errexit = 0
+inp_file = 'none'
+line_num = -1
+
+# Verbose message
+def message(msg):
+  if verbose: print >>sys.stdout, msg
+
+# Fatal error termination
+def error(msg):
+  if line_num != -1:
+    msg += ", file '" + inp_file + "', line (" + str(line_num) + ")"
+  if errexit:
+    msg = " Error: " + msg
+  else:
+    msg = " Warning: " + msg
+
+  print >>sys.stdout, msg
+  print >>sys.stderr, sys.argv[0] + msg
+
+def fatal(msg):
+  error(msg)
+  if errexit: sys.exit(1)
+
+#############################################################
+# Normalizing API arguments
+def filtr_api_args(args_str):
+  args_str = re.sub(r'^\s*', r'', args_str);
+  args_str = re.sub(r'\s*$', r'', args_str);
+  args_str = re.sub(r'\s*,\s*', r',', args_str);
+  args_str = re.sub(r'\s+', r' ', args_str);
+  args_str = re.sub(r'void \*', r'void* ', args_str);
+  args_str = re.sub(r'(enum|struct) ', '', args_str);
+  return args_str
+
+# Normalizing types
+def norm_api_types(type_str):
+  type_str = re.sub(r'uint32_t', r'unsigned int', type_str)
+  type_str = re.sub(r'^unsigned$', r'unsigned int', type_str)
+  return type_str
+
+# Creating a list of arguments [(type, name), ...]
+def list_api_args(args_str):
+  args_str = filtr_api_args(args_str)
+  args_list = []
+  if args_str != '':
+    for arg_pair in args_str.split(','):
+      if arg_pair == 'void': continue
+      arg_pair = re.sub(r'\s*=\s*\S+$','', arg_pair);
+      m = re.match("^(.*)\s(\S+)$", arg_pair);
+      if m:
+        arg_type = norm_api_types(m.group(1))
+        arg_name = m.group(2)
+        args_list.append((arg_type, arg_name))
+      else:
+        fatal("bad args: args_str: '" + args_str + "' arg_pair: '" + arg_pair + "'")
+  return args_list;
+
+# Creating arguments string "type0, type1, ..."
+def filtr_api_types(args_str):
+  args_list = list_api_args(args_str)
+  types_str = ''
+  for arg_tuple in args_list:
+    types_str += arg_tuple[0] + ', '
+  return types_str
+
+# Creating options list [opt0, opt1, ...]
+def filtr_api_opts(args_str):
+  args_list = list_api_args(args_str)
+  opts_list = []
+  for arg_tuple in args_list:
+    opts_list.append(arg_tuple[1])
+  return opts_list
+#############################################################
+# Parsing API header
+# hipError_t hipSetupArgument(const void* arg, size_t size, size_t offset);
+def parse_api(inp_file_p, out):
+  global inp_file
+  global line_num
+  inp_file = inp_file_p
+
+  beg_pattern = re.compile("^(hipError_t|const char\s*\*)\s+([^\(]+)\(");
+  api_pattern = re.compile("^(hipError_t|const char\s*\*)\s+([^\(]+)\(([^\)]*)\)");
+  end_pattern = re.compile("Texture");
+  hidden_pattern = re.compile(r'__attribute__\(\(visibility\("hidden"\)\)\)')
+  nms_open_pattern = re.compile(r'namespace hip_impl {')
+  nms_close_pattern = re.compile(r'}')
+
+  inp = open(inp_file, 'r')
+
+  found = 0
+  hidden = 0
+  nms_level = 0;
+  record = ""
+  line_num = -1 
+
+  for line in inp.readlines():
+    record += re.sub(r'^\s+', r' ', line[:-1])
+    line_num += 1
+
+    if len(record) > REC_MAX_LEN:
+      fatal("bad record \"" + record + "\"")
+
+    m = beg_pattern.match(line)
+    if m:
+      name = m.group(2)
+      if hidden != 0:
+        message("api: " + name + " - hidden")
+      elif nms_level != 0:
+        message("api: " + name + " - hip_impl")
+      else:
+        message("api: " + name)
+        found = 1
+
+    if found != 0:
+      record = re.sub("\s__dparm\([^\)]*\)", '', record);
+      m = api_pattern.match(record)
+      if m:
+        found = 0
+        if end_pattern.search(record): break
+        out[m.group(2)] = m.group(3)
+      else: continue
+
+    hidden = 0
+    if hidden_pattern.match(line): hidden = 1
+
+    if nms_open_pattern.match(line): nms_level += 1
+    if (nms_level > 0) and nms_close_pattern.match(line): nms_level -= 1
+    if nms_level < 0:
+      fatal("nms level < 0")
+
+    record = ""
+
+  inp.close()
+  line_num = -1
+#############################################################
+# Parsing API implementation
+# hipError_t hipSetupArgument(const void* arg, size_t size, size_t offset) {
+#    HIP_INIT_CB(hipSetupArgument, arg, size, offset);
+# inp_file - input implementation source file
+# api_map - input public API map [<api name>] => <api args>
+# out - output map  [<api name>] => [opt0, opt1, ...]
+def parse_content(inp_file_p, api_map, out):
+  global inp_file
+  global line_num
+  inp_file = inp_file_p
+
+  # API definition begin pattern
+  beg_pattern = re.compile("^(hipError_t|const char\s*\*)\s+[^\(]+\(");
+  # API definition complete pattern
+  api_pattern = re.compile("^(hipError_t|const char\s*\*)\s+([^\(]+)\(([^\)]*)\)\s*{");
+  # API init macro pattern
+  init_pattern = re.compile("^\s*HIP_INIT[_\w]*_API\(([^,]+)(,|\))");
+  target_pattern = re.compile("^(\s*HIP_INIT[^\(]*)(_API\()(.*)\);\s*$");
+
+  # Open input file
+  inp = open(inp_file, 'r')
+
+  # API name
+  api_name = ""
+  # Valid public API found flag
+  api_valid = 0
+
+  # Input file patched content
+  content = ''
+  # Sub content for found API defiition
+  sub_content = ''
+  # Current record, accumulating several API definition related lines
+  record = ''
+  # Current input file line number
+  line_num = -1 
+  # API beginning found flag
+  found = 0
+
+  # Reading input file
+  for line in inp.readlines():
+    # Accumulating record
+    record += re.sub(r'^\s+', r' ', line[:-1])
+    line_num += 1
+
+    if len(record) > REC_MAX_LEN:
+      fatal("bad record \"" + record + "\"")
+      break;
+
+    # Looking for API begin
+    if beg_pattern.match(record): found = 1
+
+    # Matching complete API definition
+    if found == 1:
+      record = re.sub("\s__dparm\([^\)]*\)", '', record);
+      m = api_pattern.match(record)
+      # Checking if complete API matched
+      if m:
+        found = 2
+        api_name = m.group(2);
+        # Checking if API name is in the API map
+        if api_name in api_map:
+          # Getting API arguments
+          api_args = m.group(3)
+          # Getting etalon arguments from the API map
+          eta_args = api_map[api_name]
+          if eta_args == '':
+            eta_args = api_args
+            api_map[api_name] = eta_args
+          # Normalizing API arguments
+          api_types = filtr_api_types(api_args)
+          # Normalizing etalon arguments
+          eta_types = filtr_api_types(eta_args)
+          if api_types == eta_types:
+            # API is already found
+            if api_name in out:
+              fatal("API redefined \"" + api_name + "\", record \"" + record + "\"")
+            # Set valid public API found flag
+            api_valid = 1
+            # Set output API map with API arguments list
+            out[api_name] = filtr_api_opts(api_args)
+          else:
+            # Warning about mismatched API, possible non public overloaded version
+            api_diff = '\t\t' + inp_file + " line(" + str(line_num) + ")\n\t\tapi: " + api_types + "\n\t\teta: " + eta_types
+            message("\t" + api_name + ' args mismatch:\n' + api_diff + '\n')
+
+    # API found action
+    if found == 2:
+      # Looking for INIT macro
+      m = init_pattern.match(line)
+      if m:
+        found = 0
+        if api_valid == 1:
+          api_valid = 0
+          message("\t" + api_name)
+        else:
+          # Registering dummy API for non public API if the name in INIT is not NONE
+          init_name = m.group(1)
+          # Ignore if it is initialized as NONE
+          if init_name != 'NONE':
+            # Check if init name matching API name
+            if init_name != api_name:
+              fatal("init name mismatch: '" + init_name +  "' <> '" + api_name + "'")
+            # If init name is not in public API map then it is private API
+            # else it was not identified and will be checked on finish
+            if not init_name in api_map:
+              if init_name in out:
+                fatal("API reinit \"" + api_name + "\", record \"" + record + "\"")
+              out[init_name] = []
+      elif re.search('}', line):
+        found = 0
+        # Expect INIT macro for valid public API
+        if api_valid == 1:
+          api_valid = 0
+          if api_name in out:
+            del out[api_name]
+            del api_map[api_name]
+            out['.' + api_name] = 1
+          else:
+            fatal("API is not in out \"" + api_name + "\", record \"" + record + "\"")
+
+    if found != 1: record = ""
+    content += line
+
+  inp.close()
+  line_num = -1
+
+  if len(out) != 0:
+    return content
+  else:
+    return '' 
+
+# src path walk
+def parse_src(api_map, src_path, src_patt, out):
+  pattern = re.compile(src_patt)
+  src_path = re.sub(r'\s', '', src_path)
+  for src_dir in src_path.split(':'):
+    message("Parsing " + src_dir + " for '" + src_patt + "'")
+    for root, dirs, files in os.walk(src_dir):
+      for fnm in files:
+        if pattern.search(fnm):
+          file = root + '/' + fnm
+          message(file)
+          content = parse_content(file, api_map, out);
+          if content != '':
+            f = open(file, 'w')
+            f.write(content)
+            f.close()
+#############################################################
+# Generating profiling primitives header
+# api_map - public API map [<api name>] => [(type, name), ...]
+# opts_map - opts map  [<api name>] => [opt0, opt1, ...]
+def generate_prof_header(f, api_map, opts_map):
+  # Private API list
+  priv_lst = []
+
+  f.write('// automatically generated sources\n')
+  f.write('#ifndef _HIP_PROF_STR_H\n');
+  f.write('#define _HIP_PROF_STR_H\n');
+  f.write('#include <sstream>\n');
+  f.write('#include <string>\n');
+  
+  # Generating dummy macro for non-public API
+  f.write('\n// Dummy API primitives\n')
+  f.write('#define INIT_NONE_CB_ARGS_DATA(cb_data) {};\n')
+  for name in opts_map:
+    if not name in api_map:
+      opts_lst = opts_map[name]
+      if len(opts_lst) != 0:
+        fatal("bad dummy API \"" + name + "\", args: " + str(opts_lst))
+      f.write('#define INIT_'+ name + '_CB_ARGS_DATA(cb_data) {};\n')
+      priv_lst.append(name)
+  
+  for name in priv_lst:
+    message("Private: " + name)
+  
+  # Generating the callbacks ID enumaration
+  f.write('\n// HIP API callbacks ID enumaration\n')
+  f.write('enum hip_api_id_t {\n')
+  cb_id = 0
+  for name in api_map.keys():
+    f.write('  HIP_API_ID_' + name + ' = ' + str(cb_id) + ',\n')
+    cb_id += 1
+  f.write('  HIP_API_ID_NUMBER = ' + str(cb_id) + ',\n')
+  f.write('  HIP_API_ID_ANY = ' + str(cb_id + 1) + ',\n')
+  f.write('\n')
+  f.write('  HIP_API_ID_NONE = HIP_API_ID_NUMBER,\n')
+  for name in priv_lst:
+    f.write('  HIP_API_ID_' + name + ' = HIP_API_ID_NUMBER,\n')
+  f.write('};\n')
+  
+  # Generating the callbacks ID enumaration
+  f.write('\n// Return HIP API string\n')
+  f.write('static const char* hip_api_name(const uint32_t& id) {\n')
+  f.write('  switch(id) {\n')
+  for name in api_map.keys():
+    f.write('    case HIP_API_ID_' + name + ': return "' +  name + '";\n')
+  f.write('  };\n')
+  f.write('  return "unknown";\n')
+  f.write('};\n')
+  
+  # Generating the callbacks data structure
+  f.write('\n// HIP API callbacks data structure\n')
+  f.write(
+  'struct hip_api_data_t {\n' +
+  '  uint64_t correlation_id;\n' +
+  '  uint32_t phase;\n' +
+  '  union {\n'
+  )
+  for name, args in api_map.items():
+    if len(args) != 0:
+      f.write('    struct {\n')
+      for arg_tuple in args:
+        f.write('      ' + arg_tuple[0] + ' ' + arg_tuple[1] + ';\n')
+      f.write('    } ' + name + ';\n')
+  f.write(
+  '  } args;\n' +
+  '};\n'
+  )
+  
+  # Generating the callbacks args data filling macros
+  f.write('\n// HIP API callbacks args data filling macros\n')
+  for name, args in api_map.items():
+    f.write('// ' + name + str(args) + '\n')
+    f.write('#define INIT_' + name + '_CB_ARGS_DATA(cb_data) { \\\n')
+    if name in opts_map:
+      opts_list = opts_map[name]
+      if len(args) != len(opts_list):
+        fatal("\"" + name + "\" API args and opts mismatch, args: " + str(args) + ", opts: " + str(opts_list))
+      # API args iterating:
+      #   type is args[<ind>][0]
+      #   name is args[<ind>][1]
+      for ind in range(0, len(args)):
+        arg_tuple = args[ind]
+        fld_name = arg_tuple[1]
+        arg_name = opts_list[ind]
+        f.write('  cb_data.args.' + name + '.' + fld_name + ' = ' + arg_name + '; \\\n')
+    f.write('};\n')
+  f.write('#define INIT_CB_ARGS_DATA(cb_id, cb_data) INIT_##cb_id##_CB_ARGS_DATA(cb_data)\n')
+  
+  # Generating the method for the API string, name and parameters
+  f.write('\n')
+  f.write('#if 0\n')
+  f.write('// HIP API string method, method name and parameters\n')
+  f.write('const char* hipApiString(hip_api_id_t id, const hip_api_data_t* data) {\n')
+  f.write('  std::ostringstream oss;\n')
+  f.write('  switch (id) {\n')
+  for name, args in api_map.items():
+    f.write('    case HIP_API_ID_' + name + ':\n')
+    f.write('      oss << "' + name + '("')
+    for ind in range(0, len(args)):
+      arg_tuple = args[ind]
+      arg_name = arg_tuple[1]
+      if ind != 0: f.write(' << ","')
+      f.write('\n          << " ' + arg_name  + '=" << data->args.' + name + '.' + arg_name)
+    f.write('\n          << ")";\n')
+    f.write('    break;\n')
+  f.write('    default: oss << "unknown";\n')
+  f.write('  };\n')
+  f.write('  return strdup(oss.str().c_str());\n')
+  f.write('};\n')
+  f.write('#endif\n')
+  
+  f.write('#endif  // _HIP_PROF_STR_H\n');
+
+#############################################################
+# main
+# Usage
+if (len(sys.argv) > 1) and (sys.argv[1] == '-v'):
+  verbose = 1
+  sys.argv.pop(1)
+
+if (len(sys.argv) > 1) and (sys.argv[1] == '-e'):
+  errexit = 1
+  sys.argv.pop(1)
+
+if (len(sys.argv) < 3):
+  fatal ("Usage: " + sys.argv[0] + " [-v] <input HIP API .h file> <patched srcs path>\n" +
+         "  -v - verbose messages\n" +
+         "  example:\n" +
+         "  $ hipap.py hip/include/hip/hcc_detail/hip_runtime_api.h hip/src")
+
+# API header file given as an argument
+api_hfile = sys.argv[1]
+if not os.path.isfile(api_hfile):
+  fatal("input file '" + api_hfile + "' not found")
+
+# Srcs directory given as an argument
+src_pat = "\.cpp$"
+src_dir = sys.argv[2]
+if not os.path.isdir(src_dir):
+  fatal("src directory " + src_dir + "' not found")
+
+if len(sys.argv) > 3: OUTPUT = sys.argv[3]
+
+# API declaration map
+api_map = {
+  'hipHccModuleLaunchKernel': ''
+}
+# API options map
+opts_map = {}
+
+# Parsing API header
+parse_api(api_hfile, api_map)
+
+# Parsing sources
+parse_src(api_map, src_dir, src_pat, opts_map)
+
+# Checking for non-conformant APIs
+for name in opts_map.keys():
+  m = re.match(r'\.(\S*)', name)
+  if m:
+    message("Init missing: " + m.group(1))
+    del opts_map[name]
+
+# Converting api map to map of lists
+# Checking for not found APIs
+not_found = 0
+if len(opts_map) != 0:
+  for name in api_map.keys():
+    args_str = api_map[name];
+    api_map[name] = list_api_args(args_str)
+    if not name in opts_map:
+      error("implementation not found: " + name)
+      not_found += 1
+if not_found != 0:
+  fatal(str(not_found) + " API calls missing in interception layer")
+
+# Generating output header file
+with open(OUTPUT, 'w') as f:
+  generate_prof_header(f, api_map, opts_map)
+
+# Successfull exit
+sys.exit(0)

include/hip/hcc_detail/device_functions.h

@@ -46,7 +46,7 @@ __device__ static inline unsigned int __popcll(unsigned long long int input) {
 }
 
 __device__ static inline int __clz(int input) {
-  return __ockl_clz_u32((uint)input);
+    return __ockl_clz_u32((uint)input);
 }
 
 __device__ static inline int __clzll(long long int input) {
@@ -224,56 +224,59 @@ __device__ static inline unsigned int __lane_id() { return  __mbcnt_hi(-1, __mbc
 HIP specific device functions
 */
 
-// utility union type
-union __u {
-    int i;
-    unsigned int u;
-    float f;
-};
-
 __device__ static inline unsigned __hip_ds_bpermute(int index, unsigned src) {
-    __u tmp; tmp.u = src;
+    union { int i; unsigned u; float f; } tmp; tmp.u = src;
     tmp.i = __llvm_amdgcn_ds_bpermute(index, tmp.i);
     return tmp.u;
 }
 
 __device__ static inline float __hip_ds_bpermutef(int index, float src) {
-    __u tmp; tmp.f = src;
+    union { int i; unsigned u; float f; } tmp; tmp.f = src;
     tmp.i = __llvm_amdgcn_ds_bpermute(index, tmp.i);
     return tmp.f;
 }
 
 __device__ static inline unsigned __hip_ds_permute(int index, unsigned src) {
-  __u tmp; tmp.u = src;
-  tmp.i = __llvm_amdgcn_ds_permute(index, tmp.i);
-  return tmp.u;
+    union { int i; unsigned u; float f; } tmp; tmp.u = src;
+    tmp.i = __llvm_amdgcn_ds_permute(index, tmp.i);
+    return tmp.u;
 }
 
 __device__ static inline float __hip_ds_permutef(int index, float src) {
-  __u tmp; tmp.u = src;
-  tmp.i = __llvm_amdgcn_ds_permute(index, tmp.i);
-  return tmp.u;
-}
-
-__device__ static inline unsigned __hip_ds_swizzle(unsigned int src, int pattern) {
-    __u tmp; tmp.u = src;
-    tmp.i = __llvm_amdgcn_ds_swizzle(tmp.i, pattern);
+    union { int i; unsigned u; float f; } tmp; tmp.u = src;
+    tmp.i = __llvm_amdgcn_ds_permute(index, tmp.i);
     return tmp.u;
 }
-__device__ static inline float __hip_ds_swizzlef(float src, int pattern) {
-    __u tmp; tmp.f = src;
-    tmp.i = __llvm_amdgcn_ds_swizzle(tmp.i, pattern);
+
+#define __hip_ds_swizzle(src, pattern)  __hip_ds_swizzle_N<(pattern)>((src))
+#define __hip_ds_swizzlef(src, pattern) __hip_ds_swizzlef_N<(pattern)>((src))
+
+template <int pattern>
+__device__ static inline unsigned __hip_ds_swizzle_N(unsigned int src) {
+    union { int i; unsigned u; float f; } tmp; tmp.u = src;
+    tmp.i = __builtin_amdgcn_ds_swizzle(tmp.i, pattern);
+    return tmp.u;
+}
+
+template <int pattern>
+__device__ static inline float __hip_ds_swizzlef_N(float src) {
+    union { int i; unsigned u; float f; } tmp; tmp.f = src;
+    tmp.i = __builtin_amdgcn_ds_swizzle(tmp.i, pattern);
     return tmp.f;
 }
 
-__device__ static inline int __hip_move_dpp(int src, int dpp_ctrl, int row_mask,
-                                            int bank_mask, bool bound_ctrl) {
-    return __llvm_amdgcn_move_dpp(src, dpp_ctrl, row_mask, bank_mask, bound_ctrl);
+#define __hip_move_dpp(src, dpp_ctrl, row_mask, bank_mask, bound_ctrl) \
+  __hip_move_dpp_N<(dpp_ctrl), (row_mask), (bank_mask), (bound_ctrl)>((src))
+
+template <int dpp_ctrl, int row_mask, int bank_mask, bool bound_ctrl>
+__device__ static inline int __hip_move_dpp_N(int src) {
+    return __llvm_amdgcn_move_dpp(src, dpp_ctrl, row_mask, bank_mask,
+                                  bound_ctrl);
 }
 
 static constexpr int warpSize = 64;
 
- __device__
+__device__
 inline
 int __shfl(int var, int src_lane, int width = warpSize) {
     int self = __lane_id();
@@ -283,14 +286,14 @@ int __shfl(int var, int src_lane, int width = warpSize) {
 __device__
 inline
 unsigned int __shfl(unsigned int var, int src_lane, int width = warpSize) {
-     __u tmp; tmp.u = var;
+     union { int i; unsigned u; float f; } tmp; tmp.u = var;
     tmp.i = __shfl(tmp.i, src_lane, width);
     return tmp.u;
 }
 __device__
 inline
 float __shfl(float var, int src_lane, int width = warpSize) {
-    __u tmp; tmp.f = var;
+    union { int i; unsigned u; float f; } tmp; tmp.f = var;
     tmp.i = __shfl(tmp.i, src_lane, width);
     return tmp.f;
 }
@@ -320,14 +323,14 @@ int __shfl_up(int var, unsigned int lane_delta, int width = warpSize) {
 __device__
 inline
 unsigned int __shfl_up(unsigned int var, unsigned int lane_delta, int width = warpSize) {
-    __u tmp; tmp.u = var;
+    union { int i; unsigned u; float f; } tmp; tmp.u = var;
     tmp.i = __shfl_up(tmp.i, lane_delta, width);
     return tmp.u;
 }
 __device__
 inline
 float __shfl_up(float var, unsigned int lane_delta, int width = warpSize) {
-    __u tmp; tmp.f = var;
+    union { int i; unsigned u; float f; } tmp; tmp.f = var;
     tmp.i = __shfl_up(tmp.i, lane_delta, width);
     return tmp.f;
 }
@@ -357,14 +360,14 @@ int __shfl_down(int var, unsigned int lane_delta, int width = warpSize) {
 __device__
 inline
 unsigned int __shfl_down(unsigned int var, unsigned int lane_delta, int width = warpSize) {
-    __u tmp; tmp.u = var;
+    union { int i; unsigned u; float f; } tmp; tmp.u = var;
     tmp.i = __shfl_down(tmp.i, lane_delta, width);
     return tmp.u;
 }
 __device__
 inline
 float __shfl_down(float var, unsigned int lane_delta, int width = warpSize) {
-    __u tmp; tmp.f = var;
+    union { int i; unsigned u; float f; } tmp; tmp.f = var;
     tmp.i = __shfl_down(tmp.i, lane_delta, width);
     return tmp.f;
 }
@@ -394,14 +397,14 @@ int __shfl_xor(int var, int lane_mask, int width = warpSize) {
 __device__
 inline
 unsigned int __shfl_xor(unsigned int var, int lane_mask, int width = warpSize) {
-    __u tmp; tmp.u = var;
+    union { int i; unsigned u; float f; } tmp; tmp.u = var;
     tmp.i = __shfl_xor(tmp.i, lane_mask, width);
     return tmp.u;
 }
 __device__
 inline
 float __shfl_xor(float var, int lane_mask, int width = warpSize) {
-    __u tmp; tmp.f = var;
+    union { int i; unsigned u; float f; } tmp; tmp.f = var;
     tmp.i = __shfl_xor(tmp.i, lane_mask, width);
     return tmp.f;
 }
@@ -681,9 +684,9 @@ inline  __attribute((always_inline))
 long long int __clock64() {
 // ToDo: Unify HCC and HIP implementation.
 #if __HCC__
-  return (long long int) __clock_u64();
+    return (long long int) __clock_u64();
 #else
-  return (long long int) __builtin_amdgcn_s_memrealtime();
+    return (long long int) __builtin_amdgcn_s_memrealtime();
 #endif
 }
 
@@ -881,7 +884,7 @@ __device__
 inline
 __attribute__((weak))
 void abort() {
-  return __builtin_trap();
+    return __builtin_trap();
 }
 
 

include/hip/hcc_detail/functional_grid_launch.hpp

@@ -189,6 +189,7 @@ inline
 void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
                         std::uint32_t sharedMemBytes, hipStream_t stream,
                         Args... args) {
+    hip_impl::hip_init();
     auto kernarg = hip_impl::make_kernarg(
         kernel, std::tuple<Args...>{std::move(args)...});
     std::size_t kernarg_size = kernarg.size();
@@ -212,4 +213,4 @@ inline void hipLaunchKernel(F kernel, const dim3& numBlocks, const dim3& dimBloc
                             std::uint32_t groupMemBytes, hipStream_t stream, Args... args) {
     hipLaunchKernelGGL(kernel, numBlocks, dimBlocks, groupMemBytes, stream, hipLaunchParm{},
                        std::move(args)...);
-}
+}

include/hip/hcc_detail/hip_runtime.h

@@ -190,10 +190,13 @@ __device__ float __hip_ds_bpermutef(int index, float src);
 __device__ unsigned __hip_ds_permute(int index, unsigned src);
 __device__ float __hip_ds_permutef(int index, float src);
 
-__device__ unsigned __hip_ds_swizzle(unsigned int src, int pattern);
-__device__ float __hip_ds_swizzlef(float src, int pattern);
+template <int pattern>
+__device__ unsigned __hip_ds_swizzle_N(unsigned int src);
+template <int pattern>
+__device__ float __hip_ds_swizzlef_N(float src);
 
-__device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask, bool bound_ctrl);
+template <int dpp_ctrl, int row_mask, int bank_mask, bool bound_ctrl>
+__device__ int __hip_move_dpp_N(int src);
 
 #endif  //__HIP_ARCH_GFX803__ == 1
 
@@ -331,18 +334,15 @@ extern void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, gri
 
 typedef int hipLaunchParm;
 
-template <typename... Args, typename F = void (*)(Args...)>
-inline void hipLaunchKernelGGL(F&& kernelName, const dim3& numblocks, const dim3& numthreads,
-                               unsigned memperblock, hipStream_t streamId, Args... args) {
-  kernelName<<<numblocks, numthreads, memperblock, streamId>>>(args...);
-}
+#define hipLaunchKernel(kernelName, numblocks, numthreads, memperblock, streamId, ...)             \
+    do {                                                                                           \
+        kernelName<<<(numblocks), (numthreads), (memperblock), (streamId)>>>(hipLaunchParam{}, ##__VA_ARGS__); \
+    } while (0)
 
-template <typename... Args, typename F = void (*)(hipLaunchParm, Args...)>
-inline void hipLaunchKernel(F&& kernel, const dim3& numBlocks, const dim3& dimBlocks,
-                            std::uint32_t groupMemBytes, hipStream_t stream, Args... args) {
-    hipLaunchKernelGGL(kernel, numBlocks, dimBlocks, groupMemBytes, stream, hipLaunchParm{},
-                       std::move(args)...);
-}
+#define hipLaunchKernelGGL(kernelName, numblocks, numthreads, memperblock, streamId, ...)          \
+    do {                                                                                           \
+        kernelName<<<(numblocks), (numthreads), (memperblock), (streamId)>>>(__VA_ARGS__);         \
+    } while (0)
 
 #include <hip/hip_runtime_api.h>
 

include/hip/hcc_detail/hip_runtime_api.h

@@ -78,6 +78,10 @@ THE SOFTWARE.
   #define __dparm(x)
 #endif
 
+namespace hip_impl {
+hipError_t hip_init();
+}  // namespace hip_impl
+
 // Structure definitions:
 #ifdef __cplusplus
 extern "C" {
@@ -1377,7 +1381,26 @@ hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, h
 hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes,
                               hipStream_t stream);
 
-#if !__HIP_VDI__
+#if __HIP_VDI__
+hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes,
+    hipModule_t hmod, const char* name);
+
+hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName);
+hipError_t hipGetSymbolSize(size_t* size, const void* symbolName);
+hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src,
+                             size_t sizeBytes, size_t offset __dparm(0),
+                             hipMemcpyKind kind __dparm(hipMemcpyHostToDevice));
+hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void* src,
+                                  size_t sizeBytes, size_t offset,
+                                  hipMemcpyKind kind, hipStream_t stream __dparm(0));
+hipError_t hipMemcpyFromSymbol(void* dst, const void* symbolName,
+                               size_t sizeBytes, size_t offset __dparm(0),
+                               hipMemcpyKind kind __dparm(hipMemcpyDeviceToHost));
+hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName,
+                                    size_t sizeBytes, size_t offset,
+                                    hipMemcpyKind kind,
+                                    hipStream_t stream __dparm(0));
+#else
 __attribute__((visibility("hidden")))
 hipError_t hipModuleGetGlobal(void**, size_t*, hipModule_t, const char*);
 
@@ -1396,7 +1419,7 @@ inline
 __attribute__((visibility("hidden")))
 hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName) {
     //HIP_INIT_API(hipGetSymbolAddress, devPtr, symbolName);
-
+    hip_impl::hip_init();
     size_t size = 0;
     return hipModuleGetGlobal(devPtr, &size, 0, (const char*)symbolName);
 }
@@ -1416,7 +1439,7 @@ inline
 __attribute__((visibility("hidden")))
 hipError_t hipGetSymbolSize(size_t* size, const void* symbolName) {
     // HIP_INIT_API(hipGetSymbolSize, size, symbolName);
-
+    hip_impl::hip_init();
     void* devPtr = nullptr;
     return hipModuleGetGlobal(&devPtr, size, 0, (const char*)symbolName);
 }

src/hip_clang.cpp

@@ -32,7 +32,7 @@ THE SOFTWARE.
 extern "C" std::vector<hipModule_t>*
 __hipRegisterFatBinary(const void* data)
 {
-  HIP_INIT();
+  hip_impl::hip_init();
 
   tprintf(DB_FB, "Enter __hipRegisterFatBinary(%p)\n", data);
   const __CudaFatBinaryWrapper* fbwrapper = reinterpret_cast<const __CudaFatBinaryWrapper*>(data);

src/hip_fatbin.cpp

@@ -43,7 +43,7 @@ void __hipDumpCodeObject(const std::string& image) {
 const void* __hipExtractCodeObjectFromFatBinary(const void* data,
     const char* agent_name)
 {
-  HIP_INIT();
+  hip_impl::hip_init();
 
   tprintf(DB_FB, "Enter __hipExtractCodeObjectFromFatBinary(%p, \"%s\")\n",
       data, agent_name);

src/hip_hcc.cpp

@@ -119,9 +119,6 @@ int HCC_OPT_FLUSH = 1;
 int HCC_OPT_FLUSH = 0;
 #endif
 
-
-std::once_flag hip_initialized;
-
 // Array of pointers to devices.
 ihipDevice_t** g_deviceArray;
 
@@ -1442,6 +1439,15 @@ void ihipInit() {
             g_numLogicalThreads);
 }
 
+namespace hip_impl {
+hipError_t hip_init() {
+  static std::once_flag hip_initialized;
+  std::call_once(hip_initialized, ihipInit);
+  ihipCtxStackUpdate();
+  return hipSuccess;
+}
+}
+
 hipError_t ihipStreamSynchronize(hipStream_t stream) {
     hipError_t e = hipSuccess;
 
@@ -1561,7 +1567,6 @@ void ihipPrintKernelLaunch(const char* kernelName, const grid_launch_parm* lp,
 // Allows runtime to track some information about the stream.
 hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_launch_parm* lp,
                                 const char* kernelNameStr) {
-    HIP_INIT();
     stream = ihipSyncAndResolveStream(stream);
     lp->grid_dim.x = grid.x;
     lp->grid_dim.y = grid.y;
@@ -1583,7 +1588,6 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_
 
 hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, grid_launch_parm* lp,
                                 const char* kernelNameStr) {
-    HIP_INIT();
     stream = ihipSyncAndResolveStream(stream);
     lp->grid_dim.x = grid;
     lp->grid_dim.y = 1;
@@ -1604,7 +1608,6 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, gri
 
 hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, grid_launch_parm* lp,
                                 const char* kernelNameStr) {
-    HIP_INIT();
     stream = ihipSyncAndResolveStream(stream);
     lp->grid_dim.x = grid.x;
     lp->grid_dim.y = grid.y;
@@ -1625,7 +1628,6 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, gri
 
 hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, grid_launch_parm* lp,
                                 const char* kernelNameStr) {
-    HIP_INIT();
     stream = ihipSyncAndResolveStream(stream);
     lp->grid_dim.x = grid;
     lp->grid_dim.y = 1;
@@ -2485,4 +2487,4 @@ namespace hip_impl {
             std::terminate();
         #endif
     }
-} // Namespace hip_impl.
+} // Namespace hip_impl.

src/hip_hcc_internal.h

@@ -288,19 +288,13 @@ extern uint64_t recordApiTrace(std::string* fullStr, const std::string& apiStr);
 #define API_TRACE(IS_CMD, ...) tls_tidInfo.incApiSeqNum();
 #endif
 
-
-// Just initialize the HIP runtime, but don't log any trace information.
-#define HIP_INIT()                                                                                 \
-    std::call_once(hip_initialized, ihipInit);                                                     \
-    ihipCtxStackUpdate();
 #define HIP_SET_DEVICE() ihipDeviceSetState();
 
-
 // This macro should be called at the beginning of every HIP API.
 // It initializes the hip runtime (exactly once), and
 // generates a trace string that can be output to stderr or to ATP file.
 #define HIP_INIT_API(cid, ...)                                                                     \
-    HIP_INIT()                                                                                     \
+    hip_impl::hip_init();                                                                                    \
     API_TRACE(0, __VA_ARGS__);                                                                     \
     HIP_CB_SPAWNER_OBJECT(cid);
 
@@ -309,7 +303,7 @@ extern uint64_t recordApiTrace(std::string* fullStr, const std::string& apiStr);
 // Replace HIP_INIT_API with this call inside HIP APIs that launch work on the GPU:
 // kernel launches, copy commands, memory sets, etc.
 #define HIP_INIT_SPECIAL_API(cid, tbit, ...)                                                       \
-    HIP_INIT()                                                                                     \
+    hip_impl::hip_init();                                                                                    \
     API_TRACE((HIP_TRACE_API & (1 << tbit)), __VA_ARGS__);                                         \
     HIP_CB_SPAWNER_OBJECT(cid);
 
@@ -933,7 +927,6 @@ class ihipCtx_t {
 
 //=================================================================================================
 // Global variable definition:
-extern std::once_flag hip_initialized;
 extern unsigned g_deviceCnt;
 extern hsa_agent_t g_cpu_agent;   // the CPU agent.
 extern hsa_agent_t* g_allAgents;  // CPU agents + all the visible GPU agents.