From 88979fb6641bafa91a5347345285581c329f7a56 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 24 Feb 2023 07:09:36 +0530 Subject: [PATCH 1/5] SWDEV-370742 - Check image support before using image APIs. (#3143) Change-Id: I2fc086256bca605385f2f6a94119b46741020afa [ROCm/hip commit: 24cc1bc2bba9147550b9cbfa41a2d42f917c5cfe] --- projects/hip/tests/src/runtimeApi/memory/hipMemcpy2D.cpp | 1 + projects/hip/tests/src/runtimeApi/memory/hipMemcpy3D.cpp | 1 + projects/hip/tests/src/runtimeApi/memory/hipMemcpyAtoH.cpp | 1 + projects/hip/tests/src/runtimeApi/memory/hipMemcpyHtoA.cpp | 1 + projects/hip/tests/src/runtimeApi/memory/hipMemset2D.cpp | 1 + projects/hip/tests/src/runtimeApi/memory/hipMemset3D.cpp | 1 + 6 files changed, 6 insertions(+) diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemcpy2D.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemcpy2D.cpp index 328bf4b6e7..b1ec178b4c 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipMemcpy2D.cpp +++ b/projects/hip/tests/src/runtimeApi/memory/hipMemcpy2D.cpp @@ -315,6 +315,7 @@ bool Memcpy2D::Memcpy2D_NegativeTest_SizeCheck() { int main(int argc, char* argv[]) { bool TestPassed = true; + checkImageSupport(); Memcpy2D Memcpy2DObj; HipTest::parseStandardArguments(argc, argv, false); if (p_tests == 1) { diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemcpy3D.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemcpy3D.cpp index d7e06299ab..59277a2bc0 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipMemcpy3D.cpp +++ b/projects/hip/tests/src/runtimeApi/memory/hipMemcpy3D.cpp @@ -100,6 +100,7 @@ void runTest(int width,int height,int depth, hipChannelFormatKind formatKind) int main(int argc, char **argv) { + checkImageSupport(); for(int i=1;i<25;i++) { runTest(i,i,i, hipChannelFormatKindFloat); diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyAtoH.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyAtoH.cpp index 08d3e6b66f..f7c6f34adf 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyAtoH.cpp +++ b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyAtoH.cpp @@ -202,6 +202,7 @@ bool MemcpyAtoH::hipMemcpyAtoH_NegativeTests() { int main(int argc, char **argv) { bool TestPassed = true; + checkImageSupport(); HipTest::parseStandardArguments(argc, argv, false); MemcpyAtoH AtoH_obj; if (p_tests == 1) { diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyHtoA.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyHtoA.cpp index 245ac2f13c..3dc15181bb 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipMemcpyHtoA.cpp +++ b/projects/hip/tests/src/runtimeApi/memory/hipMemcpyHtoA.cpp @@ -214,6 +214,7 @@ bool MemcpyHtoA::hipMemcpyHtoA_NegativeTests() { int main(int argc, char **argv) { bool TestPassed = true; + checkImageSupport(); HipTest::parseStandardArguments(argc, argv, false); MemcpyHtoA HtoA_obj; if (p_tests == 1) { diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemset2D.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemset2D.cpp index 140a18f4a8..15b762476f 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipMemset2D.cpp +++ b/projects/hip/tests/src/runtimeApi/memory/hipMemset2D.cpp @@ -204,6 +204,7 @@ int main(int argc, char *argv[]) { int extraArgs = 0; bool testResult = true; + checkImageSupport(); HIPCHECK(hipSetDevice(p_gpuDevice)); extraArgs = HipTest::parseStandardArguments(argc, argv, false); parseExtraArguments(extraArgs, argv); diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemset3D.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemset3D.cpp index e307303fa0..e4662f4d7d 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipMemset3D.cpp +++ b/projects/hip/tests/src/runtimeApi/memory/hipMemset3D.cpp @@ -132,6 +132,7 @@ bool testhipMemset3DAsync(int memsetval,int p_gpuDevice) int main(int argc, char *argv[]) { + checkImageSupport(); HipTest::parseStandardArguments(argc, argv, true); HIPCHECK(hipSetDevice(p_gpuDevice)); bool testResult = true; From 985bdb8abf74d5fa1c5253aa0cd0de63e3534166 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 24 Feb 2023 07:10:26 +0530 Subject: [PATCH 2/5] SWDEV-370743 - Change the hard coded CU mask to support higher CUs. (#3144) Change-Id: I87ac149a6144959b46963ed8c867e6d3b6bca462 [ROCm/hip commit: 5c3e4fba39e1690dddc11d0da1c3dccb00d52178] --- projects/hip/tests/src/runtimeApi/stream/hipStreamGetCUMask.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/tests/src/runtimeApi/stream/hipStreamGetCUMask.cpp b/projects/hip/tests/src/runtimeApi/stream/hipStreamGetCUMask.cpp index ec3722f91d..1ea05c7ec4 100644 --- a/projects/hip/tests/src/runtimeApi/stream/hipStreamGetCUMask.cpp +++ b/projects/hip/tests/src/runtimeApi/stream/hipStreamGetCUMask.cpp @@ -32,7 +32,7 @@ using namespace std; int main(int argc, char* argv[]) { hipStream_t stream; - vector cuMask(6); + vector cuMask(8); stringstream ss; int nGpu = 0; From bb5bc88311a024f0828a49fb8192ba1538560c21 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 24 Feb 2023 11:30:33 +0530 Subject: [PATCH 3/5] SWDEV-370743 - Adding 'accelerator' keyword in lspci grep for AMD supported devices. (#3142) Change-Id: I5a0c2f4a678c5c80a2f13afe74d2c723912b0858 [ROCm/hip commit: e0425b53e0ea4fbc1067ab5c4d3d649ded0c3846] --- .../hip/tests/src/runtimeApi/device/hipDeviceGetPCIBusId.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/tests/src/runtimeApi/device/hipDeviceGetPCIBusId.cpp b/projects/hip/tests/src/runtimeApi/device/hipDeviceGetPCIBusId.cpp index af25ae7228..8349375647 100644 --- a/projects/hip/tests/src/runtimeApi/device/hipDeviceGetPCIBusId.cpp +++ b/projects/hip/tests/src/runtimeApi/device/hipDeviceGetPCIBusId.cpp @@ -113,7 +113,7 @@ bool compareHipDeviceGetPCIBusIdWithLspci() { char const *command = "lspci -D | grep controller | grep NVIDIA | " "cut -d ' ' -f 1"; #else - char const *command = "lspci -D | grep controller | grep AMD/ATI | " + char const *command = "lspci -D | grep -e controller -e accelerator | grep AMD/ATI | " "cut -d ' ' -f 1"; #endif fpipe = popen(command, "r"); From 41e0f4d34d72f011892602ef9903a681d854c0bb Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 24 Feb 2023 15:10:09 +0530 Subject: [PATCH 4/5] SWDEV-372153 - Add hipStreamGetDevice API. Add documentation for hip_prof_gen.py. (#3149) Change-Id: I2cb650e3b2ccf5431f4ed7d7685c4c6532732d44 [ROCm/hip commit: 873911c60f9c7c1d13a70e414549e055b5953071] --- projects/hip/docs/markdown/hip_build.md | 20 + projects/hip/hip_prof_gen.py | 508 --------------------- projects/hip/include/hip/hip_runtime_api.h | 11 + 3 files changed, 31 insertions(+), 508 deletions(-) delete mode 100755 projects/hip/hip_prof_gen.py diff --git a/projects/hip/docs/markdown/hip_build.md b/projects/hip/docs/markdown/hip_build.md index 4bab7ee8e1..5fc7daa9b3 100755 --- a/projects/hip/docs/markdown/hip_build.md +++ b/projects/hip/docs/markdown/hip_build.md @@ -102,6 +102,26 @@ By default, release version of AMDHIP is built. After make install command, make sure HIP_PATH is pointed to $PWD/install/hip. +## Generating profiling header after adding/changing a HIP API + +When you add or change a HIP API, you might need to generate a new `hip_prof_str.h` header. This header is used by rocm tools to track HIP APIs like rocprofiler/roctracer etc. +To generate the header after your change, use the tool `hip_prof_gen.py` present in `hipamd/src`. + +Usage: + +`hip_prof_gen.py [-v] []` + +Flags: + + * -v - verbose messages + * -r - process source directory recursively + * -t - API types matching check + * --priv - private API check + * -e - on error exit mode + * -p - HIP_INIT_API macro patching mode + +Example Usage: `hip_prof_gen.py -v -p -t --priv /include/hip/hip_runtime_api.h /src /include/hip/amd_detail/hip_prof_str.h /include/hip/amd_detail/hip_prof_str.h.new` + ## Build HIP tests ### Build HIP directed tests diff --git a/projects/hip/hip_prof_gen.py b/projects/hip/hip_prof_gen.py deleted file mode 100755 index e43177b0ea..0000000000 --- a/projects/hip/hip_prof_gen.py +++ /dev/null @@ -1,508 +0,0 @@ -#!/usr/bin/python -# Copyright (c) 2019 - 2021 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 -# copies 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 -# IMPLIED, 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 CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN -# THE SOFTWARE. - -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: sys.stdout.write(msg + '\n') - -# 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 - - sys.stdout.write(msg + '\n') - sys.stderr.write(sys.argv[0] + msg +'\n') - -def fatal(msg): - error(msg) - if errexit: sys.exit(1) - -############################################################# -# Normalizing API name -def filtr_api_name(name): - name = re.sub(r'\s*$', r'', name); - return name - -# 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'\s*(\*+)\s*', r'\1 ', 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[filtr_api_name(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 [] => -# out - output map [] => [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 = filtr_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 [] => [(type, name), ...] -# opts_map - opts map [] => [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('#define HIP_PROF_VER 1\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('\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 inline 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( - 'typedef 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: - if arg_tuple[0] == "hipLimit_t": - f.write(' enum ' + arg_tuple[0] + ' ' + arg_tuple[1] + ';\n') - else: - f.write(' ' + arg_tuple[0] + ' ' + arg_tuple[1] + ';\n') - f.write(' } ' + name + ';\n') - f.write( - ' } args;\n' + - '} hip_api_data_t;\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[][0] - # name is args[][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 HIP_PROF_HIP_API_STRING\n') - f.write('#include \n'); - f.write('#include \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 // HIP_PROF_HIP_API_STRING\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] \n" + - " -v - verbose messages\n" + - " example:\n" + - " $ hipap.py hip/include/hip/amd_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': '', - 'hipExtModuleLaunchKernel': '' -} -# 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 list(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]; - args_list = list_api_args(args_str) - api_map[name] = args_list - 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) diff --git a/projects/hip/include/hip/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h index c5696d7d29..80bfd2891f 100644 --- a/projects/hip/include/hip/hip_runtime_api.h +++ b/projects/hip/include/hip/hip_runtime_api.h @@ -2229,6 +2229,17 @@ hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int* flags); * @see hipStreamCreateWithFlags */ hipError_t hipStreamGetPriority(hipStream_t stream, int* priority); +/** + * @brief Get the device assocaited with the stream + * + * @param[in] stream stream to be queried + * @param[out] hipDevice_t device associated with the stream + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorContextIsDestroyed, #hipErrorInvalidHandle, + * #hipErrorNotInitialized, #hipErrorDeinitialized, #hipErrorInvalidContext + * + * @see hipStreamCreate, hipStreamDestroy, hipDeviceGetStreamPriorityRange + */ +hipError_t hipStreamGetDevice(hipStream_t stream, hipDevice_t* device); /** * @brief Create an asynchronous stream with the specified CU mask. * From caa94251e6100115e0a5111bb43617da9e2b6985 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Sat, 25 Feb 2023 00:15:16 +0530 Subject: [PATCH 5/5] SWDEV-383461 - gfx11 not supporting wave64 check (#3170) Change-Id: I12e8ff2cb8ad9c72cba216847366f2e0ab15c497 [ROCm/hip commit: c629e630247c17dafb2bc5c85dad8770175e1d33] --- projects/hip/include/hip/hip_runtime.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/include/hip/hip_runtime.h b/projects/hip/include/hip/hip_runtime.h index a8e36d19db..0faf215176 100644 --- a/projects/hip/include/hip/hip_runtime.h +++ b/projects/hip/include/hip/hip_runtime.h @@ -37,8 +37,8 @@ THE SOFTWARE. #ifndef HIP_INCLUDE_HIP_HIP_RUNTIME_H #define HIP_INCLUDE_HIP_HIP_RUNTIME_H -#if (__gfx1010__ || __gfx1011__ || __gfx1012__ || __gfx1030__ || __gfx1031__) && __AMDGCN_WAVEFRONT_SIZE == 64 -#error HIP is not supported on GFX10 with wavefront size 64 +#if __HIP_DEVICE_COMPILE__ && !__GFX8__ && !__GFX9__ && __AMDGCN_WAVEFRONT_SIZE == 64 +#error HIP is not supported on the specified GPU ARCH with wavefront size 64 #endif #if !defined(__HIPCC_RTC__)