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.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__)
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.
*
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");
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;
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;