<commit_before>################################################################################
# Copyright (C) 2016 Advanced Micro Devices, Inc. All rights reserved.
#
# Permission is hereby granted, free of charge, to any person obtaining a copy
# of this software and associated documentation files (the "Software"), to deal
# in the Software without restriction, including without limitation the rights
# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell cop-
# ies of the Software, and to permit persons to whom the Software is furnished
# to do so, subject to the following conditions:
#
# The above copyright notice and this permission notice shall be included in all
# copies or substantial portions of the Software.
#
# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IM-
# PLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS
# FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR
# COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER
# IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNE-
# CTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
################################################################################
# This script only gets called by CMake
from Common import globalParameters, HR, print1, print2, printExit, ensurePath, CHeader, CMakeHeader, assignGlobalParameters, ProgressBar, listToInitializer
from SolutionStructs import Solution
import YAMLIO
from SolutionWriter import SolutionWriter
from KernelWriterSource import KernelWriterSource
from KernelWriterAssembly import KernelWriterAssembly
import multiprocessing

import os
import sys
import os.path
import argparse
from shutil import copy as shutil_copy
import time


################################################################################
# Process a single kernel, return results:
################################################################################
def processKernelSource(kernel, kernelWriterSource, kernelWriterAssembly):
    kernelWriter = kernelWriterSource if kernel["KernelLanguage"] == "Source" else kernelWriterAssembly
    # get kernel name
    kernelName = kernelWriter.getKernelName(kernel)
    #sys.stderr.write("kernel:%s\n"% kernelName)
    (err, src) = kernelWriter.getSourceFileString(kernel)

    header = kernelWriter.getHeaderFileString(kernel)

    return (err, src, header, kernelName)


################################################################################
# Process a range of kernels:
# kLock is for the kernel file - is this used?
# pLock is for the progress bar

################################################################################
def processKernelSourceChunk(kernels,
                             kernelWriterSource, kernelWriterAssembly, \
                             kiStart, kiStop, pipe):

    results = []

    for ki in range(kiStart, kiStop):
      kernel = kernels[ki]
      results.append (processKernelSource(kernel, kernelWriterSource, kernelWriterAssembly)) # returns err, src, header, kernelName

    if pipe != None:
      pipe.send(results)
    else:
      return results



# create and prepare the assembly directory  - called ONCE per output dir:
def prepAsm():
  asmPath = ensurePath(os.path.join(globalParameters["WorkingPath"], "assembly") )
  assemblerFileName = os.path.join(asmPath, \
      "asm.%s"%("bat" if os.name=="nt" else "sh"))
  assemblerFile = open(assemblerFileName, "w")
  if os.name == "nt":
    assemblerFile.write("echo Windows: Copying instead of Assembling\n")
    assemblerFile.write("copy %1.s %1.o\n")
    assemblerFile.write("copy %1.o %1.co\n")
  else:
    assemblerFile.write("#!/bin/sh %s\n" % ("-x" if globalParameters["PrintLevel"] >=2  else ""))
    assemblerFile.write("# usage: asm.sh kernelName ASM_ARGS\n")
    assemblerFile.write("# example: asm.sh kernelName -mcpu=gfx900\n")
    assemblerFile.write("f=$1\n")
    assemblerFile.write("shift\n")
    assemblerFile.write("ASM=%s\n"%globalParameters["AssemblerPath"])
    # cannot use globalParameters["CurrentISA"] because it might be (0,0,0)
    defaultIsa = (9,0,0)
    assemblerFile.write( \
      "${ASM} -x assembler -target amdgcn--amdhsa %s $@ -c -o $f.o $f.s\n" % \
      ("-mno-code-object-v3" if \
      globalParameters["AsmCaps"][defaultIsa]["HasCodeObjectV3"] else ""))
    assemblerFile.write("${ASM} -target amdgcn--amdhsa $f.o -o $f.co\n")
  assemblerFile.close()
  os.chmod(assemblerFileName, 0777)

################################################################################
# processResults
# input: results is list with (err, src, header, kernelName)
# Log errors and write appropriate info to kernelSourceFile and kernelHeaderFile
################################################################################
def processResults(results, outputPath, kernelsWithBuildErrs, \
      kernelSourceFile, kernelHeaderFile):

  for (err,src,header,kernelName) in results:
    if err:
      kernelsWithBuildErrs[kernelName] = err
      #print "*** warning: invalid kernel#%s"%kernelName

    # write kernel.cpp
    if not globalParameters["MergeFiles"]:
      kernelSourceFile = open(os.path.join(outputPath, \
          "Kernels", kernelName+".cpp"), "w")
      kernelSourceFile.write(CHeader)

    kernelSourceFile.write(src)

    if not globalParameters["MergeFiles"]:
      kernelSourceFile.close()
      # write kernel.h
      kernelHeaderFile = open(os.path.join(outputPath, \
          "Kernels", kernelName+".h"), "w")
      kernelHeaderFile.write(CHeader)

    kernelHeaderFile.write(header)

    if not globalParameters["MergeFiles"]:
      kernelHeaderFile.close()

################################################################################
# Write Solutions and Kernels for BenchmarkClient or LibraryClient
################################################################################
def writeSolutionsAndKernels(outputPath, problemTypes, solutions, kernels, kernelsBetaOnly, \
    solutionWriter, kernelWriterSource, kernelWriterAssembly):
  start = time.time()
  print1("# Writing Kernels...")
  if not globalParameters["MergeFiles"]:
    ensurePath(os.path.join(outputPath, "Solutions"))
    ensurePath(os.path.join(outputPath, "Kernels"))

  if globalParameters["ShowProgressBar"]:
    progressBar = ProgressBar(len(kernels))

  ##############################################################################
  # Write Kernels
  ##############################################################################
  if globalParameters["MergeFiles"]:
    kernelSourceFile = open(os.path.join(outputPath, \
        "Kernels.cpp"), "w")
    kernelHeaderFile = open(os.path.join(outputPath, \
        "Kernels.h"), "w")
    kernelSourceFile.write(CHeader)
    kernelHeaderFile.write(CHeader)
    kernelSourceFile.write("#include \"Kernels.h\"\n")
    kernelHeaderFile.write("#pragma once\n")
    if globalParameters["RuntimeLanguage"] == "HIP":
      kernelHeaderFile.write("// Also set env var HCC_ENABLE_PRINTF=1 for printf\n")
      kernelHeaderFile.write("#define HCC_ENABLE_ACCELERATOR_PRINTF\n\n")
      kernelHeaderFile.write("#include <hip/hip_runtime.h>\n")
      kernelHeaderFile.write("#include \"TensileTypes.h\"\n")
      kernelHeaderFile.write("#include \"KernelHeader.h\"\n")
      kernelHeaderFile.write("\n\n")
      kernelHeaderFile.write("__device__ inline int GenDot4(int a, int b, int c) { \n")
      kernelHeaderFile.write("  typedef struct { int c0:8,c1:8,c2:8,c3:8; } C4I8;\n")
      kernelHeaderFile.write("  typedef union { int32_t i; C4I8 z; } PkInt8x4;\n")
      kernelHeaderFile.write("  PkInt8x4 va, vb; va.i = a; vb.i = b;\n")

      kernelHeaderFile.write("//if (__oclc_ISA_version == 906)\n")
      kernelHeaderFile.write("//{\n")
      kernelHeaderFile.write("//    return (float)__llvm_amdgcn_sdot4(a, b, c, true);\n")
      kernelHeaderFile.write("//}\n")
      kernelHeaderFile.write("//else\n")
      kernelHeaderFile.write("//{\n")
      kernelHeaderFile.write("      return c + (vb.z.c3*va.z.c3 + vb.z.c2*va.z.c2 + vb.z.c1*va.z.c1 + vb.z.c0*va.z.c0); }\n")
      kernelHeaderFile.write("//}\n")
      kernelHeaderFile.write("\n\n")
    else:
      kernelHeaderFile.write("#include <string>\n")

  kernelsWithBuildErrs = {}
<commit_msg>call amd_mixed_dot in ROCm 2.2<commit_after>################################################################################
# Copyright (C) 2016 Advanced Micro Devices, Inc. All rights reserved.
#
# Permission is hereby granted, free of charge, to any person obtaining a copy
# of this software and associated documentation files (the "Software"), to deal
# in the Software without restriction, including without limitation the rights
# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell cop-
# ies of the Software, and to permit persons to whom the Software is furnished
# to do so, subject to the following conditions:
#
# The above copyright notice and this permission notice shall be included in all
# copies or substantial portions of the Software.
#
# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IM-
# PLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS
# FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR
# COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER
# IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNE-
# CTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
################################################################################
# This script only gets called by CMake
from Common import globalParameters, HR, print1, print2, printExit, ensurePath, CHeader, CMakeHeader, assignGlobalParameters, ProgressBar, listToInitializer
from SolutionStructs import Solution
import YAMLIO
from SolutionWriter import SolutionWriter
from KernelWriterSource import KernelWriterSource
from KernelWriterAssembly import KernelWriterAssembly
import multiprocessing

import os
import sys
import os.path
import argparse
from shutil import copy as shutil_copy
import time


################################################################################
# Process a single kernel, return results:
################################################################################
def processKernelSource(kernel, kernelWriterSource, kernelWriterAssembly):
    kernelWriter = kernelWriterSource if kernel["KernelLanguage"] == "Source" else kernelWriterAssembly
    # get kernel name
    kernelName = kernelWriter.getKernelName(kernel)
    #sys.stderr.write("kernel:%s\n"% kernelName)
    (err, src) = kernelWriter.getSourceFileString(kernel)

    header = kernelWriter.getHeaderFileString(kernel)

    return (err, src, header, kernelName)


################################################################################
# Process a range of kernels:
# kLock is for the kernel file - is this used?
# pLock is for the progress bar

################################################################################
def processKernelSourceChunk(kernels,
                             kernelWriterSource, kernelWriterAssembly, \
                             kiStart, kiStop, pipe):

    results = []

    for ki in range(kiStart, kiStop):
      kernel = kernels[ki]
      results.append (processKernelSource(kernel, kernelWriterSource, kernelWriterAssembly)) # returns err, src, header, kernelName

    if pipe != None:
      pipe.send(results)
    else:
      return results



# create and prepare the assembly directory  - called ONCE per output dir:
def prepAsm():
  asmPath = ensurePath(os.path.join(globalParameters["WorkingPath"], "assembly") )
  assemblerFileName = os.path.join(asmPath, \
      "asm.%s"%("bat" if os.name=="nt" else "sh"))
  assemblerFile = open(assemblerFileName, "w")
  if os.name == "nt":
    assemblerFile.write("echo Windows: Copying instead of Assembling\n")
    assemblerFile.write("copy %1.s %1.o\n")
    assemblerFile.write("copy %1.o %1.co\n")
  else:
    assemblerFile.write("#!/bin/sh %s\n" % ("-x" if globalParameters["PrintLevel"] >=2  else ""))
    assemblerFile.write("# usage: asm.sh kernelName ASM_ARGS\n")
    assemblerFile.write("# example: asm.sh kernelName -mcpu=gfx900\n")
    assemblerFile.write("f=$1\n")
    assemblerFile.write("shift\n")
    assemblerFile.write("ASM=%s\n"%globalParameters["AssemblerPath"])
    # cannot use globalParameters["CurrentISA"] because it might be (0,0,0)
    defaultIsa = (9,0,0)
    assemblerFile.write( \
      "${ASM} -x assembler -target amdgcn--amdhsa %s $@ -c -o $f.o $f.s\n" % \
      ("-mno-code-object-v3" if \
      globalParameters["AsmCaps"][defaultIsa]["HasCodeObjectV3"] else ""))
    assemblerFile.write("${ASM} -target amdgcn--amdhsa $f.o -o $f.co\n")
  assemblerFile.close()
  os.chmod(assemblerFileName, 0777)

################################################################################
# processResults
# input: results is list with (err, src, header, kernelName)
# Log errors and write appropriate info to kernelSourceFile and kernelHeaderFile
################################################################################
def processResults(results, outputPath, kernelsWithBuildErrs, \
      kernelSourceFile, kernelHeaderFile):

  for (err,src,header,kernelName) in results:
    if err:
      kernelsWithBuildErrs[kernelName] = err
      #print "*** warning: invalid kernel#%s"%kernelName

    # write kernel.cpp
    if not globalParameters["MergeFiles"]:
      kernelSourceFile = open(os.path.join(outputPath, \
          "Kernels", kernelName+".cpp"), "w")
      kernelSourceFile.write(CHeader)

    kernelSourceFile.write(src)

    if not globalParameters["MergeFiles"]:
      kernelSourceFile.close()
      # write kernel.h
      kernelHeaderFile = open(os.path.join(outputPath, \
          "Kernels", kernelName+".h"), "w")
      kernelHeaderFile.write(CHeader)

    kernelHeaderFile.write(header)

    if not globalParameters["MergeFiles"]:
      kernelHeaderFile.close()

################################################################################
# Write Solutions and Kernels for BenchmarkClient or LibraryClient
################################################################################
def writeSolutionsAndKernels(outputPath, problemTypes, solutions, kernels, kernelsBetaOnly, \
    solutionWriter, kernelWriterSource, kernelWriterAssembly):
  start = time.time()
  print1("# Writing Kernels...")
  if not globalParameters["MergeFiles"]:
    ensurePath(os.path.join(outputPath, "Solutions"))
    ensurePath(os.path.join(outputPath, "Kernels"))

  if globalParameters["ShowProgressBar"]:
    progressBar = ProgressBar(len(kernels))

  ##############################################################################
  # Write Kernels
  ##############################################################################
  if globalParameters["MergeFiles"]:
    kernelSourceFile = open(os.path.join(outputPath, \
        "Kernels.cpp"), "w")
    kernelHeaderFile = open(os.path.join(outputPath, \
        "Kernels.h"), "w")
    kernelSourceFile.write(CHeader)
    kernelHeaderFile.write(CHeader)
    kernelSourceFile.write("#include \"Kernels.h\"\n")
    kernelHeaderFile.write("#pragma once\n")
    if globalParameters["RuntimeLanguage"] == "HIP":
      kernelHeaderFile.write("// Also set env var HCC_ENABLE_PRINTF=1 for printf\n")
      kernelHeaderFile.write("#define HCC_ENABLE_ACCELERATOR_PRINTF\n\n")
      kernelHeaderFile.write("#include <hip/hip_runtime.h>\n")
      kernelHeaderFile.write("#include \"TensileTypes.h\"\n")
      kernelHeaderFile.write("#include \"KernelHeader.h\"\n")
      kernelHeaderFile.write("\n\n")
      kernelHeaderFile.write("__device__ inline int GenDot4(int a, int b, int c) { \n")
      kernelHeaderFile.write("  typedef union { int32_t i; char4 z; } PkInt8x4;\n")
      kernelHeaderFile.write("  PkInt8x4 va, vb; va.i = a; vb.i = b;\n")

      kernelHeaderFile.write("//if (__oclc_ISA_version == 906)\n")
      kernelHeaderFile.write("//{\n")
      kernelHeaderFile.write("//    return (float)__llvm_amdgcn_sdot4(a, b, c, true);\n")
      kernelHeaderFile.write("//}\n")
      kernelHeaderFile.write("//else\n")
      kernelHeaderFile.write("//{\n")
      kernelHeaderFile.write("      return amd_mixed_dot(va.z, vb.z, c, true); }\n")
      kernelHeaderFile.write("//}\n")
      kernelHeaderFile.write("\n\n")
    else:
      kernelHeaderFile.write("#include <string>\n")

  kernelsWithBuildErrs = {}

{'fraction_non_alphanumeric': 0.21932413710017468, 'fraction_numerical': 0.006204445515330402, 'mean_word_length': 4.623983739837398, 'pattern_counts': {'":': 4, '<': 7, '<?xml version=': 0, '>': 9, 'https://': 0, 'lorem ipsum': 0, 'www.': 0, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 0, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '753805', 'n_tokens_mistral': 4746, 'n_tokens_neox': 4533, 'n_words': 1221}