AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge branch 'master' into amd-master-next

Browse Source 
Change-Id: Ibb7775e7acd263d2ece40a241517bbd15976fdd4
This commit is contained in:
Laurent Morichetti
2020-02-12 00:14:07 -08:00
parent 6462732e49 ab7526f64c
commit 2cf3257795
38 changed files with 602 additions and 1010 deletions
Download Patch File Download Diff File Expand all files Collapse all files
tests/README.md
+15 -6
View File
@@ -47,7 +47,7 @@ In the above, BUILD commands provide instructions on how to build the test case
The supported syntax for the BUILD command is:
```
BUILD: %t %s HIPCC_OPTIONS <hipcc_specific_options> HCC_OPTIONS <hcc_specific_options> NVCC_OPTIONS <nvcc_specific_options> EXCLUDE_HIP_PLATFORM <hcc|nvcc|all> DEPENDS <dependencies>
BUILD: %t %s HIPCC_OPTIONS <hipcc_specific_options> HCC_OPTIONS <hcc_specific_options> NVCC_OPTIONS <nvcc_specific_options> EXCLUDE_HIP_PLATFORM <hcc|nvcc|all> EXCLUDE_HIP_RUNTIME <HCC|VDI> EXCLUDE_HIP_COMPILER <hcc|clang> DEPENDS <dependencies>
```
%s: refers to current source file name. Additional source files needed for the test can be specified by name (including relative path).
%t: refers to target executable named derived by removing the extension from the current source file. Alternatively a target executable name can be specified.
@@ -55,6 +55,9 @@ HIPCC_OPTIONS: All options specified after this delimiter are passed to hipcc on
HCC_OPTIONS: All options specified after this delimiter are passed to hipcc on HCC platform only.
NVCC_OPTIONS: All options specified after this delimiter are passed to hipcc on NVCC platform only.
EXCLUDE_HIP_PLATFORM: This can be used to exclude a test case from HCC, NVCC or both platforms.
EXCLUDE_HIP_RUNTIME: This can be used to exclude a test case from HCC or VDI runtime.
EXCLUDE_HIP_COMPILER: This can be used to exclude a test case from hcc or clang compiler.
EXCLUDE_HIP_RUNTIME AND EXCLUDE_HIP_COMPILER: when both options are specified it excludes test case from particular runtime and compiler.
DEPENDS: This can be used to specify dependencies that need to be built before building the current target.
@@ -62,7 +65,7 @@ DEPENDS: This can be used to specify dependencies that need to be built before b
The supported syntax for the BUILD_CMD command is:
```
BUILD_CMD: <targetname> <build_command> EXCLUDE_HIP_PLATFORM <hcc|nvcc|all> DEPENDS <dependencies>
BUILD_CMD: <targetname> <build_command> EXCLUDE_HIP_PLATFORM <hcc|nvcc|all> EXCLUDE_HIP_RUNTIME <HCC|VDI> EXCLUDE_HIP_COMPILER <hcc|clang> DEPENDS <dependencies>
```
%s: refers to current source file name. Additional source files needed for the test can be specified by name (including relative path).
%t: refers to target executable named derived by removing the extension from the current source file. Alternatively a target executable name can be specified.
@@ -73,6 +76,9 @@ BUILD_CMD: <targetname> <build_command> EXCLUDE_HIP_PLATFORM <hcc|nvcc|all> DEPE
%S: refers to path to current source file.
%T: refers to path to current build target.
EXCLUDE_HIP_PLATFORM: This can be used to exclude a test case from HCC, NVCC or both platforms.
EXCLUDE_HIP_RUNTIME: This can be used to exclude a test case from HCC or VDI runtime.
EXCLUDE_HIP_COMPILER: This can be used to exclude a test case from hcc or clang compiler.
EXCLUDE_HIP_RUNTIME AND EXCLUDE_HIP_COMPILER: when both options are specified it excludes test from particular runtime and compiler.
DEPENDS: This can be used to specify dependencies that need to be built before building the current target.
@@ -80,17 +86,20 @@ DEPENDS: This can be used to specify dependencies that need to be built before b
The supported syntax for the TEST command is:
```
TEST: %t <arguments_to_test_executable> EXCLUDE_HIP_PLATFORM <hcc|nvcc|all>
TEST: %t <arguments_to_test_executable> EXCLUDE_HIP_PLATFORM <hcc|nvcc|all> EXCLUDE_HIP_RUNTIME <HCC|VDI> EXCLUDE_HIP_COMPILER <hcc|clang>
```
%t: refers to target executable named derived by removing the extension from the current source file. Alternatively a target executable name can be specified.
EXCLUDE_HIP_PLATFORM: This can be used to exclude a test case from HCC, NVCC or both platforms. Note that if the test has been excluded for a specific platform in the BUILD command, it is automatically excluded from the TEST command as well for the same platform.
EXCLUDE_HIP_PLATFORM: This can be used to exclude a test case from HCC, NVCC or both platforms.
EXCLUDE_HIP_RUNTIME: This can be used to exclude a test case from HCC or VDI runtime.
EXCLUDE_HIP_COMPILER: This can be used to exclude a test case from hcc or clang compiler.
EXCLUDE_HIP_RUNTIME AND EXCLUDE_HIP_COMPILER: when both options are specified it excludes test from particular runtime and compiler.
Note that if the test has been excluded for a specific platform/runtime/compiler in the BUILD command, it is automatically excluded from the TEST command as well for the sameplatform.
#### TEST_NAMED command
When using the TEST command, HIT will squash and append the arguments specified to the test executable name to generate the CMAKE test name. Sometimes we might want to specify a more descriptive name. The TEST_NAMED command is used for that. The supported syntax for the TEST_NAMED command is:
```
TEST: %t CMAKE_TEST_NAME <arguments_to_test_executable> EXCLUDE_HIP_PLATFORM <hcc|nvcc|all>
TEST: %t CMAKE_TEST_NAME <arguments_to_test_executable> EXCLUDE_HIP_PLATFORM <hcc|nvcc|all> EXCLUDE_HIP_RUNTIME <HCC|VDI> EXCLUDE_HIP_COMPILER <hcc|clang>
```
tests/hit/HIT.cmake
+126 -16
View File
@@ -3,7 +3,7 @@ find_package(HIP REQUIRED)
#-------------------------------------------------------------------------------
# Helper macro to parse BUILD instructions
macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_options _link_options _exclude_platforms _depends _dir)
macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_options _link_options _exclude_platforms _exclude_runtime _exclude_compiler _depends _dir)
set(${_target})
set(${_sources})
set(${_hipcc_options})
@@ -11,6 +11,8 @@ macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_opt
set(${_nvcc_options})
set(${_link_options})
set(${_exclude_platforms})
set(${_exclude_runtime})
set(${_exclude_compiler})
set(${_depends})
set(_target_found FALSE)
set(_hipcc_options_found FALSE)
@@ -18,6 +20,8 @@ macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_opt
set(_nvcc_options_found FALSE)
set(_link_options_found FALSE)
set(_exclude_platforms_found FALSE)
set(_exclude_runtime_found FALSE)
set(_exclude_compiler_found FALSE)
set(_depends_found FALSE)
foreach(arg ${ARGN})
if(NOT _target_found)
@@ -29,6 +33,8 @@ macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_opt
set(_nvcc_options_found FALSE)
set(_link_options_found FALSE)
set(_exclude_platforms_found FALSE)
set(_exclude_runtime_found FALSE)
set(_exclude_compiler_found FALSE)
set(_depends_found FALSE)
elseif("x${arg}" STREQUAL "xHCC_OPTIONS")
set(_hipcc_options_found FALSE)
@@ -36,6 +42,8 @@ macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_opt
set(_nvcc_options_found FALSE)
set(_link_options_found FALSE)
set(_exclude_platforms_found FALSE)
set(_exclude_runtime_found FALSE)
set(_exclude_compiler_found FALSE)
set(_depends_found FALSE)
elseif("x${arg}" STREQUAL "xNVCC_OPTIONS")
set(_hipcc_options_found FALSE)
@@ -43,6 +51,8 @@ macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_opt
set(_nvcc_options_found TRUE)
set(_link_options_found FALSE)
set(_exclude_platforms_found FALSE)
set(_exclude_runtime_found FALSE)
set(_exclude_compiler_found FALSE)
set(_depends_found FALSE)
elseif("x${arg}" STREQUAL "xLINK_OPTIONS")
set(_hipcc_options_found FALSE)
@@ -50,6 +60,8 @@ macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_opt
set(_nvcc_options_found FALSE)
set(_link_options_found TRUE)
set(_exclude_platforms_found FALSE)
set(_exclude_runtime_found FALSE)
set(_exclude_compiler_found FALSE)
set(_depends_found FALSE)
elseif("x${arg}" STREQUAL "xEXCLUDE_HIP_PLATFORM")
set(_hipcc_options_found FALSE)
@@ -57,6 +69,26 @@ macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_opt
set(_nvcc_options_found FALSE)
set(_link_options_found FALSE)
set(_exclude_platforms_found TRUE)
set(_exclude_runtime_found FALSE)
set(_exclude_compiler_found FALSE)
set(_depends_found FALSE)
elseif("x${arg}" STREQUAL "xEXCLUDE_HIP_RUNTIME")
set(_hipcc_options_found FALSE)
set(_hcc_options_found FALSE)
set(_nvcc_options_found FALSE)
set(_link_options_found FALSE)
set(_exclude_platforms_found FALSE)
set(_exclude_runtime_found TRUE)
set(_exclude_compiler_found FALSE)
set(_depends_found FALSE)
elseif("x${arg}" STREQUAL "xEXCLUDE_HIP_COMPILER")
set(_hipcc_options_found FALSE)
set(_hcc_options_found FALSE)
set(_nvcc_options_found FALSE)
set(_link_options_found FALSE)
set(_exclude_platforms_found FALSE)
set(_exclude_runtime_found FALSE)
set(_exclude_compiler_found TRUE)
set(_depends_found FALSE)
elseif("x${arg}" STREQUAL "xDEPENDS")
set(_hipcc_options_found FALSE)
@@ -64,6 +96,8 @@ macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_opt
set(_nvcc_options_found FALSE)
set(_link_options_found FALSE)
set(_exclude_platforms_found FALSE)
set(_exclude_runtime_found FALSE)
set(_exclude_compiler_found FALSE)
set(_depends_found TRUE)
else()
if(_hipcc_options_found)
@@ -76,6 +110,10 @@ macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_opt
list(APPEND ${_link_options} ${arg})
elseif(_exclude_platforms_found)
list(APPEND ${_exclude_platforms} ${arg})
elseif(_exclude_runtime_found)
list(APPEND ${_exclude_runtime} ${arg})
elseif(_exclude_compiler_found)
list(APPEND ${_exclude_compiler} ${arg})
elseif(_depends_found)
list(APPEND ${_depends} ${arg})
else()
@@ -86,13 +124,15 @@ macro(PARSE_BUILD_COMMAND _target _sources _hipcc_options _hcc_options _nvcc_opt
endmacro()
# Helper macro to parse CUSTOM BUILD instructions
macro(PARSE_CUSTOMBUILD_COMMAND _target _buildcmd _exclude_platforms _depends)
macro(PARSE_CUSTOMBUILD_COMMAND _target _buildcmd _exclude_platforms _exclude_runtime _exclude_compiler _depends)
set(${_target})
set(${_buildcmd} " ")
set(${_exclude_platforms})
set(${_depends})
set(_target_found FALSE)
set(_exclude_platforms_found FALSE)
set(_exclude_runtime_found FALSE)
set(_exclude_compiler_found FALSE)
set(_depends_found FALSE)
foreach(arg ${ARGN})
if(NOT _target_found)
@@ -100,13 +140,31 @@ macro(PARSE_CUSTOMBUILD_COMMAND _target _buildcmd _exclude_platforms _depends)
set(${_target} ${arg})
elseif("x${arg}" STREQUAL "xEXCLUDE_HIP_PLATFORM")
set(_exclude_platforms_found TRUE)
set(_exclude_runtime_found FALSE)
set(_exclude_compiler_found FALSE)
set(_depends_found FALSE)
elseif("x${arg}" STREQUAL "xEXCLUDE_HIP_RUNTIME")
set(_exclude_platforms_found FALSE)
set(_exclude_runtime_found TRUE)
set(_exclude_compiler_found FALSE)
set(_depends_found FALSE)
elseif("x${arg}" STREQUAL "xEXCLUDE_HIP_COMPILER")
set(_exclude_platforms_found FALSE)
set(_exclude_runtime_found FALSE)
set(_exclude_compiler_found TRUE)
set(_depends_found FALSE)
elseif("x${arg}" STREQUAL "xDEPENDS")
set(_exclude_platforms_found FALSE)
set(_exclude_runtime_found FALSE)
set(_exclude_compiler_found FALSE)
set(_depends_found TRUE)
else()
if(_exclude_platforms_found)
list(APPEND ${_exclude_platforms} ${arg})
elseif(_exclude_runtime_found)
list(APPEND ${_exclude_runtime} ${arg})
elseif(_exclude_compiler_found)
list(APPEND ${_exclude_compiler} ${arg})
elseif(_depends_found)
list(APPEND ${_depends} ${arg})
else()
@@ -117,21 +175,39 @@ macro(PARSE_CUSTOMBUILD_COMMAND _target _buildcmd _exclude_platforms _depends)
endmacro()
# Helper macro to parse TEST instructions
macro(PARSE_TEST_COMMAND _target _arguments _exclude_platforms)
macro(PARSE_TEST_COMMAND _target _arguments _exclude_platforms _exclude_runtime _exclude_compiler)
set(${_target})
set(${_arguments} " ")
set(${_exclude_platforms})
set(${_exclude_runtime})
set(${_exclude_compiler})
set(_target_found FALSE)
set(_exclude_platforms_found FALSE)
set(_exclude_runtime_found FALSE)
set(_exclude_compiler_found FALSE)
foreach(arg ${ARGN})
if(NOT _target_found)
set(_target_found TRUE)
set(${_target} ${arg})
elseif("x${arg}" STREQUAL "xEXCLUDE_HIP_PLATFORM")
set(_exclude_platforms_found TRUE)
set(_exclude_runtime_found FALSE)
set(_exclude_compiler_found FALSE)
elseif("x${arg}" STREQUAL "xEXCLUDE_HIP_RUNTIME")
set(_exclude_platforms_found FALSE)
set(_exclude_runtime_found TRUE)
set(_exclude_compiler_found FALSE)
elseif("x${arg}" STREQUAL "xEXCLUDE_HIP_COMPILER")
set(_exclude_platforms_found FALSE)
set(_exclude_runtime_found FALSE)
set(_exclude_compiler_found TRUE)
else()
if(_exclude_platforms_found)
list(APPEND ${_exclude_platforms} ${arg})
elseif(_exclude_runtime_found)
list(APPEND ${_exclude_runtime} ${arg})
elseif(_exclude_compiler_found)
list(APPEND ${_exclude_compiler} ${arg})
else()
list(APPEND ${_arguments} ${arg})
endif()
@@ -140,13 +216,15 @@ macro(PARSE_TEST_COMMAND _target _arguments _exclude_platforms)
endmacro()
# Helper macro to parse TEST_NAMED instructions
macro(PARSE_TEST_NAMED_COMMAND _target _testname _arguments _exclude_platforms)
macro(PARSE_TEST_NAMED_COMMAND _target _testname _arguments _exclude_platforms _exclude_runtime _exclude_compiler)
set(${_target})
set(${_arguments} " ")
set(${_exclude_platforms})
set(_target_found FALSE)
set(_testname_found FALSE)
set(_exclude_platforms_found FALSE)
set(_exclude_runtime_found FALSE)
set(_exclude_compiler_found FALSE)
foreach(arg ${ARGN})
if(NOT _target_found)
set(_target_found TRUE)
@@ -156,9 +234,23 @@ macro(PARSE_TEST_NAMED_COMMAND _target _testname _arguments _exclude_platforms)
set(${_testname} ${arg})
elseif("x${arg}" STREQUAL "xEXCLUDE_HIP_PLATFORM")
set(_exclude_platforms_found TRUE)
set(_exclude_runtime_found FALSE)
set(_exclude_compiler_found FALSE)
elseif("x${arg}" STREQUAL "xEXCLUDE_HIP_RUNTIME")
set(_exclude_platforms_found FALSE)
set(_exclude_runtime_found TRUE)
set(_exclude_compiler_found FALSE)
elseif("x${arg}" STREQUAL "xEXCLUDE_HIP_COMPILER")
set(_exclude_platforms_found FALSE)
set(_exclude_runtime_found FALSE)
set(_exclude_compiler_found TRUE)
else()
if(_exclude_platforms_found)
list(APPEND ${_exclude_platforms} ${arg})
elseif(_exclude_runtime_found)
list(APPEND ${_exclude_runtime} ${arg})
elseif(_exclude_compiler_found)
list(APPEND ${_exclude_compiler} ${arg})
else()
list(APPEND ${_arguments} ${arg})
endif()
@@ -203,10 +295,16 @@ macro(HIT_ADD_FILES _dir _label _parent)
string(REGEX REPLACE "\n" ";" _contents "${_contents}")
foreach(_cmd ${_contents})
string(REGEX REPLACE " " ";" _cmd "${_cmd}")
parse_build_command(_target _sources _hipcc_options _hcc_options _nvcc_options _link_options _exclude_platforms _depends ${_dir} ${_cmd})
parse_build_command(_target _sources _hipcc_options _hcc_options _nvcc_options _link_options _exclude_platforms _exclude_runtime _exclude_compiler _depends ${_dir} ${_cmd})
string(REGEX REPLACE "/" "." target ${_label}/${_target})
insert_into_map("_exclude" "${target}" "${_exclude_platforms}")
if("all" IN_LIST _exclude_platforms OR ${HIP_PLATFORM} IN_LIST _exclude_platforms)
insert_into_map("_exclude" "${target}" TRUE)
elseif(NOT _exclude_runtime AND ${HIP_COMPILER} IN_LIST _exclude_compiler)
insert_into_map("_exclude" "${target}" TRUE)
elseif(NOT _exclude_compiler AND ${HIP_RUNTIME} IN_LIST _exclude_runtime)
insert_into_map("_exclude" "${target}" TRUE)
elseif(_exclude_runtime STREQUAL ${HIP_RUNTIME} AND ${HIP_COMPILER} IN_LIST _exclude_compiler)
insert_into_map("_exclude" "${target}" TRUE)
else()
set_source_files_properties(${_sources} PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1)
hip_reset_flags()
@@ -236,10 +334,16 @@ macro(HIT_ADD_FILES _dir _label _parent)
string(REGEX REPLACE "%T" ${_label} _contents "${_contents}")
foreach(_cmd ${_contents})
string(REGEX REPLACE " " ";" _cmd "${_cmd}")
parse_custombuild_command(_target _buildcmd _exclude_platforms _depends ${_cmd})
parse_custombuild_command(_target _buildcmd _exclude_platforms _exclude_runtime _exclude_compiler _depends ${_cmd})
string(REGEX REPLACE "/" "." target ${_label}/${_target})
insert_into_map("_exclude" "${target}" "${_exclude_platforms}")
if("all" IN_LIST _exclude_platforms OR ${HIP_PLATFORM} IN_LIST _exclude_platforms)
insert_into_map("_exclude" "${target}" TRUE)
elseif(NOT _exclude_runtime AND ${HIP_COMPILER} IN_LIST _exclude_compiler)
insert_into_map("_exclude" "${target}" TRUE)
elseif(NOT _exclude_compiler AND ${HIP_RUNTIME} IN_LIST _exclude_runtime)
insert_into_map("_exclude" "${target}" TRUE)
elseif(_exclude_runtime STREQUAL ${HIP_RUNTIME} AND ${HIP_COMPILER} IN_LIST _exclude_compiler)
insert_into_map("_exclude" "${target}" TRUE)
else()
string(REGEX REPLACE ";" " " _buildcmd "${_buildcmd}")
#string(CONCAT buildscript ${CMAKE_CURRENT_BINARY_DIR}/${target} ".sh")
@@ -263,11 +367,14 @@ macro(HIT_ADD_FILES _dir _label _parent)
string(REGEX REPLACE "\n" ";" _contents "${_contents}")
foreach(_cmd ${_contents})
string(REGEX REPLACE " " ";" _cmd "${_cmd}")
parse_test_command(_target _arguments _exclude_platforms ${_cmd})
parse_test_command(_target _arguments _exclude_platforms _exclude_runtime _exclude_compiler ${_cmd})
string(REGEX REPLACE "/" "." target ${_label}/${_target})
read_from_map("_exclude" "${target}" _exclude_platforms_from_build)
if("all" IN_LIST _exclude_platforms OR ${HIP_PLATFORM} IN_LIST _exclude_platforms OR
_exclude_platforms_from_build STREQUAL "all" OR _exclude_platforms_from_build STREQUAL ${HIP_PLATFORM})
read_from_map("_exclude" "${target}" _exclude_test_from_build)
if("all" IN_LIST _exclude_platforms OR ${HIP_PLATFORM} IN_LIST _exclude_platforms)
elseif(NOT _exclude_runtime AND ${HIP_COMPILER} IN_LIST _exclude_compiler)
elseif(NOT _exclude_compiler AND ${HIP_RUNTIME} IN_LIST _exclude_runtime)
elseif(${HIP_RUNTIME} IN_LIST _exclude_runtime AND ${HIP_COMPILER} IN_LIST _exclude_compiler)
elseif(_exclude_test_from_build STREQUAL TRUE)
else()
make_test(${_label}/${_target} ${_arguments})
endif()
@@ -282,11 +389,14 @@ macro(HIT_ADD_FILES _dir _label _parent)
string(REGEX REPLACE "\n" ";" _contents "${_contents}")
foreach(_cmd ${_contents})
string(REGEX REPLACE " " ";" _cmd "${_cmd}")
parse_test_named_command(_target _testname _arguments _exclude_platforms ${_cmd})
parse_test_named_command(_target _testname _arguments _exclude_platforms _exclude_runtime _exclude_compiler ${_cmd})
string(REGEX REPLACE "/" "." target ${_label}/${_target})
read_from_map("_exclude" "${target}" _exclude_platforms_from_build)
if("all" IN_LIST _exclude_platforms OR ${HIP_PLATFORM} IN_LIST _exclude_platforms OR
_exclude_platforms_from_build STREQUAL "all" OR _exclude_platforms_from_build STREQUAL ${HIP_PLATFORM})
read_from_map("_exclude" "${target}" _exclude_test_from_build)
if("all" IN_LIST _exclude_platforms OR ${HIP_PLATFORM} IN_LIST _exclude_platforms)
elseif(NOT _exclude_runtime AND ${HIP_COMPILER} IN_LIST _exclude_compiler)
elseif(NOT _exclude_compiler AND ${HIP_RUNTIME} IN_LIST _exclude_runtime)
elseif(${HIP_RUNTIME} IN_LIST _exclude_runtime AND ${HIP_COMPILER} IN_LIST _exclude_compiler)
elseif(_exclude_test_from_build STREQUAL TRUE)
else()
make_named_test(${_label}/${_target} ${_label}/${_testname}.tst ${_arguments})
endif()
tests/src/gcc/LaunchKernel.c
+41 -12
View File
@@ -36,7 +36,7 @@ bool LaunchKernelArg()
dim3 blocks = {1,1,1};
dim3 threads = {1,1,1};
HIPCHECK(hipLaunchKernel(kernel, blocks, threads,NULL, 0, 0));
HIPCHECK(hipLaunchKernel(kernel, blocks, threads, NULL, 0, 0));
return true;
}
@@ -50,9 +50,9 @@ bool LaunchKernelArg1()
// Allocate Device memory
HIPCHECK(hipMalloc((void**)&A_d, sizeof(int)));
void* Args[]={A_d};
HIPCHECK(hipLaunchKernel(kernel1, blocks, threads, Args,0,0));
void* Args[]={&A_d};
HIPCHECK(hipLaunchKernel(kernel1, blocks, threads, Args, 0, 0));
// Get the result back to host memory
HIPCHECK(hipMemcpy(&A, A_d, sizeof(int), hipMemcpyDeviceToHost));
@@ -60,7 +60,7 @@ bool LaunchKernelArg1()
HIPCHECK(hipFree(A_d));
if(A != 333)
return false;
return false;
return true;
}
@@ -81,9 +81,9 @@ bool LaunchKernelArg2()
HIPCHECK(hipMalloc((void**)&B_d, sizeof(int)));
// Copy data from host memory to device memory
HIPCHECK(hipMemcpy(B_d,&B, sizeof(int), hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy(B_d, &B, sizeof(int), hipMemcpyHostToDevice));
void* Args[]={A_d,B_d};
void* Args[]={&A_d, &B_d};
HIPCHECK(hipLaunchKernel(kernel2, blocks, threads, Args,0,0));
// Get the result back to host memory
@@ -118,11 +118,11 @@ bool LaunchKernelArg3()
HIPCHECK(hipMalloc((void**)&C_d, sizeof(int)));
// Copy data from host memory to device memory
HIPCHECK(hipMemcpy(A_d,&A, sizeof(int), hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy(A_d, &A, sizeof(int), hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy(B_d,&B, sizeof(int), hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy(B_d, &B, sizeof(int), hipMemcpyHostToDevice));
void* Args[]={A_d,B_d,C_d};
void* Args[]={&A_d, &B_d, &C_d};
HIPCHECK(hipLaunchKernel(kernel3, blocks, threads, Args,0,0));
// Get the result back to host memory
@@ -138,14 +138,43 @@ bool LaunchKernelArg3()
return true;
}
bool LaunchKernelArg4()
{
int A = 0;
int *A_d = NULL;
dim3 blocks = {1,1,1};
dim3 threads = {1,1,1};
// Allocate Device memory
HIPCHECK(hipMalloc((void**)&A_d, sizeof(int)));
char c = 1;
short s = 10;
int i = 100;
struct things t = {2,20,200};
void* Args[]={&A_d, &c, &s, &i, &t};
HIPCHECK(hipLaunchKernel(kernel4, blocks, threads, Args, 0, 0));
// Get the result back to host memory
HIPCHECK(hipMemcpy(&A, A_d, sizeof(int), hipMemcpyDeviceToHost));
HIPCHECK(hipFree(A_d));
if (A != (c + s + i + t.c + t.s + t.i))
return false;
return true;
}
int main()
{
if( LaunchKernelArg() &&
LaunchKernelArg1() &&
LaunchKernelArg2() &&
LaunchKernelArg3())
LaunchKernelArg3() &&
LaunchKernelArg4())
{
printf("PASSED!\n");
}
tests/src/gcc/LaunchKernel.h
+16
View File
@@ -17,7 +17,23 @@
* THE SOFTWARE.
*/
#ifdef __cplusplus
extern "C"
{
#endif
extern __global__ void kernel();
extern __global__ void kernel1(int*);
extern __global__ void kernel2(int*,int*);
extern __global__ void kernel3(int*,int*,int*);
struct things {
char c;
short s;
int i;
};
extern __global__ void kernel4(int*, char, short, int, struct things);
#ifdef __cplusplus
}
#endif
tests/src/gcc/gpu.cpp
+6
View File
@@ -19,6 +19,7 @@
#include<hip/hip_runtime.h>
#include "LaunchKernel.h"
extern "C"
{
@@ -43,4 +44,9 @@ __global__ void kernel3(int *a, int*b, int* c)
*c = *a+*b;
}
__global__ void kernel4(int *a, char c, short s, int i, struct things t)
{
*a = c + s + i + t.c + t.s + t.i;
}
}//extern "C"
tests/src/runtimeApi/module/hipLaunchCoopMultiKernel.cpp
+1 -1
View File
@@ -156,7 +156,7 @@ int main() {
for (int i = 0; i < nGpu; i++) {
HIPCHECK(hipSetDevice(i));
dimBlock.x = workgroups[set];
HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocks,
HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(reinterpret_cast<uint32_t*>(&numBlocks),
(hipFunction_t)test_gws, dimBlock.x * dimBlock.y * dimBlock.z, dimBlock.x * sizeof(long)));
std::cout << "GPU(" << i << ") Block size: " << dimBlock.x << " Num blocks per CU: " << numBlocks << "\n";
tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp
+1 -1
View File
@@ -116,7 +116,7 @@ int main() {
dimBlock.x = workgroups[i];
// Calculate the device occupancy to know how many blocks can be run concurrently
hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocks,
hipOccupancyMaxActiveBlocksPerMultiprocessor(reinterpret_cast<uint32_t*>(&numBlocks),
test_gws, dimBlock.x * dimBlock.y * dimBlock.z, dimBlock.x * sizeof(long));
dimGrid.x = deviceProp.multiProcessorCount * std::min(numBlocks, 32);
tests/src/runtimeApi/occupancy/hipOccupancyMaxActiveBlocksPerMultiprocessor.cpp → tests/src/runtimeApi/module/hipOccupancyMaxActiveBlocksPerMultiprocessor.cpp
+18 -3
View File
@@ -30,6 +30,10 @@ THE SOFTWARE.
#include "hip/hip_runtime.h"
#include "test_common.h"
#define fileName "vcpy_kernel.code"
#define kernel_name "hello_world"
__global__ void f1(float *a) { *a = 1.0; }
template <typename T>
@@ -45,10 +49,11 @@ int main(int argc, char* argv[]) {
hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f1, 0, 0);
assert(gridSize != 0 && blockSize != 0);
int numBlock = 0;
hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, f1, (int)blockSize, 0);
uint32_t numBlock = 0;
hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, f1, blockSize, 0);
assert(numBlock != 0);
// test case for using kernel function pointer with template
gridSize = 0;
blockSize = 0;
@@ -56,7 +61,17 @@ int main(int argc, char* argv[]) {
assert(gridSize != 0 && blockSize != 0);
numBlock = 0;
hipOccupancyMaxActiveBlocksPerMultiprocessor<void(*)(int *)>(&numBlock, f2, (int)blockSize, 0);
hipOccupancyMaxActiveBlocksPerMultiprocessor<void(*)(int *)>(&numBlock, f2, blockSize, 0);
assert(numBlock != 0);
// test case for using kernel with hipFunction_t type
numBlock = 0;
hipModule_t Module;
hipFunction_t Function;
HIPCHECK(hipModuleLoad(&Module, fileName));
HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name));
HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, Function, blockSize, 0));
assert(numBlock != 0);
passed();
tests/src/runtimeApi/occupancy/hipOccupancyMaxPotentialBlockSize.cpp → tests/src/runtimeApi/module/hipOccupancyMaxPotentialBlockSize.cpp
+1
View File
@@ -33,6 +33,7 @@ THE SOFTWARE.
#define fileName "vcpy_kernel.code"
#define kernel_name "hello_world"
__global__ void f1(float *a) { *a = 1.0; }
template <typename T>
tests/src/texture/hipNormalizedFloatValueTex.cpp
+19 -3
View File
@@ -27,8 +27,24 @@ THE SOFTWARE.
*/
#include "test_common.h"
#define SIZE 10
static float getNormalizedValue(const float value,
const enum hipArray_Format texFormat) {
switch (texFormat) {
case HIP_AD_FORMAT_SIGNED_INT8:
return (value / SCHAR_MAX);
case HIP_AD_FORMAT_UNSIGNED_INT8:
return (value / UCHAR_MAX);
case HIP_AD_FORMAT_SIGNED_INT16:
return (value / SHRT_MAX);
case HIP_AD_FORMAT_UNSIGNED_INT16:
return (value / USHRT_MAX);
default:
return value;
}
}
texture<float, hipTextureType1D, hipReadModeElementType> textureNormalizedVal_1D;
__global__ void normalizedValTextureTest(unsigned int numElements, float* pDst)
@@ -47,7 +63,6 @@ bool textureTest(enum hipArray_Format texFormat)
T *dData = NULL;
HIPCHECK(hipMalloc((void **) &dData, sizeof(T)*SIZE));
HIPCHECK(hipMemcpyHtoD((hipDeviceptr_t)dData, hData, sizeof(T)*SIZE));
textureReference* texRef = &textureNormalizedVal_1D;
HIPCHECK(hipTexRefSetAddressMode(texRef, 0, hipAddressModeClamp));
HIPCHECK(hipTexRefSetAddressMode(texRef, 1, hipAddressModeClamp));
@@ -73,7 +88,8 @@ bool textureTest(enum hipArray_Format texFormat)
for(int i = 0; i < SIZE; i++)
{
if((float)hData[i]/texFormatToSize[texFormat] != hOutputData[i])
float expected = getNormalizedValue(float(hData[i]), texFormat);
if(expected != hOutputData[i])
{
printf("mismatch at index:%d for texType:%d output:%f\n",i,texFormat,hOutputData[i]);
testResult = false;