SWDEV-299127 - Merge 'develop' into 'amd-staging'

Change-Id: I2580ed5a8d7904e8fcb47fa6f996fabee04f869d
Este commit está contenido en:
Jenkins
2022-04-04 13:12:44 -04:00
Se han modificado 55 ficheros con 6582 adiciones y 259 borrados
+21 -5
Ver fichero
@@ -22,7 +22,6 @@
# Need perl > 5.10 to use logic-defined or
use 5.006; use v5.10.1;
use strict;
use warnings;
use File::Basename;
@@ -35,13 +34,30 @@ my $SCRIPT_DIR=dirname(__FILE__);
if ($HIPCC_USE_PERL_SCRIPT) {
#Invoke hipcc.pl
my $HIPCC_PERL=catfile($SCRIPT_DIR, '/hipcc.pl');
exec($^X, $HIPCC_PERL, @ARGV);
system("$^X $HIPCC_PERL @ARGV");
} else {
#Invoke hipcc.bin
my $HIPCC_BIN=catfile($SCRIPT_DIR, '/hipcc.bin');
$isWindows = ($^O eq 'MSWin32' or $^O eq 'msys');
$BIN_NAME="/hipcc.bin";
if ($isWindows) {
$BIN_NAME="/hipcc.bin.exe";
}
my $HIPCC_BIN=catfile($SCRIPT_DIR, $BIN_NAME);
if ( -e $HIPCC_BIN ) {
exec($HIPCC_BIN, @ARGV);
#Invoke hipcc.bin
my $output = qx($HIPCC_BIN @ARGV);
print ("$output\n");
} else {
print "hipcc.bin not present; Install HIPCC binaries before proceeding";
exit(-1);
}
}
if ($? == -1) {
exit($?);
}
elsif ($? & 127) {
exit($?);
}
else {
$CMD_EXIT_CODE = $? >> 8;
}
exit($CMD_EXIT_CODE);
+21 -5
Ver fichero
@@ -22,7 +22,6 @@
# Need perl > 5.10 to use logic-defined or
use 5.006; use v5.10.1;
use strict;
use warnings;
use File::Basename;
@@ -35,13 +34,30 @@ my $SCRIPT_DIR=dirname(__FILE__);
if ($HIPCONFIG_USE_PERL_SCRIPT) {
#Invoke hipconfig.pl
my $HIPCONFIG_PERL=catfile($SCRIPT_DIR, '/hipconfig.pl');
exec($^X, $HIPCONFIG_PERL, @ARGV);
system("$^X $HIPCONFIG_PERL @ARGV");
} else {
#Invoke hipconfig.bin
my $HIPCONFIG_BIN=catfile($SCRIPT_DIR, '/hipconfig.bin');
$isWindows = ($^O eq 'MSWin32' or $^O eq 'msys');
$BIN_NAME="/hipconfig.bin";
if ($isWindows) {
$BIN_NAME="/hipconfig.bin.exe";
}
my $HIPCONFIG_BIN=catfile($SCRIPT_DIR, $BIN_NAME);
if ( -e $HIPCONFIG_BIN ) {
exec($HIPCONFIG_BIN, @ARGV);
#Invoke hipconfig.bin
my $output = qx($HIPCONFIG_BIN @ARGV);
print ("$output\n");
} else {
print "hipconfig.bin not present; Install HIPCC binaries before proceeding";
exit(-1);
}
}
if ($? == -1) {
exit($?);
}
elsif ($? & 127) {
exit($?);
}
else {
$CMD_EXIT_CODE = $? >> 8;
}
exit($CMD_EXIT_CODE);
+10 -1
Ver fichero
@@ -4,6 +4,12 @@
HIP allows you to compile kernels at runtime with its ```hiprtc*``` APIs.
Kernels can be store as a text string and can be passed on to hiprtc APIs alongside options to guide the compilation.
NOTE:
- This library can be used on systems without HIP install nor AMD GPU driver installed at all (offline compilation). Therefore it does not depend on any HIP runtime library.
- But it does depend on COMGr. We may try to statically link COMGr into hipRTC to avoid any ambiguity.
- Developers can decide to bundle this library with their application.
## Example
To use hiprtc functionality, hiprtc header needs to be included first.
```#include <hip/hiprtc.h>```
@@ -84,6 +90,9 @@ Please have a look at saxpy.cpp and hiprtcGetLoweredName.cpp files for a detaile
## HIPRTC specific options
HIPRTC provides a few hiprtc specific flags
- ```--gpu-architecture``` : This flag can guide the code object generation for a specific gpu arch. Example: ```--gpu-architecture=gfx906:sramecc+:xnack-```, its equivalent to ```--offload-arch```.
- This option is compulsory if compilation is done on a system without AMD GPUs supported by HIP runtime.
- Otherwise, hipRTC will load the hip runtime and gather the current device and its architecture info and use it as option.
## Deprecation notice
Users will be required to link to libhiprtc.so/libhiprtc.dll in future releases. Currently all symbols are present in libhipamd64.so/libhipamd64.dll and there is a plan in action to separate HIPRTC APIs from HIP APIs.
Currently HIPRTC APIs are separated from HIP APIs and HIPRTC is available as a separate library libhiprtc.so/libhiprtc.dll. But hiprtc symbols are present in libhipamd64.so/libhipamd64.dll in order to support the existing applications. Gradually, these symbols will be removed from HIP library and applications using HIPRTC will be required to explictly link to HIPRTC library.
+2 -1
Ver fichero
@@ -98,7 +98,8 @@ if (HSA_HEADER-NOTFOUND)
message (FATAL_ERROR "HSA header not found! ROCM_PATH environment not set")
endif()
file(GLOB HIP_CLANGRT_LIB_SEARCH_PATHS "${CMAKE_HIP_COMPILER}/../lib/clang/*/lib/*")
get_filename_component(HIP_COMPILER_INSTALL_PATH ${CMAKE_HIP_COMPILER} DIRECTORY)
file(GLOB HIP_CLANGRT_LIB_SEARCH_PATHS "${HIP_COMPILER_INSTALL_PATH}/../lib/clang/*/lib/*")
find_library(CLANGRT_BUILTINS
NAMES
clang_rt.builtins
+374 -97
Ver fichero
@@ -1085,6 +1085,104 @@ typedef struct _hipGraphicsResource hipGraphicsResource;
typedef hipGraphicsResource* hipGraphicsResource_t;
/**
* An opaque value that represents a hip graph
*/
typedef struct ihipGraph* hipGraph_t;
/**
* An opaque value that represents a hip graph node
*/
typedef struct hipGraphNode* hipGraphNode_t;
/**
* An opaque value that represents a hip graph Exec
*/
typedef struct hipGraphExec* hipGraphExec_t;
/**
* @brief hipGraphNodeType
* @enum
*
*/
typedef enum hipGraphNodeType {
hipGraphNodeTypeKernel = 1, ///< GPU kernel node
hipGraphNodeTypeMemcpy = 2, ///< Memcpy 3D node
hipGraphNodeTypeMemset = 3, ///< Memset 1D node
hipGraphNodeTypeHost = 4, ///< Host (executable) node
hipGraphNodeTypeGraph = 5, ///< Node which executes an embedded graph
hipGraphNodeTypeEmpty = 6, ///< Empty (no-op) node
hipGraphNodeTypeWaitEvent = 7, ///< External event wait node
hipGraphNodeTypeEventRecord = 8, ///< External event record node
hipGraphNodeTypeMemcpy1D = 9, ///< Memcpy 1D node
hipGraphNodeTypeMemcpyFromSymbol = 10, ///< MemcpyFromSymbol node
hipGraphNodeTypeMemcpyToSymbol = 11, ///< MemcpyToSymbol node
hipGraphNodeTypeCount
} hipGraphNodeType;
typedef void (*hipHostFn_t)(void* userData);
typedef struct hipHostNodeParams {
hipHostFn_t fn;
void* userData;
} hipHostNodeParams;
typedef struct hipKernelNodeParams {
dim3 blockDim;
void** extra;
void* func;
dim3 gridDim;
void** kernelParams;
unsigned int sharedMemBytes;
} hipKernelNodeParams;
typedef struct hipMemsetParams {
void* dst;
unsigned int elementSize;
size_t height;
size_t pitch;
unsigned int value;
size_t width;
} hipMemsetParams;
/**
* @brief hipGraphExecUpdateResult
* @enum
*
*/
typedef enum hipGraphExecUpdateResult {
hipGraphExecUpdateSuccess = 0x0, ///< The update succeeded
hipGraphExecUpdateError = 0x1, ///< The update failed for an unexpected reason which is described
///< in the return value of the function
hipGraphExecUpdateErrorTopologyChanged = 0x2, ///< The update failed because the topology changed
hipGraphExecUpdateErrorNodeTypeChanged = 0x3, ///< The update failed because a node type changed
hipGraphExecUpdateErrorFunctionChanged =
0x4, ///< The update failed because the function of a kernel node changed
hipGraphExecUpdateErrorParametersChanged =
0x5, ///< The update failed because the parameters changed in a way that is not supported
hipGraphExecUpdateErrorNotSupported =
0x6, ///< The update failed because something about the node is not supported
hipGraphExecUpdateErrorUnsupportedFunctionChange = 0x7
} hipGraphExecUpdateResult;
typedef enum hipStreamCaptureMode {
hipStreamCaptureModeGlobal = 0,
hipStreamCaptureModeThreadLocal,
hipStreamCaptureModeRelaxed
} hipStreamCaptureMode;
typedef enum hipStreamCaptureStatus {
hipStreamCaptureStatusNone = 0, ///< Stream is not capturing
hipStreamCaptureStatusActive, ///< Stream is actively capturing
hipStreamCaptureStatusInvalidated ///< Stream is part of a capture sequence that has been
///< invalidated, but not terminated
} hipStreamCaptureStatus;
typedef enum hipStreamUpdateCaptureDependenciesFlags {
hipStreamAddCaptureDependencies = 0, ///< Add new nodes to the dependency set
hipStreamSetCaptureDependencies, ///< Replace the dependency set with the new nodes
} hipStreamUpdateCaptureDependenciesFlags;
typedef enum hipGraphInstantiateFlags {
hipGraphInstantiateFlagAutoFreeOnLaunch =
1, ///< Automatically free memory allocated in a graph before relaunching.
} hipGraphInstantiateFlags;
#include <hip/amd_detail/amd_hip_runtime_pt_api.h>
// Doxygen end group GlobalDefs
/** @} */
//-------------------------------------------------------------------------------------------------
@@ -5112,103 +5210,6 @@ int hipGetStreamDeviceId(hipStream_t stream);
* This section describes the graph management types & functions of HIP runtime API.
*/
/**
* An opaque value that represents a hip graph
*/
typedef struct ihipGraph* hipGraph_t;
/**
* An opaque value that represents a hip graph node
*/
typedef struct hipGraphNode* hipGraphNode_t;
/**
* An opaque value that represents a hip graph Exec
*/
typedef struct hipGraphExec* hipGraphExec_t;
/**
* @brief hipGraphNodeType
* @enum
*
*/
typedef enum hipGraphNodeType {
hipGraphNodeTypeKernel = 1, ///< GPU kernel node
hipGraphNodeTypeMemcpy = 2, ///< Memcpy 3D node
hipGraphNodeTypeMemset = 3, ///< Memset 1D node
hipGraphNodeTypeHost = 4, ///< Host (executable) node
hipGraphNodeTypeGraph = 5, ///< Node which executes an embedded graph
hipGraphNodeTypeEmpty = 6, ///< Empty (no-op) node
hipGraphNodeTypeWaitEvent = 7, ///< External event wait node
hipGraphNodeTypeEventRecord = 8, ///< External event record node
hipGraphNodeTypeMemcpy1D = 9, ///< Memcpy 1D node
hipGraphNodeTypeMemcpyFromSymbol = 10, ///< MemcpyFromSymbol node
hipGraphNodeTypeMemcpyToSymbol = 11, ///< MemcpyToSymbol node
hipGraphNodeTypeCount
} hipGraphNodeType;
typedef void (*hipHostFn_t)(void* userData);
typedef struct hipHostNodeParams {
hipHostFn_t fn;
void* userData;
} hipHostNodeParams;
typedef struct hipKernelNodeParams {
dim3 blockDim;
void** extra;
void* func;
dim3 gridDim;
void** kernelParams;
unsigned int sharedMemBytes;
} hipKernelNodeParams;
typedef struct hipMemsetParams {
void* dst;
unsigned int elementSize;
size_t height;
size_t pitch;
unsigned int value;
size_t width;
} hipMemsetParams;
/**
* @brief hipGraphExecUpdateResult
* @enum
*
*/
typedef enum hipGraphExecUpdateResult {
hipGraphExecUpdateSuccess = 0x0, ///< The update succeeded
hipGraphExecUpdateError = 0x1, ///< The update failed for an unexpected reason which is described
///< in the return value of the function
hipGraphExecUpdateErrorTopologyChanged = 0x2, ///< The update failed because the topology changed
hipGraphExecUpdateErrorNodeTypeChanged = 0x3, ///< The update failed because a node type changed
hipGraphExecUpdateErrorFunctionChanged =
0x4, ///< The update failed because the function of a kernel node changed
hipGraphExecUpdateErrorParametersChanged =
0x5, ///< The update failed because the parameters changed in a way that is not supported
hipGraphExecUpdateErrorNotSupported =
0x6, ///< The update failed because something about the node is not supported
hipGraphExecUpdateErrorUnsupportedFunctionChange = 0x7
} hipGraphExecUpdateResult;
typedef enum hipStreamCaptureMode {
hipStreamCaptureModeGlobal = 0,
hipStreamCaptureModeThreadLocal,
hipStreamCaptureModeRelaxed
} hipStreamCaptureMode;
typedef enum hipStreamCaptureStatus {
hipStreamCaptureStatusNone = 0, ///< Stream is not capturing
hipStreamCaptureStatusActive, ///< Stream is actively capturing
hipStreamCaptureStatusInvalidated ///< Stream is part of a capture sequence that has been
///< invalidated, but not terminated
} hipStreamCaptureStatus;
typedef enum hipStreamUpdateCaptureDependenciesFlags {
hipStreamAddCaptureDependencies = 0, ///< Add new nodes to the dependency set
hipStreamSetCaptureDependencies, ///< Replace the dependency set with the new nodes
} hipStreamUpdateCaptureDependenciesFlags;
typedef enum hipGraphInstantiateFlags {
hipGraphInstantiateFlagAutoFreeOnLaunch =
1, ///< Automatically free memory allocated in a graph before relaunching.
} hipGraphInstantiateFlags;
/**
* @brief Begins graph capture on a stream.
*
@@ -6172,6 +6173,282 @@ hipError_t hipGraphExecEventWaitNodeSetEvent(hipGraphExec_t hGraphExec, hipGraph
*/
/**
* Memory allocation properties
*/
typedef struct hipMemAllocationProp {
unsigned char compressionType; ///< Compression type
hipMemLocation location; ///< Memory location
hipMemAllocationHandleType requestedHandleType; ///< Requested handle type
hipMemAllocationType type; ///< Memory allocation type
unsigned short usage; ///< Usage
void* win32HandleMetaData; ///< Metadata for Win32 handles
} hipMemAllocationProp;
/**
* Generic handle for memory allocation
*/
typedef struct ihipMemGenericAllocationHandle* hipMemGenericAllocationHandle_t;
/**
* @brief Flags for granularity
* @enum
* @ingroup Enumerations
*/
typedef enum hipMemAllocationGranularity_flags {
hipMemAllocationGranularityMinimum = 0x0, ///< Minimum granularity
hipMemAllocationGranularityRecommended = 0x1 ///< Recommended granularity for performance
} hipMemAllocationGranularity_flags;
/**
* @brief Memory handle type
* @enum
* @ingroup Enumerations
*/
typedef enum hipMemHandleType {
hipMemHandleTypeGeneric = 0x0 ///< Generic handle type
} hipMemHandleType;
/**
* @brief Memory operation types
* @enum
* @ingroup Enumerations
*/
typedef enum hipMemOperationType {
hipMemOperationTypeMap = 0x1, ///< Map operation
hipMemOperationTypeUnmap = 0x2 ///< Unmap operation
} hipMemOperationType;
/**
* @brief Subresource types for sparse arrays
* @enum
* @ingroup Enumerations
*/
typedef enum hipArraySparseSubresourceType {
hipArraySparseSubresourceTypeSparseLevel = 0x0, ///< Sparse level
hipArraySparseSubresourceTypeMiptail = 0x1 ///< Miptail
} hipArraySparseSubresourceType;
/**
* Map info for arrays
*/
typedef struct hipArrayMapInfo {
hipResourceType resourceType; ///< Resource type
union {
hipMipmappedArray mipmap;
hipArray_t array;
} resource;
hipArraySparseSubresourceType subresourceType; ///< Sparse subresource type
union {
struct {
unsigned int level; ///< For mipmapped arrays must be a valid mipmap level. For arrays must be zero
unsigned int layer; ///< For layered arrays must be a valid layer index. Otherwise, must be zero
unsigned int offsetX; ///< X offset in elements
unsigned int offsetY; ///< Y offset in elements
unsigned int offsetZ; ///< Z offset in elements
unsigned int extentWidth; ///< Width in elements
unsigned int extentHeight; ///< Height in elements
unsigned int extentDepth; ///< Depth in elements
} sparseLevel;
struct {
unsigned int layer; ///< For layered arrays must be a valid layer index. Otherwise, must be zero
unsigned long long offset; ///< Offset within mip tail
unsigned long long size; ///< Extent in bytes
} miptail;
} subresource;
hipMemOperationType memOperationType; ///< Memory operation type
hipMemHandleType memHandleType; ///< Memory handle type
union {
hipMemGenericAllocationHandle_t memHandle;
} memHandle;
unsigned long long offset; ///< Offset within the memory
unsigned int deviceBitMask; ///< Device ordinal bit mask
unsigned int flags; ///< flags for future use, must be zero now.
unsigned int reserved[2]; ///< Reserved for future use, must be zero now.
} hipArrayMapInfo;
/**
*-------------------------------------------------------------------------------------------------
*-------------------------------------------------------------------------------------------------
* @defgroup Virtual Memory Management
* @{
* This section describes the virtual memory management functions of HIP runtime API.
*/
/**
* @brief Frees an address range reservation made via hipMemAddressReserve
*
* @param [in] devPtr - starting address of the range.
* @param [in] size - size of the range.
* @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorNotSupported
* @warning : This API is marked as beta, meaning, while this is feature complete,
* it is still open to changes and may have outstanding issues.
*/
hipError_t hipMemAddressFree(void* devPtr, size_t size);
/**
* @brief Reserves an address range
*
* @param [out] ptr - starting address of the reserved range.
* @param [in] size - size of the reservation.
* @param [in] alignment - alignment of the address.
* @param [in] addr - requested starting address of the range.
* @param [in] flags - currently unused, must be zero.
* @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorNotSupported
* @warning : This API is marked as beta, meaning, while this is feature complete,
* it is still open to changes and may have outstanding issues.
*/
hipError_t hipMemAddressReserve(void** ptr, size_t size, size_t alignment, void* addr, unsigned long long flags);
/**
* @brief Creates a memory allocation described by the properties and size
*
* @param [out] handle - value of the returned handle.
* @param [in] size - size of the allocation.
* @param [in] prop - properties of the allocation.
* @param [in] flags - currently unused, must be zero.
* @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorNotSupported
* @warning : This API is marked as beta, meaning, while this is feature complete,
* it is still open to changes and may have outstanding issues.
*/
hipError_t hipMemCreate(hipMemGenericAllocationHandle_t* handle, size_t size, const hipMemAllocationProp* prop, unsigned long long flags);
/**
* @brief Exports an allocation to a requested shareable handle type.
*
* @param [out] shareableHandle - value of the returned handle.
* @param [in] handle - handle to share.
* @param [in] handleType - type of the shareable handle.
* @param [in] flags - currently unused, must be zero.
* @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorNotSupported
* @warning : This API is marked as beta, meaning, while this is feature complete,
* it is still open to changes and may have outstanding issues.
*/
hipError_t hipMemExportToShareableHandle(void* shareableHandle, hipMemGenericAllocationHandle_t handle, hipMemAllocationHandleType handleType, unsigned long long flags);
/**
* @brief Get the access flags set for the given location and ptr.
*
* @param [out] flags - flags for this location.
* @param [in] location - target location.
* @param [in] ptr - address to check the access flags.
* @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorNotSupported
* @warning : This API is marked as beta, meaning, while this is feature complete,
* it is still open to changes and may have outstanding issues.
*/
hipError_t hipMemGetAccess(unsigned long long* flags, const hipMemLocation* location, void* ptr);
/**
* @brief Calculates either the minimal or recommended granularity.
*
* @param [out] granularity - returned granularity.
* @param [in] prop - location properties.
* @param [in] option - determines which granularity to return.
* @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorNotSupported
* @warning : This API is marked as beta, meaning, while this is feature complete,
* it is still open to changes and may have outstanding issues.
*/
hipError_t hipMemGetAllocationGranularity(size_t* granularity, const hipMemAllocationProp* prop, hipMemAllocationGranularity_flags option);
/**
* @brief Retrieve the property structure of the given handle.
*
* @param [out] prop - properties of the given handle.
* @param [in] handle - handle to perform the query on.
* @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorNotSupported
* @warning : This API is marked as beta, meaning, while this is feature complete,
* it is still open to changes and may have outstanding issues.
*/
hipError_t hipMemGetAllocationPropertiesFromHandle(hipMemAllocationProp* prop, hipMemGenericAllocationHandle_t handle);
/**
* @brief Imports an allocation from a requested shareable handle type.
*
* @param [out] handle - returned value.
* @param [in] osHandle - shareable handle representing the memory allocation.
* @param [in] shHandleType - handle type.
* @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorNotSupported
* @warning : This API is marked as beta, meaning, while this is feature complete,
* it is still open to changes and may have outstanding issues.
*/
hipError_t hipMemImportFromShareableHandle(hipMemGenericAllocationHandle_t* handle, void* osHandle, hipMemAllocationHandleType shHandleType);
/**
* @brief Maps an allocation handle to a reserved virtual address range.
*
* @param [in] ptr - address where the memory will be mapped.
* @param [in] size - size of the mapping.
* @param [in] offset - offset into the memory, currently must be zero.
* @param [in] handle - memory allocation to be mapped.
* @param [in] flags - currently unused, must be zero.
* @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorNotSupported
* @warning : This API is marked as beta, meaning, while this is feature complete,
* it is still open to changes and may have outstanding issues.
*/
hipError_t hipMemMap(void* ptr, size_t size, size_t offset, hipMemGenericAllocationHandle_t handle, unsigned long long flags);
/**
* @brief Maps or unmaps subregions of sparse HIP arrays and sparse HIP mipmapped arrays.
*
* @param [in] mapInfoList - list of hipArrayMapInfo.
* @param [in] count - number of hipArrayMapInfo in mapInfoList.
* @param [in] stream - stream identifier for the stream to use for map or unmap operations.
* @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorNotSupported
* @warning : This API is marked as beta, meaning, while this is feature complete,
* it is still open to changes and may have outstanding issues.
*/
hipError_t hipMemMapArrayAsync(hipArrayMapInfo* mapInfoList, unsigned int count, hipStream_t stream);
/**
* @brief Release a memory handle representing a memory allocation which was previously allocated through hipMemCreate.
*
* @param [in] handle - handle of the memory allocation.
* @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorNotSupported
* @warning : This API is marked as beta, meaning, while this is feature complete,
* it is still open to changes and may have outstanding issues.
*/
hipError_t hipMemRelease(hipMemGenericAllocationHandle_t handle);
/**
* @brief Returns the allocation handle of the backing memory allocation given the address.
*
* @param [out] handle - handle representing addr.
* @param [in] addr - address to look up.
* @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorNotSupported
* @warning : This API is marked as beta, meaning, while this is feature complete,
* it is still open to changes and may have outstanding issues.
*/
hipError_t hipMemRetainAllocationHandle(hipMemGenericAllocationHandle_t* handle, void* addr);
/**
* @brief Set the access flags for each location specified in desc for the given virtual address range.
*
* @param [in] ptr - starting address of the virtual address range.
* @param [in] size - size of the range.
* @param [in] desc - array of hipMemAccessDesc.
* @param [in] count - number of hipMemAccessDesc in desc.
* @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorNotSupported
* @warning : This API is marked as beta, meaning, while this is feature complete,
* it is still open to changes and may have outstanding issues.
*/
hipError_t hipMemSetAccess(void* ptr, size_t size, const hipMemAccessDesc* desc, size_t count);
/**
* @brief Unmap memory allocation of a given address range.
*
* @param [in] ptr - starting address of the range to unmap.
* @param [in] size - size of the virtual address range.
* @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorNotSupported
* @warning : This API is marked as beta, meaning, while this is feature complete,
* it is still open to changes and may have outstanding issues.
*/
hipError_t hipMemUnmap(void* ptr, size_t size);
// doxygen end virtual memory management API
/**
* @}
*/
/**
*-------------------------------------------------------------------------------------------------
*-------------------------------------------------------------------------------------------------
+56 -2
Ver fichero
@@ -1,5 +1,5 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2021 - 2022 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
@@ -23,6 +23,10 @@ THE SOFTWARE.
#pragma once
#include "hip_test_common.hh"
#include <iostream>
#include<fstream>
#include<regex>
#include <type_traits>
#define guarantee(cond, str) \
{ \
if (!(cond)) { \
@@ -235,5 +239,55 @@ unsigned setNumBlocks(T blocksPerCU, T threadsPerBlock,
}
return blocks;
}
template<typename T>
static bool assemblyFile_Verification(std::string assemfilename, std::string inst) {
std::string filePath = "./catch/unit/deviceLib/";
bool result = false;
std::string filename;
filename = filePath + assemfilename;
std::ifstream file(filename.c_str(), std::ios::out);
if (file) {
std::string line;
int line_pos = 0, start_pos = 0;
int last_pos = 0;
int start_match = 0;
while (getline(file, line)) {
line_pos++;
if ((std::is_same<T, float>::value)) {
if (!start_pos &&
std::regex_search(line,
std::regex("Begin function (.*)AtomicCheck"))) {
start_pos = line_pos;
}
if (!last_pos &&
std::regex_search(line,
std::regex(".Lfunc_end0-(.*)AtomicCheck"))) {
last_pos = line_pos;
break;
}
} else {
if ((start_match != 2) && std::regex_search(line,
std::regex("Begin function (.*)AtomicCheck"))) {
start_match++;
if (start_match == 2)
start_pos = line_pos;
}
if (!last_pos && std::regex_search(line,
std::regex("func_end1-(.*)AtomicCheck"))) {
last_pos = line_pos;
break;
}
}
if (start_pos) {
result = std::regex_search(line, std::regex(inst));
if (result)
break;
}
}
} else {
result = true;
SUCCEED("Assembly file does not exist");
}
return result;
}
} // namespace HipTest
+2 -1
Ver fichero
@@ -23,9 +23,9 @@ if (UNIX)
add_subdirectory(memory)
add_subdirectory(graph)
add_subdirectory(rtc)
add_subdirectory(deviceLib)
endif()
add_subdirectory(deviceLib)
add_subdirectory(stream)
add_subdirectory(event)
add_subdirectory(occupancy)
@@ -34,3 +34,4 @@ add_subdirectory(printf)
add_subdirectory(printfExe)
add_subdirectory(texture)
add_subdirectory(streamperthread)
add_subdirectory(kernel)
@@ -0,0 +1,91 @@
/*
Copyright (c) 2022 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.
*/
/*
AtomicAdd on FineGrainMemory
1. The following test scenario verifies
atomicAdd on fineGrain memory with -mno-unsafe-atomics flag
This testcase works only on gfx90a.
*/
#include<hip_test_checkers.hh>
#include<hip_test_common.hh>
#define INC_VAL 10
#define INITIAL_VAL 5
template<typename T>
static __global__ void AtomicCheck(T* Ad, T* result) {
T inc_val = 10;
*result = atomicAdd(Ad, inc_val);
}
/*atomicAdd API for the fine grained memory variable
with -mno-unsafe-atomics flag
Input: Ad{5}, INC_VAL{10}
Output: atomicAdd API would work and the 0/P is INITIAL_VAL + INC_VAL
Generate the assembly file and check whether
global_atomic_cmpswap instruction is generated
or not */
TEMPLATE_TEST_CASE("Unit_AtomicAdd_Coherentwithnounsafeflag", "",
float, double) {
hipDeviceProp_t prop;
int device;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
TestType *A_h{nullptr}, *result{nullptr};
TestType *A_d{nullptr}, *result_d{nullptr};
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(TestType),
hipHostMallocCoherent));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&result),
sizeof(TestType),
hipHostMallocCoherent));
result[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&result_d),
result, 0));
hipLaunchKernelGGL(AtomicCheck<TestType>, dim3(1), dim3(1),
0, 0, A_d,
result_d);
HIP_CHECK(hipDeviceSynchronize());
bool testResult;
testResult = HipTest::assemblyFile_Verification<TestType>(
"AtomicAdd_Coherent_withnoUnsafeflag-hip-amdgcn(.*)\\.s",
"global_atomic_cmpswap");
REQUIRE(testResult == true);
REQUIRE(A_h[0] == INITIAL_VAL + INC_VAL);
REQUIRE(result[0] == INITIAL_VAL);
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipHostFree(result));
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
@@ -0,0 +1,90 @@
/*
Copyright (c) 2022 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.
*/
/*
AtomicAdd on FineGrainMemory
1. The following test scenario verifies
atomicAdd on fineGrain memory without any unsafeatomics flag
This testcase works only on gfx90a.
*/
#include<hip_test_checkers.hh>
#include<hip_test_common.hh>
#define INC_VAL 10
#define INITIAL_VAL 5
template<typename T>
static __global__ void AtomicCheck(T* Ad, T* result) {
T inc_val = 10;
*result = atomicAdd(Ad, inc_val);
}
/*atomicAdd API for the fine grained memory variable
without any flag
Input: Ad{5}, INC_VAL{10}
Output: atomicAdd API would work and the 0/P is INITIAL_VAL + INC_VAL
Generate the assembly file and check whether
global_atomic_cmpswap instruction is generated
or not */
TEMPLATE_TEST_CASE("Unit_AtomicAdd_Coherentwithoutflag", "",
float, double) {
hipDeviceProp_t prop;
int device;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
TestType *A_h{nullptr}, *result{nullptr};
TestType *A_d{nullptr}, *result_d{nullptr};
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(TestType),
hipHostMallocCoherent));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&result),
sizeof(TestType),
hipHostMallocCoherent));
result[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&result_d),
result, 0));
hipLaunchKernelGGL(AtomicCheck<TestType>, dim3(1), dim3(1),
0, 0, A_d,
result_d);
HIP_CHECK(hipDeviceSynchronize());
bool testResult;
testResult = HipTest::assemblyFile_Verification<TestType>(
"AtomicAdd_Coherent_withoutflag-hip-amdgcn(.*)\\.s",
"global_atomic_cmpswap");
REQUIRE(result[0] == INITIAL_VAL);
REQUIRE(A_h[0] == INITIAL_VAL + INC_VAL);
REQUIRE(testResult == true);
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipHostFree(result));
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
@@ -0,0 +1,99 @@
/*
Copyright (c) 2022 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.
*/
/*
AtomicAdd on FineGrainMemory
1. The following test scenario verifies
atomicAdd on fineGrain memory with -munsafe-fp-atomics flag
This testcase works only on gfx90a.
*/
#include<hip_test_checkers.hh>
#include<hip_test_common.hh>
#define INC_VAL 10
#define INITIAL_VAL 5
template<typename T>
static __global__ void AtomicCheck(T* Ad, T* result) {
T inc_val = 10;
*result = atomicAdd(Ad, inc_val);
}
/*atomicAdd API for the fine grained memory variable
with -m-unsafe-atomics flag
Input: Ad{5}, INC_VAL{10}
Output: atomicAdd API would return 0 and the 0/P is 5
Generate the assembly file and check whether
global_atomic_cmpswap instruction is generated
or not */
TEMPLATE_TEST_CASE("Unit_AtomicAdd_CoherentwithUnsafeflag", "",
float, double) {
hipDeviceProp_t prop;
int device;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
TestType *A_h{nullptr}, *result{nullptr};
TestType *A_d{nullptr}, *result_d{nullptr};
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(TestType),
hipHostMallocCoherent));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&result),
sizeof(TestType),
hipHostMallocCoherent));
result[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&result_d),
result, 0));
hipLaunchKernelGGL(AtomicCheck<TestType>, dim3(1), dim3(1),
0, 0, A_d,
result_d);
HIP_CHECK(hipDeviceSynchronize());
bool testResult;
if ((std::is_same<TestType, float>::value)) {
testResult = HipTest::assemblyFile_Verification<TestType>(
"AtomicAdd_Coherent_withunsafeflag-hip-amdgcn(.*)\\.s",
"global_atomic_add_f32");
REQUIRE(testResult == true);
} else {
testResult = HipTest::assemblyFile_Verification<TestType>(
"AtomicAdd_Coherent_withunsafeflag-hip-amdgcn(.*)\\.s",
"global_atomic_add_f64");
REQUIRE(testResult == true);
}
REQUIRE(A_h[0] == INITIAL_VAL);
REQUIRE(result[0] == 0);
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipHostFree(result));
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
@@ -0,0 +1,91 @@
/*
Copyright (c) 2022 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.
*/
/*
AtomicAdd on CoarseGrainMemory
1. The following test scenario verifies
atomicAdd on CoarseGrain memory with -mno-unsafe-atomics flag
This testcase works only on gfx90a.
*/
#include<hip_test_checkers.hh>
#include<hip_test_common.hh>
#define INC_VAL 10
#define INITIAL_VAL 5
template<typename T>
static __global__ void AtomicCheck(T* Ad, T* result) {
T inc_val = 10;
*result = atomicAdd(Ad, inc_val);
}
/*atomicAdd API for the coarse grained memory variable
with -mno-unsafe-atomics flag
Input: Ad{5}, INC_VAL{10}
Output: atomicAdd API would work and the 0/P is INITIAL_VAL + INC_VAL
Generate the assembly file and check whether
global_atomic_cmpswap instruction is generated
or not */
TEMPLATE_TEST_CASE("Unit_AtomicAdd_NonCoherentwithnounsafeflag", "",
float, double) {
hipDeviceProp_t prop;
int device;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
TestType *A_h{nullptr}, *result{nullptr};
TestType *A_d{nullptr}, *result_d{nullptr};
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(TestType),
hipHostMallocNonCoherent));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&result),
sizeof(TestType),
hipHostMallocNonCoherent));
result[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&result_d),
result, 0));
hipLaunchKernelGGL(AtomicCheck<TestType>,
dim3(1), dim3(1),
0, 0, A_d,
result_d);
HIP_CHECK(hipDeviceSynchronize());
bool testResult;
REQUIRE(A_h[0] == INITIAL_VAL + INC_VAL);
REQUIRE(result[0] == INITIAL_VAL);
testResult = HipTest::assemblyFile_Verification<TestType>(
"AtomicAdd_NonCoherent_withnounsafeflag-hip-amdgcn(.*)\\.s",
"global_atomic_cmpswap");
REQUIRE(testResult == true);
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipHostFree(result));
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
@@ -0,0 +1,91 @@
/*
Copyright (c) 2022 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.
*/
/*
AtomicAdd on CoarseGrainMemory
1. The following test scenario verifies
atomicAdd on CoarseGrain memory without any unsafeatomics flag
This testcase works only on gfx90a.
*/
#include<hip_test_checkers.hh>
#include<hip_test_common.hh>
#define INC_VAL 10
#define INITIAL_VAL 5
template<typename T>
static __global__ void AtomicCheck(T* Ad, T* result) {
T inc_val = 10;
*result = atomicAdd(Ad, inc_val);
}
/*atomicAdd API for the coarse grained memory variable
without any flag
Input: Ad{5}, INC_VAL{10}
Output: atomicAdd API would work and the 0/P is INITIAL_VAL + INC_VAL
Generate the assembly file and check whether
global_atomic_cmpswap instruction is generated
or not */
TEMPLATE_TEST_CASE("Unit_AtomicAdd_NonCoherentwithoutflag", "",
float, double) {
hipDeviceProp_t prop;
int device;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
TestType *A_h{nullptr}, *result{nullptr};
TestType *A_d{nullptr}, *result_d{nullptr};
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(TestType),
hipHostMallocNonCoherent));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&result),
sizeof(TestType),
hipHostMallocNonCoherent));
result[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&result_d),
result, 0));
hipLaunchKernelGGL(AtomicCheck<TestType>,
dim3(1), dim3(1),
0, 0, A_d,
result_d);
HIP_CHECK(hipDeviceSynchronize());
bool testResult;
REQUIRE(A_h[0] == INITIAL_VAL + INC_VAL);
REQUIRE(result[0] == INITIAL_VAL);
testResult = HipTest::assemblyFile_Verification<TestType>(
"AtomicAdd_NonCoherent_withoutflag-hip-amdgcn(.*)\\.s",
"global_atomic_cmpswap");
REQUIRE(testResult == true);
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipHostFree(result));
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
@@ -0,0 +1,98 @@
/*
Copyright (c) 2022 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.
*/
/*
AtomicAdd on CoarseGrainMemory
1. The following test scenario verifies
atomicAdd on CoarseGrain memory with -munsafe-fp-atomics flag
This testcase works only on gfx90a.
*/
#include<hip_test_checkers.hh>
#include<hip_test_common.hh>
#define INC_VAL 10
#define INITIAL_VAL 5
template<typename T>
static __global__ void AtomicCheck(T* Ad, T* result) {
T inc_val = 10;
*result = atomicAdd(Ad, inc_val);
}
/*atomicAdd API for the fine grained memory variable
with -m-unsafe-atomics flag
Input: Ad{5}, INC_VAL{10}
Output: atomicAdd API would work and the 0/P is 15
Generate the assembly file and check whether
global_atomic_add_float/double instruction is generated
or not */
TEMPLATE_TEST_CASE("Unit_AtomicAdd_NonCoherentwithUnsafeflag", "",
float, double) {
hipDeviceProp_t prop;
int device;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
TestType *A_h{nullptr}, *result{nullptr};
TestType *A_d{nullptr}, *result_d{nullptr};
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(TestType),
hipHostMallocNonCoherent));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&result),
sizeof(TestType),
hipHostMallocNonCoherent));
result[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&result_d),
result, 0));
hipLaunchKernelGGL(AtomicCheck<TestType>, dim3(1), dim3(1),
0, 0, A_d,
result_d);
HIP_CHECK(hipDeviceSynchronize());
bool testResult;
REQUIRE(A_h[0] == INITIAL_VAL + INC_VAL);
REQUIRE(result[0] == INITIAL_VAL);
if ((std::is_same<TestType, float>::value)) {
testResult = HipTest::assemblyFile_Verification<TestType>(
"AtomicAdd_NonCoherent_withunsafeflag-hip-amdgcn(.*)\\.s",
"global_atomic_add_f32");
REQUIRE(testResult == true);
} else {
testResult = HipTest::assemblyFile_Verification<TestType>(
"AtomicAdd_NonCoherent_withunsafeflag-hip-amdgcn(.*)\\.s",
"global_atomic_add_f64");
REQUIRE(testResult == true);
}
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipHostFree(result));
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
+292
Ver fichero
@@ -0,0 +1,292 @@
/*
Copyright (c) 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.
*/
/*
This testfile verifies __builtin_amdgcn_global_atomic_fadd_f64 API scenarios
1. AtomicAdd on Coherent Memory
2. AtomicAdd on Non-Coherent Memory
3. AtomicAdd on Coherent Memory with RTC
4. AtomicAdd on Non-Coherent Memory with RTC
*/
#include<hip_test_checkers.hh>
#include<hip_test_common.hh>
#include <hip/hiprtc.h>
#define INC_VAL 10
#define INITIAL_VAL 5
__global__ void AtomicAdd_GlobalMem(double* addr, double* result) {
double inc_val = 10;
*result = __builtin_amdgcn_global_atomic_fadd_f64(addr, inc_val);
}
static constexpr auto AtomicAddGlobalMem{
R"(
extern "C"
__global__ void AtomicAdd_GlobalMem(double* addr, double* result) {
double inc_val = 10;
*result = __builtin_amdgcn_global_atomic_fadd_f64(addr, inc_val);
}
)"};
/*
This test verifies the built in atomic add API on Coherent Memory
Input: A_h with INITIAL_VAL
Output: A_h will not get updated with Coherent Memory
A_h will be INITIAL_VAL
ret value would be 0, B_h would be 0
*/
TEST_CASE("Unit_BuiltInAtomicAdd_CoherentGlobalMem") {
hipDeviceProp_t prop;
int device;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does support HostPinned Memory");
} else {
double *A_h, *result_h, *result;
double *A_d;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(double),
hipHostMallocCoherent));
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&result_h),
sizeof(double), hipHostMallocCoherent));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&result),
result_h, 0));
std::cout << "test" << std::endl;
hipLaunchKernelGGL(AtomicAdd_GlobalMem, dim3(1), dim3(1),
0, 0, A_d,
result);
std::cout << "test 1" << std::endl;
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(A_h[0] == INITIAL_VAL);
REQUIRE(*result_h == 0);
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipFree(result));
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a Hence"
"skipping the testcase for this GPU " << device);
}
}
/*
This test verifies the built in atomic add API on Non-Coherent Memory
Input: A_h with INITIAL_VAL
Output: A_h will not get updated with Coherent Memory
A_h will be INITIAL_VAL+INC_VAL
B_h would be initial value of A_h, B_h would be INITIAL_VAL
*/
TEST_CASE("Unit_BuiltInAtomicAdd_NonCoherentGlobalMem") {
hipDeviceProp_t prop;
int device;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
double *A_h, *result, *B_h;
double *A_d;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(double),
hipHostMallocNonCoherent));
B_h = reinterpret_cast<double*>(malloc(sizeof(double)));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&result), sizeof(double)));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
hipLaunchKernelGGL(AtomicAdd_GlobalMem, dim3(1), dim3(1),
0, 0, static_cast<double* >(A_d),
static_cast<double* >(result));
HIP_CHECK(hipDeviceSynchronize());
HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost));
REQUIRE(A_h[0] == INITIAL_VAL + INC_VAL);
REQUIRE(*B_h == INITIAL_VAL);
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipFree(result));
free(B_h);
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a"
"Hence skipping the testcase for GPU-0");
}
}
/*
This test verifies the built in atomic add API on Coherent Memory with RTC
Input: A_h with INITIAL_VAL
Output: A_h will not get updated with Coherent Memory
A_h will be INITIAL_VAL
ret value would be 0, B_h would be 0
*/
TEST_CASE("Unit_BuiltInAtomicAdd_CoherentGlobalMemWithRtc") {
hipDeviceProp_t prop;
int device;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
hiprtcProgram prog;
hiprtcCreateProgram(&prog, // prog
AtomicAddGlobalMem, // buffer
"kernel.cu", // name
0, nullptr, nullptr);
std::string sarg = std::string("--gpu-architecture=") + prop.gcnArchName;
const char* options[] = {sarg.c_str()};
hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, options)};
size_t logSize;
HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize));
if (logSize) {
std::string log(logSize, '\0');
HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0]));
INFO(log);
}
REQUIRE(compileResult == HIPRTC_SUCCESS);
size_t codeSize;
HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize));
std::vector<char> code(codeSize);
HIPRTC_CHECK(hiprtcGetCode(prog, code.data()));
HIPRTC_CHECK(hiprtcDestroyProgram(&prog));
hipModule_t module;
hipFunction_t fmaxkernel;
HIP_CHECK(hipModuleLoadData(&module, code.data()));
HIP_CHECK(hipModuleGetFunction(&fmaxkernel, module,
"AtomicAdd_GlobalMem"));
double *A_h, *result, *B_h;
double *A_d;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(double),
hipHostMallocCoherent));
B_h = reinterpret_cast<double*>(malloc(sizeof(double)));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&result), sizeof(double)));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
struct {
double* p;
double* res;
} args_f{A_d, result};
auto size = sizeof(args_f);
void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f,
HIP_LAUNCH_PARAM_BUFFER_SIZE,
&size, HIP_LAUNCH_PARAM_END};
hipModuleLaunchKernel(fmaxkernel, 1, 1, 1, 1, 1, 1, 0,
nullptr, nullptr, config_d);
HIP_CHECK(hipDeviceSynchronize());
HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost));
REQUIRE(A_h[0] == INITIAL_VAL);
REQUIRE(*B_h == 0);
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipFree(result));
free(B_h);
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
/*
This test verifies the built in atomic add API on Non-Coherent Memory
Input: A_h with INITIAL_VAL
Output: A_h will not get updated with Coherent Memory
A_h will be INITIAL_VAL+INC_VAL
B_h would be initial value of A_h, B_h would be INITIAL_VAL
*/
TEST_CASE("Unit_BuiltInAtomicAdd_NonCoherentGlobalMemWithRtc") {
hipDeviceProp_t prop;
int device;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does support HostPinned Memory");
} else {
hiprtcProgram prog;
hiprtcCreateProgram(&prog, // prog
AtomicAddGlobalMem, // buffer
"kernel.cu", // name
0, nullptr, nullptr);
std::string sarg = std::string("--gpu-architecture=") + prop.gcnArchName;
const char* options[] = {sarg.c_str()};
hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, options)};
size_t logSize;
HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize));
if (logSize) {
std::string log(logSize, '\0');
HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0]));
WARN(log);
}
REQUIRE(compileResult == HIPRTC_SUCCESS);
size_t codeSize;
HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize));
std::vector<char> code(codeSize);
HIPRTC_CHECK(hiprtcGetCode(prog, code.data()));
HIPRTC_CHECK(hiprtcDestroyProgram(&prog));
hipModule_t module;
hipFunction_t fmaxkernel;
HIP_CHECK(hipModuleLoadData(&module, code.data()));
HIP_CHECK(hipModuleGetFunction(&fmaxkernel, module,
"AtomicAdd_GlobalMem"));
double *A_h, *result, *B_h;
double *A_d;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(double),
hipHostMallocNonCoherent));
B_h = reinterpret_cast<double*>(malloc(sizeof(double)));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&result), sizeof(double)));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
struct {
double* p;
double* res;
} args_f{A_d, result};
auto size = sizeof(args_f);
void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f,
HIP_LAUNCH_PARAM_BUFFER_SIZE,
&size, HIP_LAUNCH_PARAM_END};
hipModuleLaunchKernel(fmaxkernel, 1, 1, 1, 1, 1, 1, 0,
nullptr, nullptr, config_d);
HIP_CHECK(hipDeviceSynchronize());
HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost));
REQUIRE(A_h[0] == INITIAL_VAL + INC_VAL);
REQUIRE(*B_h == INITIAL_VAL);
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipFree(result));
free(B_h);
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
+354
Ver fichero
@@ -0,0 +1,354 @@
/*
Copyright (c) 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.
*/
/*
This testfile verifies Built fmax API scenarios
1. Builtin fmax on Coherent Memory with memory type as global
2. Builtin fmax on Non-Coherent Memory with memory type as global
3. Builtin fmax with memory type as flat
4. Builtin fmax on Coherent Memory with RTC and memory type as global
5. Builtin fmax on Non-Coherent Memory with RTC and memory type as global
6. Builtin fmax with RTC and memory type as flat
*/
#include<hip_test_checkers.hh>
#include<hip_test_common.hh>
#include <hip/hiprtc.h>
#define INITIAL_VAL 5
__global__ void unsafeAtomicMax_FlatMem(double* addr, double* result) {
__shared__ double int_val;
int_val = 5;
double comp = 10;
if (__builtin_amdgcn_is_shared(
(const __attribute__((address_space(0))) void*)(&int_val)))
*result = __builtin_amdgcn_flat_atomic_fmax_f64(&int_val, comp);
else
*result = __builtin_amdgcn_global_atomic_fmax_f64(&int_val, comp);
*addr = int_val;
}
__global__ void unsafeAtomicMax_GlobalMem(double* addr, double* result) {
double comp = 10;
if (__builtin_amdgcn_is_shared(
(const __attribute__((address_space(0))) void*)(addr)))
*result = __builtin_amdgcn_flat_atomic_fmax_f64(addr, comp);
else
*result = __builtin_amdgcn_global_atomic_fmax_f64(addr, comp);
}
static constexpr auto fmaxFlatMem {
R"(
extern "C"
__global__ void unsafeAtomicMax_FlatMem(double* addr, double* result) {
__shared__ double int_val;
int_val = 5;
double comp = 10;
if (__builtin_amdgcn_is_shared(
(const __attribute__((address_space(0))) void*)(&int_val)))
*result = __builtin_amdgcn_flat_atomic_fmax_f64(&int_val, comp);
else
*result = __builtin_amdgcn_global_atomic_fmax_f64(&int_val, comp);
*addr = int_val;
}
)"};
static constexpr auto fmaxGlobalMem {
R"(
extern "C"
__global__ void unsafeAtomicMax_GlobalMem(double* addr, double* result) {
double comp = 10;
if (__builtin_amdgcn_is_shared(
(const __attribute__((address_space(0))) void*)(addr)))
*result = __builtin_amdgcn_flat_atomic_fmax_f64(addr, comp);
else
*result = __builtin_amdgcn_global_atomic_fmax_f64(addr, comp);
}
)"};
/*
This testcase verifies the builtinAtomic fmax API on Coherent memory
with memory type as global
Input: A_h with INITIAL_VAL
Output: Return val would be 0 and the input value to API will not
get updated. A_h would be INITIAL_VAL, B_h is 0
*/
TEST_CASE("Unit_BuiltinAtomics_fmaxCoherentGlobalMem") {
hipDeviceProp_t prop;
int device;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
double *A_h, *B_h;
double *A_d;
double *result;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(double),
hipHostMallocCoherent));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
B_h = reinterpret_cast<double*>(malloc(sizeof(double)));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&result), sizeof(double)));
hipLaunchKernelGGL(unsafeAtomicMax_GlobalMem, dim3(1), dim3(1),
0, 0, static_cast<double* >(A_d), result);
HIP_CHECK(hipDeviceSynchronize());
HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost));
REQUIRE(*B_h == 0);
REQUIRE(A_h[0] == INITIAL_VAL);
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipFree(result));
free(B_h);
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
/*
This testcase verifies the builtinAtomic fmax API
1. Non Coherent memory with memory type as global
2. Memory type as flat
Input: A_h with INITIAL_VAL
Output: Return val would be initial val of A_h and the input value of
API would be updated with the max value
A_h would be 10, B_h would be INITIAL_VAL
*/
TEST_CASE("Unit_BuiltinAtomics_fmaxNonCoherentGlobalFlatMem") {
int mem_type = GENERATE(0, 1);
hipDeviceProp_t prop;
int device;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
double *A_h, *B_h;
double *A_d;
double *result;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(double),
hipHostMallocNonCoherent));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
B_h = reinterpret_cast<double*>(malloc(sizeof(double)));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&result), sizeof(double)));
if (mem_type) {
hipLaunchKernelGGL(unsafeAtomicMax_GlobalMem, dim3(1), dim3(1),
0, 0, static_cast<double* >(A_d), result);
} else {
hipLaunchKernelGGL(unsafeAtomicMax_FlatMem, dim3(1), dim3(1),
0, 0, static_cast<double* >(A_d), result);
}
HIP_CHECK(hipDeviceSynchronize());
HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost));
REQUIRE(*B_h == INITIAL_VAL);
REQUIRE(A_h[0] == 10);
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipFree(result));
free(B_h);
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
/*
This testcase verifies the builtinAtomic fmax API on Coherent memory
with RTC and memory type as global
Input: A_h with INITIAL_VAL
Output: Return val would be 0 and the input value to API will not
get updated. A_h would be INITIAL_VAL, B_h is 0
*/
TEST_CASE("Unit_BuiltinAtomicsRTC_fmaxCoherentGlobalMem") {
hipDeviceProp_t prop;
int device;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
hiprtcProgram prog;
hiprtcCreateProgram(&prog, // prog
fmaxGlobalMem, // buffer
"kernel.cu", // name
0, nullptr, nullptr);
std::string sarg = std::string("--gpu-architecture=") + prop.gcnArchName;
const char* options[] = {sarg.c_str()};
hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, options)};
size_t logSize;
HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize));
if (logSize) {
std::string log(logSize, '\0');
HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0]));
INFO(log);
}
REQUIRE(compileResult == HIPRTC_SUCCESS);
size_t codeSize;
HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize));
std::vector<char> code(codeSize);
HIPRTC_CHECK(hiprtcGetCode(prog, code.data()));
HIPRTC_CHECK(hiprtcDestroyProgram(&prog));
hipModule_t module;
hipFunction_t fmaxkernel;
HIP_CHECK(hipModuleLoadData(&module, code.data()));
HIP_CHECK(hipModuleGetFunction(&fmaxkernel, module,
"unsafeAtomicMax_GlobalMem"));
double *A_h, *B_h;
double *A_d;
double *result;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(double),
hipHostMallocCoherent));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
B_h = reinterpret_cast<double*>(malloc(sizeof(double)));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&result), sizeof(double)));
struct {
double* p;
double* res;
} args_f{A_d, result};
auto size = sizeof(args_f);
void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f,
HIP_LAUNCH_PARAM_BUFFER_SIZE,
&size, HIP_LAUNCH_PARAM_END};
hipModuleLaunchKernel(fmaxkernel, 1, 1, 1, 1, 1, 1, 0,
nullptr, nullptr, config_d);
HIP_CHECK(hipDeviceSynchronize());
HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost));
REQUIRE(*B_h == 0);
REQUIRE(A_h[0] == INITIAL_VAL);
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipFree(result));
free(B_h);
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
/*
This testcase verifies the builtinAtomic fmax API with RTC
1. Non Coherent memory with memory type as global
2. Memory type as flat
Input: A_h with INITIAL_VAL
Output: Return val would be initial val of A_h and the input value of
API would be updated with the max value
A_h would be 10, B_h would be INITIAL_VAL
*/
TEST_CASE("Unit_BuiltinAtomicsRTC_fmaxNonCoherentGlobalFlatMem") {
int mem_type = GENERATE(0, 1);
hipDeviceProp_t prop;
int device;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
hiprtcProgram prog;
if (mem_type) {
hiprtcCreateProgram(&prog, // prog
fmaxGlobalMem, // buffer
"kernel.cu", // name
0, nullptr, nullptr);
} else {
hiprtcCreateProgram(&prog, // prog
fmaxFlatMem, // buffer
"kernel.cu", // name
0, nullptr, nullptr);
}
std::string sarg = std::string("--gpu-architecture=") + prop.gcnArchName;
const char* options[] = {sarg.c_str()};
hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, options)};
size_t logSize;
HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize));
if (logSize) {
std::string log(logSize, '\0');
HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0]));
INFO(log);
}
REQUIRE(compileResult == HIPRTC_SUCCESS);
size_t codeSize;
HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize));
std::vector<char> code(codeSize);
HIPRTC_CHECK(hiprtcGetCode(prog, code.data()));
HIPRTC_CHECK(hiprtcDestroyProgram(&prog));
hipModule_t module;
hipFunction_t fmaxkernel;
HIP_CHECK(hipModuleLoadData(&module, code.data()));
if (mem_type) {
HIP_CHECK(hipModuleGetFunction(&fmaxkernel, module,
"unsafeAtomicMax_GlobalMem"));
} else {
HIP_CHECK(hipModuleGetFunction(&fmaxkernel, module,
"unsafeAtomicMax_FlatMem"));
}
double *A_h, *B_h;
double *A_d;
double *result;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(double),
hipHostMallocNonCoherent));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
B_h = reinterpret_cast<double*>(malloc(sizeof(double)));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&result), sizeof(double)));
struct {
double* p;
double* res;
} args_f{A_d, result};
auto size = sizeof(args_f);
void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f,
HIP_LAUNCH_PARAM_BUFFER_SIZE,
&size, HIP_LAUNCH_PARAM_END};
hipModuleLaunchKernel(fmaxkernel, 1, 1, 1, 1, 1, 1, 0,
nullptr, nullptr, config_d);
HIP_CHECK(hipDeviceSynchronize());
HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost));
REQUIRE(*B_h == INITIAL_VAL);
REQUIRE(A_h[0] == 10);
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipFree(result));
free(B_h);
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
+357
Ver fichero
@@ -0,0 +1,357 @@
/*
Copyright (c) 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.
*/
/*
This testfile verifies Built fmin API scenarios
1. Builtin fmin on Coherent Memory with memory type as global
2. Builtin fmin on Non-Coherent Memory with memory type as global
3. Builtin fmin with memory type as flat
4. Builtin fmin on Coherent Memory with RTC and memory type as global
5. Builtin fmin on Non-Coherent Memory with RTC and memory type as global
6. Builtin fmin with RTC and memory type as flat
*/
#include<hip_test_checkers.hh>
#include<hip_test_common.hh>
#include <hip/hiprtc.h>
#define INITIAL_VAL 5
static constexpr auto fminFlatMem{
R"(
extern "C"
__global__ void unsafeAtomicMin_FlatMem(double* addr, double* result) {
__shared__ double int_val;
int_val = 5;
double comp = 10;
if (__builtin_amdgcn_is_shared(
(const __attribute__((address_space(0))) void*)(&int_val)))
*result = __builtin_amdgcn_flat_atomic_fmin_f64(&int_val, comp);
else
*result = __builtin_amdgcn_global_atomic_fmin_f64(&int_val, comp);
*addr = int_val;
}
)"};
static constexpr auto fminGlobalMem{
R"(
extern "C"
__global__ void unsafeAtomicMin_GlobalMem(double* addr, double* result) {
double comp = 10;
if (__builtin_amdgcn_is_shared(
(const __attribute__((address_space(0))) void*)(addr)))
*result = __builtin_amdgcn_flat_atomic_fmin_f64(addr, comp);
else
*result = __builtin_amdgcn_global_atomic_fmin_f64(addr, comp);
}
)"};
__global__ void unsafeAtomicMin_FlatMem(double* addr, double* result) {
__shared__ double int_val;
int_val = 5;
double comp = 10;
if (__builtin_amdgcn_is_shared(
(const __attribute__((address_space(0))) void*)(&int_val)))
*result = __builtin_amdgcn_flat_atomic_fmin_f64(&int_val, comp);
else
*result = __builtin_amdgcn_global_atomic_fmin_f64(&int_val, comp);
*addr = int_val;
}
__global__ void unsafeAtomicMin_GlobalMem(double* addr, double* result) {
double comp = 10;
if (__builtin_amdgcn_is_shared(
(const __attribute__((address_space(0))) void*)(addr)))
*result = __builtin_amdgcn_flat_atomic_fmin_f64(addr, comp);
else
*result = __builtin_amdgcn_global_atomic_fmin_f64(addr, comp);
}
/*
This testcase verifies the builtinAtomic fmin API on Coherent memory
with memory type as global
Input: A_h with INITIAL_VAL
Output: Return val would be 0 and the input value to API will not
get updated. A_h would be INITIAL_VAL, B_h is 0
*/
TEST_CASE("Unit_BuiltinAtomics_fminCoherentGlobalMem") {
hipDeviceProp_t prop;
int device;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
double *A_h, *B_h;
double *A_d;
double *result;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(double),
hipHostMallocCoherent));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
B_h = reinterpret_cast<double*>(malloc(sizeof(double)));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&result), sizeof(double)));
hipLaunchKernelGGL(unsafeAtomicMin_GlobalMem, dim3(1), dim3(1),
0, 0, static_cast<double* >(A_d), result);
HIP_CHECK(hipDeviceSynchronize());
HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost));
REQUIRE(*B_h == 0);
REQUIRE(A_h[0] == INITIAL_VAL);
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipFree(result));
free(B_h);
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
/*
This testcase verifies the builtinAtomic fmin API
1. Non Coherent memory with memory type as global
2. Memory type as flat
Input: A_h with INITIAL_VAL
Output: Return val would be initial val of A_h and the input value of
API would be updated with the min value
A_h would be INITIAL_VAL, B_h would be INITIAL_VAL
*/
TEST_CASE("Unit_BuiltinAtomics_fminNonCoherentGlobalFlatMem") {
auto mem_type = GENERATE(0, 1);
hipDeviceProp_t prop;
int device;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
double *A_h, *B_h;
double *A_d;
double *result;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(double),
hipHostMallocNonCoherent));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
B_h = reinterpret_cast<double*>(malloc(sizeof(double)));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&result), sizeof(double)));
if (mem_type) {
hipLaunchKernelGGL(unsafeAtomicMin_GlobalMem, dim3(1), dim3(1),
0, 0, static_cast<double* >(A_d), result);
} else {
hipLaunchKernelGGL(unsafeAtomicMin_FlatMem, dim3(1), dim3(1),
0, 0, static_cast<double* >(A_d), result);
}
HIP_CHECK(hipDeviceSynchronize());
HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost));
REQUIRE(*B_h == INITIAL_VAL);
REQUIRE(A_h[0] == INITIAL_VAL);
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipFree(result));
free(B_h);
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
/*
This testcase verifies the builtinAtomic fmin API on Coherent memory
with RTC and memory type as global
Input: A_h with INITIAL_VAL
Output: Return val would be 0 and the input value to API will not
get updated. A_h would be INITIAL_VAL, B_h is 0
*/
TEST_CASE("Unit_BuiltinAtomicsRTC__fminCoherentGlobalMem") {
hipDeviceProp_t prop;
int device;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
hiprtcProgram prog;
hiprtcCreateProgram(&prog, // prog
fminGlobalMem, // buffer
"kernel.cu", // name
0, nullptr, nullptr);
std::string sarg = std::string("--gpu-architecture=") + prop.gcnArchName;
const char* options[] = {sarg.c_str()};
hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, options)};
size_t logSize;
HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize));
if (logSize) {
std::string log(logSize, '\0');
HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0]));
INFO(log);
}
REQUIRE(compileResult == HIPRTC_SUCCESS);
size_t codeSize;
HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize));
std::vector<char> code(codeSize);
HIPRTC_CHECK(hiprtcGetCode(prog, code.data()));
HIPRTC_CHECK(hiprtcDestroyProgram(&prog));
hipModule_t module;
hipFunction_t fmaxkernel;
HIP_CHECK(hipModuleLoadData(&module, code.data()));
HIP_CHECK(hipModuleGetFunction(&fmaxkernel, module,
"unsafeAtomicMin_GlobalMem"));
double *A_h, *B_h;
double *A_d;
double *result;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(double),
hipHostMallocCoherent));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
B_h = reinterpret_cast<double*>(malloc(sizeof(double)));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&result), sizeof(double)));
struct {
double* p;
double* res;
} args_f{A_d, result};
auto size = sizeof(args_f);
void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f,
HIP_LAUNCH_PARAM_BUFFER_SIZE,
&size, HIP_LAUNCH_PARAM_END};
hipModuleLaunchKernel(fmaxkernel, 1, 1, 1, 1, 1, 1, 0,
nullptr, nullptr, config_d);
HIP_CHECK(hipDeviceSynchronize());
HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost));
REQUIRE(*B_h == 0);
REQUIRE(A_h[0] == INITIAL_VAL);
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipFree(result));
free(B_h);
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
/*
This testcase verifies the builtinAtomic fmin API with RTC
1. Non Coherent memory with memory type as global
2. Memory type as flat
Input: A_h with INITIAL_VAL
Output: Return val would be initial val of A_h and the input value of
API would be updated with the max value
A_h would be 10, B_h would be INITIAL_VAL
*/
TEST_CASE("Unit_BuiltinAtomicsRTC_fminNonCoherentGlobalFlatMem") {
int mem_type = GENERATE(0, 1);
hipDeviceProp_t prop;
int device;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
hiprtcProgram prog;
if (mem_type) {
hiprtcCreateProgram(&prog, // prog
fminGlobalMem, // buffer
"kernel.cu", // name
0, nullptr, nullptr);
} else {
hiprtcCreateProgram(&prog, // prog
fminFlatMem, // buffer
"kernel.cu", // name
0, nullptr, nullptr);
}
std::string sarg = std::string("--gpu-architecture=") + prop.gcnArchName;
const char* options[] = {sarg.c_str()};
hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, options)};
size_t logSize;
HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize));
if (logSize) {
std::string log(logSize, '\0');
HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0]));
INFO(log);
}
REQUIRE(compileResult == HIPRTC_SUCCESS);
size_t codeSize;
HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize));
std::vector<char> code(codeSize);
HIPRTC_CHECK(hiprtcGetCode(prog, code.data()));
HIPRTC_CHECK(hiprtcDestroyProgram(&prog));
hipModule_t module;
hipFunction_t fmaxkernel;
HIP_CHECK(hipModuleLoadData(&module, code.data()));
if (mem_type) {
HIP_CHECK(hipModuleGetFunction(&fmaxkernel, module,
"unsafeAtomicMin_GlobalMem"));
} else {
HIP_CHECK(hipModuleGetFunction(&fmaxkernel, module,
"unsafeAtomicMin_FlatMem"));
}
double *A_h, *B_h;
double *A_d;
double *result;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(double),
hipHostMallocNonCoherent));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
B_h = reinterpret_cast<double*>(malloc(sizeof(double)));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&result), sizeof(double)));
struct {
double* p;
double* res;
} args_f{A_d, result};
auto size = sizeof(args_f);
void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f,
HIP_LAUNCH_PARAM_BUFFER_SIZE,
&size, HIP_LAUNCH_PARAM_END};
hipModuleLaunchKernel(fmaxkernel, 1, 1, 1, 1, 1, 1, 0,
nullptr, nullptr, config_d);
HIP_CHECK(hipDeviceSynchronize());
HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost));
REQUIRE(*B_h == INITIAL_VAL);
REQUIRE(A_h[0] == INITIAL_VAL);
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipFree(result));
free(B_h);
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
+34 -1
Ver fichero
@@ -10,7 +10,6 @@ set(TEST_SRC
popc.cc
ldg.cc
threadfence_system.cc
hipTestDeviceSymbol.cc
)
# skipped for windows compiler issue - Illegal instruction detected
@@ -31,10 +30,44 @@ set(AMD_TEST_SRC
bitInsert.cc
floatTM.cc
)
set(AMD_ARCH_SPEC_TEST_SRC
AtomicAdd_Coherent_withunsafeflag.cc
AtomicAdd_Coherent_withoutflag.cc
AtomicAdd_Coherent_withnoUnsafeflag.cc
AtomicAdd_NonCoherent_withoutflag.cc
AtomicAdd_NonCoherent_withnoUnsafeflag.cc
AtomicAdd_NonCoherent_withunsafeflag.cc
BuiltIns_fmax.cc
BuiltIns_fmin.cc
BuiltIns_fadd.cc
unsafeAtomicAdd_RTC.cc
unsafeAtomicAdd_Coherent_withunsafeflag.cc
unsafeAtomicAdd_Coherent_withoutflag.cc
unsafeAtomicAdd_Coherent_withnounsafeflag.cc
unsafeAtomicAdd_NonCoherent_withoutflag.cc
unsafeAtomicAdd_NonCoherent_withnounsafeflag.cc
unsafeAtomicAdd_NonCoherent_withunsafeflag.cc
)
if(HIP_PLATFORM MATCHES "amd")
string(FIND ${OFFLOAD_ARCH_STR} "gfx90a" ARCH_CHECK)
set(TEST_SRC ${TEST_SRC} ${AMD_TEST_SRC})
set_source_files_properties(floatTM.cc PROPERTIES COMPILE_FLAGS -std=c++17)
if(${ARCH_CHECK} GREATER_EQUAL 0)
set(TEST_SRC ${TEST_SRC} ${AMD_ARCH_SPEC_TEST_SRC})
set_source_files_properties(AtomicAdd_Coherent_withunsafeflag.cc PROPERTIES COMPILE_OPTIONS "-munsafe-fp-atomics")
set_source_files_properties(AtomicAdd_NonCoherent_withunsafeflag.cc PROPERTIES COMPILE_OPTIONS "-munsafe-fp-atomics")
set_source_files_properties(AtomicAdd_Coherent_withnoUnsafeflag.cc PROPERTIES COMPILE_OPTIONS "-mno-unsafe-fp-atomics")
set_source_files_properties(AtomicAdd_NonCoherent_withnoUnsafeflag.cc PROPERTIES COMPILE_OPTIONS "-mno-unsafe-fp-atomics")
set_source_files_properties(unsafeAtomicAdd_Coherent_withunsafeflag.cc PROPERTIES COMPILE_OPTIONS "-munsafe-fp-atomics")
set_source_files_properties(unsafeAtomicAdd_NonCoherent_withunsafeflag.cc PROPERTIES COMPILE_OPTIONS "-munsafe-fp-atomics")
set_source_files_properties(unsafeAtomicAdd_Coherent_withnounsafeflag.cc PROPERTIES COMPILE_OPTIONS "-mno-unsafe-fp-atomics")
set_source_files_properties(unsafeAtomicAdd_NonCoherent_withnounsafeflag.cc PROPERTIES COMPILE_OPTIONS "-mno-unsafe-fp-atomics")
file(GLOB AtomicAdd_files *AtomicAdd_*_*.cc)
set_property(SOURCE ${AtomicAdd_files} PROPERTY COMPILE_FLAGS --save-temps)
file(GLOB unsafeAtomicAdd_files *unsafeAtomicAdd_*_*.cc)
set_property(SOURCE ${unsafeAtomicAdd_files} PROPERTY COMPILE_FLAGS --save-temps)
endif()
hip_add_exe_to_target(NAME UnitDeviceTests
TEST_SRC ${TEST_SRC}
TEST_TARGET_NAME build_tests)
+72 -31
Ver fichero
@@ -1,5 +1,5 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2022 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
@@ -18,8 +18,12 @@ THE SOFTWARE.
*/
/* Test Case Description: Calling hipMemcpyTo/FromSymbolAsync() using user
declared stream obj and hipStreamPerThread*/
/*
Test Scenarios :
1) Calling hipMemcpyTo/FromSymbolAsync() using user declared stream obj and hipStreamPerThread.
2) Validate get symbol address/size for global const array.
3) Validate get symbol address/size for static const variable.
*/
#include <hip_test_common.hh>
#define NUM 1024
@@ -34,12 +38,20 @@ __global__ void Assign(int* Out) {
globalOut[tid] = globalIn[tid];
}
__device__ __constant__ int globalConst[NUM];
__device__ __constant__ int globalConstArr[NUM];
__device__ __constant__ static float statConstVar = 1.0f;
__global__ void checkAddress(int* addr, bool* out) {
*out = (globalConst == addr);
__global__ void checkGlobalConstAddress(int* addr, bool* out) {
*out = (globalConstArr == addr);
}
__global__ void checkStaticConstVarAddress(float* addr, bool* out) {
*out = (&statConstVar == addr);
}
/**
Calling hipMemcpyTo/FromSymbolAsync() using user declared stream obj and hipStreamPerThread.
*/
TEST_CASE("Unit_hipMemcpyToSymbolAsync_ToNFrom") {
int *A, *Am, *B, *Ad, *C, *Cm;
A = new int[NUM];
@@ -51,9 +63,9 @@ TEST_CASE("Unit_hipMemcpyToSymbolAsync_ToNFrom") {
C[i] = 0;
}
HIP_CHECK(hipMalloc((void**)&Ad, SIZE));
HIP_CHECK(hipHostMalloc((void**)&Am, SIZE));
HIP_CHECK(hipHostMalloc((void**)&Cm, SIZE));
HIP_CHECK(hipMalloc(&Ad, SIZE));
HIP_CHECK(hipHostMalloc(&Am, SIZE));
HIP_CHECK(hipHostMalloc(&Cm, SIZE));
for (int i = 0; i < NUM; i++) {
Am[i] = -1 * i;
Cm[i] = 0;
@@ -70,8 +82,8 @@ TEST_CASE("Unit_hipMemcpyToSymbolAsync_ToNFrom") {
hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipStreamSynchronize(stream));
for (int i = 0; i < NUM; i++) {
assert(Am[i] == B[i]);
assert(Am[i] == Cm[i]);
REQUIRE(Am[i] == B[i]);
REQUIRE(Am[i] == Cm[i]);
}
for (int i = 0; i < NUM; i++) {
@@ -86,8 +98,8 @@ TEST_CASE("Unit_hipMemcpyToSymbolAsync_ToNFrom") {
HIP_CHECK(hipMemcpyFromSymbol(C, HIP_SYMBOL(globalOut), SIZE, 0,
hipMemcpyDeviceToHost));
for (int i = 0; i < NUM; i++) {
assert(A[i] == B[i]);
assert(A[i] == C[i]);
REQUIRE(A[i] == B[i]);
REQUIRE(A[i] == C[i]);
}
for (int i = 0; i < NUM; i++) {
@@ -106,32 +118,19 @@ TEST_CASE("Unit_hipMemcpyToSymbolAsync_ToNFrom") {
}
SECTION("Calling hipMemcpyTo/FromSymbol using hipStreamPerThread") {
HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), A, SIZE, 0,
hipMemcpyHostToDevice, hipStreamPerThread));
hipMemcpyHostToDevice, hipStreamPerThread));
HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpyFromSymbolAsync(C, HIP_SYMBOL(globalOut), SIZE, 0,
hipMemcpyDeviceToHost, hipStreamPerThread));
hipMemcpyDeviceToHost, hipStreamPerThread));
HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
}
for (int i = 0; i < NUM; i++) {
assert(A[i] == B[i]);
assert(A[i] == C[i]);
REQUIRE(A[i] == B[i]);
REQUIRE(A[i] == C[i]);
}
bool *checkOkD;
bool checkOk = false;
size_t symbolSize = 0;
int *symbolAddress;
HIP_CHECK(hipGetSymbolSize(&symbolSize, HIP_SYMBOL(globalConst)));
HIP_CHECK(hipGetSymbolAddress((void**) &symbolAddress, HIP_SYMBOL(globalConst)));
HIP_CHECK(hipMalloc((void**)&checkOkD, sizeof(bool)));
hipLaunchKernelGGL(checkAddress, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, symbolAddress, checkOkD);
HIP_CHECK(hipMemcpy(&checkOk, checkOkD, sizeof(bool), hipMemcpyDeviceToHost));
HIP_CHECK(hipFree(checkOkD));
HIP_ASSERT(checkOk);
HIP_ASSERT((symbolSize == SIZE));
HIP_CHECK(hipStreamDestroy(stream));
HIP_CHECK(hipHostFree(Am));
HIP_CHECK(hipHostFree(Cm));
HIP_CHECK(hipFree(Ad));
@@ -139,3 +138,45 @@ TEST_CASE("Unit_hipMemcpyToSymbolAsync_ToNFrom") {
delete[] B;
delete[] C;
}
/**
1) Validate get symbol address/size for global const array.
2) Validate get symbol address/size for static const variable.
*/
TEST_CASE("Unit_hipGetSymbolAddressAndSize_Validation") {
bool *checkOkD;
bool checkOk = false;
size_t symbolSize{};
int *symbolArrAddress{};
float *symbolVarAddress{};
SECTION("Validate symbol size/address of global const array") {
HIP_CHECK(hipGetSymbolSize(&symbolSize, HIP_SYMBOL(globalConstArr)));
HIP_CHECK(hipGetSymbolAddress(
reinterpret_cast<void **>(&symbolArrAddress),
HIP_SYMBOL(globalConstArr)));
HIP_CHECK(hipMalloc(&checkOkD, sizeof(bool)));
hipLaunchKernelGGL(checkGlobalConstAddress, dim3(1, 1, 1), dim3(1, 1, 1),
0, 0, symbolArrAddress, checkOkD);
HIP_CHECK(hipMemcpy(&checkOk, checkOkD, sizeof(bool),
hipMemcpyDeviceToHost));
HIP_CHECK(hipFree(checkOkD));
HIP_ASSERT(checkOk);
HIP_ASSERT(symbolSize == SIZE);
}
SECTION("Validate symbol size/address of static const variable") {
HIP_CHECK(hipGetSymbolSize(&symbolSize, HIP_SYMBOL(statConstVar)));
HIP_CHECK(hipGetSymbolAddress(
reinterpret_cast<void **>(&symbolVarAddress),
HIP_SYMBOL(statConstVar)));
HIP_CHECK(hipMalloc(&checkOkD, sizeof(bool)));
hipLaunchKernelGGL(checkStaticConstVarAddress, dim3(1, 1, 1),
dim3(1, 1, 1), 0, 0, symbolVarAddress, checkOkD);
HIP_CHECK(hipMemcpy(&checkOk, checkOkD, sizeof(bool),
hipMemcpyDeviceToHost));
HIP_CHECK(hipFree(checkOkD));
HIP_ASSERT(checkOk);
HIP_ASSERT(symbolSize == sizeof(float));
}
}
@@ -0,0 +1,99 @@
/*
Copyright (c) 2022 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.
*/
/*
AtomicAdd on FineGrainMemory
1. The following test scenario verifies
unsafeatomicAdd on fineGrain memory with -mno-unsafe-fp-atomics flag
This testcase works only on gfx90a.
*/
#include<hip_test_common.hh>
#include<hip_test_checkers.hh>
#define INC_VAL 10
#define INITIAL_VAL 5
template<typename T>
static __global__ void AtomicCheck(T* Ad, T* result) {
T inc_val = 10;
*result = unsafeAtomicAdd(Ad, inc_val);
}
/*unsafeatomicAdd API for the fine grained memory variable
with -mno-unsafe-fp-atomics flag
Input: Ad{5}, INC_VAL{10}
Output: unsafeatomicAdd API would return 0 and the 0/P is 5
Generate the assembly file and check whether
atomic add instruction is generated
or not */
TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_CoherentwithnoUnsafeflag", "",
float, double) {
hipDeviceProp_t prop;
int device;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
TestType *A_h{nullptr}, *result{nullptr};
TestType *A_d{nullptr}, *result_d{nullptr};
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(TestType),
hipHostMallocCoherent));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&result),
sizeof(TestType),
hipHostMallocCoherent));
result[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&result_d),
result, 0));
hipLaunchKernelGGL(AtomicCheck<TestType>, dim3(1), dim3(1),
0, 0, A_d,
result_d);
HIP_CHECK(hipDeviceSynchronize());
bool testResult;
if ((std::is_same<TestType, float>::value)) {
testResult = HipTest::assemblyFile_Verification<TestType>(
"unsafeAtomicAdd_Coherent_withnounsafeflag-hip-amdgcn(.*)\\.s",
"global_atomic_add_f32");
REQUIRE(testResult == true);
} else {
testResult = HipTest::assemblyFile_Verification<TestType>(
"unsafeAtomicAdd_Coherent_withnounsafeflag-hip-amdgcn(.*)\\.s",
"global_atomic_add_f64");
REQUIRE(testResult == true);
}
REQUIRE(A_h[0] == INITIAL_VAL);
REQUIRE(result[0] == 0);
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipHostFree(result));
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
@@ -0,0 +1,99 @@
/*
Copyright (c) 2022 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.
*/
/*
AtomicAdd on FineGrainMemory
1. The following test scenario verifies
unsafeatomicAdd on fineGrain memory without atomics flag
This testcase works only on gfx90a.
*/
#include<hip_test_checkers.hh>
#include<hip_test_common.hh>
#define INC_VAL 10
#define INITIAL_VAL 5
template<typename T>
static __global__ void AtomicCheck(T* Ad, T* result) {
T inc_val = 10;
*result = unsafeAtomicAdd(Ad, inc_val);
}
/*unsafeatomicAdd API for the fine grained memory variable
without atomics flag
Input: Ad{5}, INC_VAL{10}
Output: unsafeatomicAdd API would return 0 and the 0/P is 5
Generate the assembly file and check whether
atomic add instruction is generated
or not */
TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_Coherentwithoutflag", "",
float, double) {
hipDeviceProp_t prop;
int device;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
TestType *A_h{nullptr}, *result{nullptr};
TestType *A_d{nullptr}, *result_d{nullptr};
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(TestType),
hipHostMallocCoherent));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&result),
sizeof(TestType),
hipHostMallocCoherent));
result[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&result_d),
result, 0));
hipLaunchKernelGGL(AtomicCheck<TestType>, dim3(1), dim3(1),
0, 0, A_d,
result_d);
HIP_CHECK(hipDeviceSynchronize());
bool testResult;
if ((std::is_same<TestType, float>::value)) {
testResult = HipTest::assemblyFile_Verification<TestType>(
"unsafeAtomicAdd_Coherent_withoutflag-hip-amdgcn(.*)\\.s",
"global_atomic_add_f32");
REQUIRE(testResult == true);
} else {
testResult = HipTest::assemblyFile_Verification<TestType>(
"unsafeAtomicAdd_Coherent_withoutflag-hip-amdgcn(.*)\\.s",
"global_atomic_add_f64");
REQUIRE(testResult == true);
}
REQUIRE(A_h[0] == INITIAL_VAL);
REQUIRE(result[0] == 0);
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipHostFree(result));
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
@@ -0,0 +1,100 @@
/*
Copyright (c) 2022 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.
*/
/*
AtomicAdd on FineGrainMemory
1. The following test scenario verifies
unsafeatomicAdd on fineGrain memory with -munsafe-fp-atomics flag
This testcase works only on gfx90a.
*/
#include<hip_test_checkers.hh>
#include<hip_test_common.hh>
#define INC_VAL 10
#define INITIAL_VAL 5
template<typename T>
static __global__ void AtomicCheck(T* Ad, T* result) {
T inc_val = 10;
*result = unsafeAtomicAdd(Ad, inc_val);
}
/*unsafeatomicAdd API for the fine grained memory variable
with -m-unsafe-atomics flag
Input: Ad{5}, INC_VAL{10}
Output: atomicAdd API would return 0 and the 0/P is 5
Generate the assembly file and check whether
atomic add instruction is generated
or not */
TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_CoherentwithUnsafeflag", "",
float, double) {
hipDeviceProp_t prop;
int device;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
TestType *A_h{nullptr}, *result{nullptr};
TestType *A_d{nullptr}, *result_d{nullptr};
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(TestType),
hipHostMallocCoherent));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&result),
sizeof(TestType),
hipHostMallocCoherent));
result[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&result_d),
result, 0));
hipLaunchKernelGGL(AtomicCheck<TestType>, dim3(1), dim3(1),
0, 0, A_d,
result_d);
HIP_CHECK(hipDeviceSynchronize());
bool testResult;
if ((std::is_same<TestType, float>::value)) {
testResult = HipTest::assemblyFile_Verification<TestType>(
"unsafeAtomicAdd_Coherent_withunsafeflag-hip-amdgcn(.*)\\.s",
"global_atomic_add_f32");
REQUIRE(testResult == true);
} else {
testResult = HipTest::assemblyFile_Verification<TestType>(
"unsafeAtomicAdd_Coherent_withunsafeflag-hip-amdgcn(.*)\\.s",
"global_atomic_add_f64");
REQUIRE(testResult == true);
}
REQUIRE(A_h[0] == INITIAL_VAL);
REQUIRE(result[0] == 0);
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipHostFree(result));
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
@@ -0,0 +1,97 @@
/*
Copyright (c) 2022 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.
*/
/*
AtomicAdd on CoarseGrainMemory
1. The following test scenario verifies
unsafeAtomicAdd on CoarseGrain memory with -mno-unsafe-fp-atomics flag
This testcase works only on gfx90a.
*/
#include<hip_test_checkers.hh>
#include<hip_test_common.hh>
#define INC_VAL 10
#define INITIAL_VAL 5
template<typename T>
static __global__ void AtomicCheck(T* Ad, T* result) {
T inc_val = 10;
*result = unsafeAtomicAdd(Ad, inc_val);
}
/*unsafeAtomicAdd API for the coarse grained memory variable
with -mno-unsafe-fp-atomics flag
Input: Ad{5}, INC_VAL{10}
Output: unsafeAtomicAdd API would work and the 0/P is INITIAL_VAL + INC_VAL
Generate the assembly file and check whether
global_atomic_add instruction is generated
or not */
TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_NonCoherentnounsafeatomicsflag", "",
float, double) {
hipDeviceProp_t prop;
int device;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
TestType *A_h{nullptr}, *result{nullptr};
TestType *A_d{nullptr}, *result_d{nullptr};
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(TestType),
hipHostMallocNonCoherent));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&result),
sizeof(TestType),
hipHostMallocNonCoherent));
result[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&result_d),
result, 0));
hipLaunchKernelGGL(AtomicCheck<TestType>,
dim3(1), dim3(1),
0, 0, A_d,
result_d);
HIP_CHECK(hipDeviceSynchronize());
bool testResult;
REQUIRE(A_h[0] == INITIAL_VAL + INC_VAL);
REQUIRE(result[0] == INITIAL_VAL);
if ((std::is_same<TestType, float>::value)) {
testResult = HipTest::assemblyFile_Verification<TestType>(
"unsafeAtomicAdd_NonCoherent_withnounsafeflag-hip-amdgcn(.*)\\.s",
"global_atomic_add_f32");
REQUIRE(testResult == true);
} else {
testResult = HipTest::assemblyFile_Verification<TestType>(
"unsafeAtomicAdd_NonCoherent_withnounsafeflag-hip-amdgcn(.*)\\.s",
"global_atomic_add_f64");
REQUIRE(testResult == true);
}
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipHostFree(result));
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
@@ -0,0 +1,97 @@
/*
Copyright (c) 2022 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.
*/
/*
unsafeAtomicAdd on CoarseGrainMemory
1. The following test scenario verifies
unsafeAtomicAdd on CoarseGrain memory without any unsafeatomics flag
This testcase works only on gfx90a.
*/
#include<hip_test_checkers.hh>
#include<hip_test_common.hh>
#define INC_VAL 10
#define INITIAL_VAL 5
template<typename T>
static __global__ void AtomicCheck(T* Ad, T* result) {
T inc_val = 10;
*result = unsafeAtomicAdd(Ad, inc_val);
}
/*unsafeAtomicAdd API for the coarse grained memory variable
without any flag
Input: Ad{5}, INC_VAL{10}
Output: unsafeAtomicAdd API would work and the 0/P is INITIAL_VAL + INC_VAL
Generate the assembly file and check whether
global_atomic_cmpswap instruction is generated
or not */
TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_NonCoherentwithoutflag", "",
float, double) {
hipDeviceProp_t prop;
int device;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
TestType *A_h{nullptr}, *result{nullptr};
TestType *A_d{nullptr}, *result_d{nullptr};
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(TestType),
hipHostMallocNonCoherent));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&result),
sizeof(TestType),
hipHostMallocNonCoherent));
result[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&result_d),
result, 0));
hipLaunchKernelGGL(AtomicCheck<TestType>,
dim3(1), dim3(1),
0, 0, A_d,
result_d);
HIP_CHECK(hipDeviceSynchronize());
bool testResult;
REQUIRE(A_h[0] == INITIAL_VAL + INC_VAL);
REQUIRE(result[0] == INITIAL_VAL);
if ((std::is_same<TestType, float>::value)) {
testResult = HipTest::assemblyFile_Verification<TestType>(
"unsafeAtomicAdd_NonCoherent_withoutflag-hip-amdgcn(.*)\\.s",
"global_atomic_add_f32");
REQUIRE(testResult == true);
} else {
testResult = HipTest::assemblyFile_Verification<TestType>(
"unsafeAtomicAdd_NonCoherent_withoutflag-hip-amdgcn(.*)\\.s",
"global_atomic_add_f64");
REQUIRE(testResult == true);
}
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipHostFree(result));
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
@@ -0,0 +1,97 @@
/*
Copyright (c) 2022 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.
*/
/*
unsafeAtomicAdd on CoarseGrainMemory
1. The following test scenario verifies
unsafeAtomicAdd on CoarseGrain memory with unsafeatomics flag
This testcase works only on gfx90a.
*/
#include<hip_test_checkers.hh>
#include<hip_test_common.hh>
#define INC_VAL 10
#define INITIAL_VAL 5
template<typename T>
static __global__ void AtomicCheck(T* Ad, T* result) {
T inc_val = 10;
*result = unsafeAtomicAdd(Ad, inc_val);
}
/*unsafeAtomicAdd API for the coarse grained memory variable
with -munsafe-fp-atomics flag
Input: Ad{5}, INC_VAL{10}
Output: unsafeAtomicAdd API would work and the 0/P is INITIAL_VAL + INC_VAL
Generate the assembly file and check whether
global_atomic_add instruction is generated
or not */
TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_NonCoherentwithunsafeatomicsflag", "",
float, double) {
hipDeviceProp_t prop;
int device;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
TestType *A_h{nullptr}, *result{nullptr};
TestType *A_d{nullptr}, *result_d{nullptr};
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(TestType),
hipHostMallocNonCoherent));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&result),
sizeof(TestType),
hipHostMallocNonCoherent));
result[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&result_d),
result, 0));
hipLaunchKernelGGL(AtomicCheck<TestType>,
dim3(1), dim3(1),
0, 0, A_d,
result_d);
HIP_CHECK(hipDeviceSynchronize());
bool testResult;
REQUIRE(A_h[0] == INITIAL_VAL + INC_VAL);
REQUIRE(result[0] == INITIAL_VAL);
if ((std::is_same<TestType, float>::value)) {
testResult = HipTest::assemblyFile_Verification<TestType>(
"unsafeAtomicAdd_NonCoherent_withunsafeflag-hip-amdgcn(.*)\\.s",
"global_atomic_add_f32");
REQUIRE(testResult == true);
} else {
testResult = HipTest::assemblyFile_Verification<TestType>(
"unsafeAtomicAdd_NonCoherent_withunsafeflag-hip-amdgcn(.*)\\.s",
"global_atomic_add_f64");
REQUIRE(testResult == true);
}
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipHostFree(result));
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
@@ -0,0 +1,583 @@
/*
Copyright (c) 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.
*/
/*
unsafeAtomicAdd Scenarios with hipRTC:
1. FineGrainMemory with -m-nounsafe-fp-atomics flag
2. FineGrainMemory without compilation flag
3. FineGrainMemory without -munsafe-fp-atomics flag
4. CoarseGrainMemory with -m-nounsafe-fp-atomics flag
5. CoarseGrainMemory without compilation flag
6. CoarseGrainMemory without -munsafe-fp-atomics flag
*/
#include<hip_test_checkers.hh>
#include<hip_test_common.hh>
#include <hip/hiprtc.h>
#define INCREMENT_VAL 10
#define INITIAL_VAL 5
static constexpr auto fkernel{
R"(
extern "C"
__global__ void AtomicCheck(float* Ad, float *result) {
*result = unsafeAtomicAdd(Ad, 10);
}
)"};
static constexpr auto dkernel{
R"(
extern "C"
__global__ void AtomicCheck(double* Ad, double *result) {
*result = unsafeAtomicAdd(Ad, 10);
}
)"};
/*
Test unsafeAtomicAdd API for the fine grained memory variable
where kernel is compiled using hipRTC and with
compilation flag -mno-unsafe-fp-atomics.
Input: Ad{5}, INCREMENT_VAL{10}
Output: unsafeAtomicAdd API will not work and returns 0 so
the initial value will be intact. expected O/P is 5
*/
TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_CoherentRTCnounsafeatomicflag", "",
float, double) {
int device = 0;
hipDeviceProp_t props;
HIP_CHECK(hipGetDeviceProperties(&props, device));
std::string gfxName(props.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
hiprtcProgram prog;
if (std::is_same<TestType, float>::value) {
hiprtcCreateProgram(&prog, // prog
fkernel, // buffer
"kernel.cu", // name
0, nullptr, nullptr);
} else {
hiprtcCreateProgram(&prog, // prog
dkernel, // buffer
"kernel.cu", // name
0, nullptr, nullptr);
}
std::string sarg = std::string("--gpu-architecture=") + props.gcnArchName;
const char* options[] = {sarg.c_str(), "-mno-unsafe-fp-atomics"};
hiprtcResult compileResult{hiprtcCompileProgram(prog, 2, options)};
size_t logSize;
HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize));
if (logSize) {
std::string log(logSize, '\0');
HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0]));
INFO(log);
}
REQUIRE(compileResult == HIPRTC_SUCCESS);
size_t codeSize;
HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize));
std::vector<char> code(codeSize);
HIPRTC_CHECK(hiprtcGetCode(prog, code.data()));
HIPRTC_CHECK(hiprtcDestroyProgram(&prog));
hipModule_t module;
hipFunction_t f_kernel;
HIP_CHECK(hipModuleLoadData(&module, code.data()));
HIP_CHECK(hipModuleGetFunction(&f_kernel, module, "AtomicCheck"));
if (props.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
TestType *A_h, *result;
TestType *A_d, *result_d;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(TestType),
hipHostMallocCoherent));
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&result),
sizeof(TestType),
hipHostMallocCoherent));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&result_d),
result, 0));
struct {
TestType* p;
TestType* result;
} args_f{A_d, result_d};
auto size = sizeof(args_f);
void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f,
HIP_LAUNCH_PARAM_BUFFER_SIZE,
&size, HIP_LAUNCH_PARAM_END};
hipModuleLaunchKernel(f_kernel, 1, 1, 1, 1, 1, 1, 0,
nullptr, nullptr, config_d);
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(A_h[0] == INITIAL_VAL);
REQUIRE(*result == 0);
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipHostFree(result));
}
HIP_CHECK(hipModuleUnload(module));
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
/*
Test unsafeAtomicAdd API for the fine grained memory variable
where kernel is compiled using hipRTC and with
compilation flag -munsafe-fp-atomics.
Input: Ad{5}, INCREMENT_VAL{10}
Output: unsafeAtomicAdd API will not work and r`eturns 0 so
the initial value will be intact. expected O/P is 5
*/
TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_CoherentRTCunsafeatomicflag", "",
float, double) {
int device = 0;
hipDeviceProp_t props;
HIP_CHECK(hipGetDeviceProperties(&props, device));
std::string gfxName(props.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
hiprtcProgram prog;
if (std::is_same<TestType, float>::value) {
hiprtcCreateProgram(&prog, // prog
fkernel, // buffer
"kernel.cu", // name
0, nullptr, nullptr);
} else {
hiprtcCreateProgram(&prog, // prog
dkernel, // buffer
"kernel.cu", // name
0, nullptr, nullptr);
}
std::string sarg = std::string("--gpu-architecture=") + props.gcnArchName;
const char* options[] = {sarg.c_str(), "-munsafe-fp-atomics"};
hiprtcResult compileResult{hiprtcCompileProgram(prog, 2, options)};
size_t logSize;
HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize));
if (logSize) {
std::string log(logSize, '\0');
HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0]));
INFO(log);
}
REQUIRE(compileResult == HIPRTC_SUCCESS);
size_t codeSize;
HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize));
std::vector<char> code(codeSize);
HIPRTC_CHECK(hiprtcGetCode(prog, code.data()));
HIPRTC_CHECK(hiprtcDestroyProgram(&prog));
hipModule_t module;
hipFunction_t f_kernel;
HIP_CHECK(hipModuleLoadData(&module, code.data()));
HIP_CHECK(hipModuleGetFunction(&f_kernel, module, "AtomicCheck"));
if (props.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
TestType *A_h, *result;
TestType *A_d, *result_d;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(TestType),
hipHostMallocCoherent));
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&result),
sizeof(TestType),
hipHostMallocCoherent));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&result_d),
result, 0));
struct {
TestType* p;
TestType* result;
} args_f{A_d, result_d};
auto size = sizeof(args_f);
void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f,
HIP_LAUNCH_PARAM_BUFFER_SIZE,
&size, HIP_LAUNCH_PARAM_END};
hipModuleLaunchKernel(f_kernel, 1, 1, 1, 1, 1, 1, 0,
nullptr, nullptr, config_d);
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(A_h[0] == INITIAL_VAL);
REQUIRE(*result == 0);
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipHostFree(result));
}
HIP_CHECK(hipModuleUnload(module));
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
/* Test unsafeAtomicAdd API for the fine grained memory variable
where kernel is compiled using hipRTC and without compilation flag
Input: Ad{5}, INCREMENT_VAL{10}
Output: unsafeAtomicAdd API will not work and returns 0 so
the initial value will be intact. expected O/P is 5*/
TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_CoherentRTCwithoutflag", "",
float, double) {
int device = 0;
hipDeviceProp_t props;
HIP_CHECK(hipGetDeviceProperties(&props, device));
std::string gfxName(props.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
hiprtcProgram prog;
if (std::is_same<TestType, float>::value) {
hiprtcCreateProgram(&prog, // prog
fkernel, // buffer
"kernel.cu", // name
0, nullptr, nullptr);
} else {
hiprtcCreateProgram(&prog, // prog
dkernel, // buffer
"kernel.cu", // name
0, nullptr, nullptr);
}
std::string sarg = std::string("--gpu-architecture=") + props.gcnArchName;
const char* options[] = {sarg.c_str()};
hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, options)};
size_t logSize;
HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize));
if (logSize) {
std::string log(logSize, '\0');
HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0]));
INFO(log);
}
REQUIRE(compileResult == HIPRTC_SUCCESS);
size_t codeSize;
HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize));
std::vector<char> code(codeSize);
HIPRTC_CHECK(hiprtcGetCode(prog, code.data()));
HIPRTC_CHECK(hiprtcDestroyProgram(&prog));
hipModule_t module;
hipFunction_t f_kernel;
HIP_CHECK(hipModuleLoadData(&module, code.data()));
HIP_CHECK(hipModuleGetFunction(&f_kernel, module, "AtomicCheck"));
if (props.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
TestType *A_h, *result;
TestType *A_d, *result_d;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(float),
hipHostMallocCoherent));
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&result), sizeof(float),
hipHostMallocCoherent));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&result_d),
result, 0));
struct {
TestType* p;
TestType* result;
} args_f{A_d, result_d};
auto size = sizeof(args_f);
void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f,
HIP_LAUNCH_PARAM_BUFFER_SIZE,
&size, HIP_LAUNCH_PARAM_END};
hipModuleLaunchKernel(f_kernel, 1, 1, 1, 1, 1,
1, 0, nullptr, nullptr, config_d);
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(A_h[0] == INITIAL_VAL);
REQUIRE(*result == 0);
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipHostFree(result));
}
HIP_CHECK(hipModuleUnload(module));
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
/*
Test unsafeAtomicAdd API for the coarse grained memory variable where kernel
is compiled using hipRTC and with compilation flag -mno-unsafe-fp-atomics
Input: Ad{5}, INCREMENT_VAL{10}
Output: Expected O/P is 15 */
TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_NonCoherentRTCnounsafeatomicflag", "",
float, double) {
int device = 0;
hipDeviceProp_t props;
HIP_CHECK(hipGetDeviceProperties(&props, device));
std::string gfxName(props.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
hiprtcProgram prog;
if (std::is_same<TestType, float>::value) {
hiprtcCreateProgram(&prog, // prog
fkernel, // buffer
"kernel.cu", // name
0, nullptr, nullptr);
} else {
hiprtcCreateProgram(&prog, // prog
dkernel, // buffer
"kernel.cu", // name
0, nullptr, nullptr);
}
std::string sarg = std::string("--gpu-architecture=") + props.gcnArchName;
const char* options[] = {sarg.c_str(), "-mno-unsafe-fp-atomics"};
hiprtcResult compileResult{hiprtcCompileProgram(prog, 2, options)};
size_t logSize;
HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize));
if (logSize) {
std::string log(logSize, '\0');
HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0]));
INFO(log);
}
REQUIRE(compileResult == HIPRTC_SUCCESS);
size_t codeSize;
HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize));
std::vector<char> code(codeSize);
HIPRTC_CHECK(hiprtcGetCode(prog, code.data()));
HIPRTC_CHECK(hiprtcDestroyProgram(&prog));
hipModule_t module;
hipFunction_t f_kernel;
HIP_CHECK(hipModuleLoadData(&module, code.data()));
HIP_CHECK(hipModuleGetFunction(&f_kernel, module, "AtomicCheck"));
if (props.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
TestType *A_h, *result;
TestType *A_d, *result_d;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(TestType),
hipHostMallocNonCoherent));
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&result),
sizeof(TestType)));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&result_d),
result, 0));
struct {
TestType* p;
TestType* result;
} args_f{A_d, result_d};
auto size = sizeof(args_f);
void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f,
HIP_LAUNCH_PARAM_BUFFER_SIZE,
&size, HIP_LAUNCH_PARAM_END};
hipModuleLaunchKernel(f_kernel, 1, 1, 1, 1, 1, 1, 0,
nullptr, nullptr, config_d);
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(A_h[0] == INITIAL_VAL + INCREMENT_VAL);
REQUIRE(*result == INITIAL_VAL);
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipHostFree(result));
}
HIP_CHECK(hipModuleUnload(module));
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
/*
Test unsafeAtomicAdd API for the coarse grained memory variable where kernel
is compiled using hipRTC and with compilation flag -munsafe-fp-atomics
Input: Ad{5}, INCREMENT_VAL{10}
Output: Expected O/P is 15 */
TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_NonCoherentRTCunsafeatomicflag", "",
float, double) {
int device = 0;
hipDeviceProp_t props;
HIP_CHECK(hipGetDeviceProperties(&props, device));
std::string gfxName(props.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
hiprtcProgram prog;
if (std::is_same<TestType, float>::value) {
hiprtcCreateProgram(&prog, // prog
fkernel, // buffer
"kernel.cu", // name
0, nullptr, nullptr);
} else {
hiprtcCreateProgram(&prog, // prog
dkernel, // buffer
"kernel.cu", // name
0, nullptr, nullptr);
}
std::string sarg = std::string("--gpu-architecture=") + props.gcnArchName;
const char* options[] = {sarg.c_str(), "-munsafe-fp-atomics"};
hiprtcResult compileResult{hiprtcCompileProgram(prog, 2, options)};
size_t logSize;
HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize));
if (logSize) {
std::string log(logSize, '\0');
HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0]));
INFO(log);
}
REQUIRE(compileResult == HIPRTC_SUCCESS);
size_t codeSize;
HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize));
std::vector<char> code(codeSize);
HIPRTC_CHECK(hiprtcGetCode(prog, code.data()));
HIPRTC_CHECK(hiprtcDestroyProgram(&prog));
hipModule_t module;
hipFunction_t f_kernel;
HIP_CHECK(hipModuleLoadData(&module, code.data()));
HIP_CHECK(hipModuleGetFunction(&f_kernel, module, "AtomicCheck"));
if (props.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
TestType *A_h, *result;
TestType *A_d, *result_d;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(TestType),
hipHostMallocNonCoherent));
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&result),
sizeof(TestType)));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&result_d),
result, 0));
struct {
TestType* p;
TestType* result;
} args_f{A_d, result_d};
auto size = sizeof(args_f);
void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f,
HIP_LAUNCH_PARAM_BUFFER_SIZE,
&size, HIP_LAUNCH_PARAM_END};
hipModuleLaunchKernel(f_kernel, 1, 1, 1, 1, 1, 1, 0,
nullptr, nullptr, config_d);
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(A_h[0] == INITIAL_VAL + INCREMENT_VAL);
REQUIRE(*result == INITIAL_VAL);
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipHostFree(result));
}
HIP_CHECK(hipModuleUnload(module));
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
/*
Test unsafeAtomicAdd API for the coarse grained memory variable
where kernel is compiled using hipRTC and without compilation flag
Input: Ad{5}, INCREMENT_VAL{10}
Output: O/P is 15 */
TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_NonCoherentRTC", "",
float, double) {
int device = 0;
hipDeviceProp_t props;
HIP_CHECK(hipGetDeviceProperties(&props, device));
std::string gfxName(props.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
hiprtcProgram prog;
if (std::is_same<TestType, float>::value) {
hiprtcCreateProgram(&prog, // prog
fkernel, // buffer
"kernel.cu", // name
0, nullptr, nullptr);
} else {
hiprtcCreateProgram(&prog, // prog
dkernel, // buffer
"kernel.cu", // name
0, nullptr, nullptr);
}
std::string sarg = std::string("--gpu-architecture=") + props.gcnArchName;
const char* options[] = {sarg.c_str()};
hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, options)};
size_t logSize;
HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize));
if (logSize) {
std::string log(logSize, '\0');
HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0]));
INFO(log);
}
REQUIRE(compileResult == HIPRTC_SUCCESS);
size_t codeSize;
HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize));
std::vector<char> code(codeSize);
HIPRTC_CHECK(hiprtcGetCode(prog, code.data()));
HIPRTC_CHECK(hiprtcDestroyProgram(&prog));
hipModule_t module;
hipFunction_t f_kernel;
HIP_CHECK(hipModuleLoadData(&module, code.data()));
HIP_CHECK(hipModuleGetFunction(&f_kernel, module, "AtomicCheck"));
if (props.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
TestType *A_h, *result;
TestType *A_d, *result_d;
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), sizeof(TestType),
hipHostMallocNonCoherent));
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&result),
sizeof(TestType)));
A_h[0] = INITIAL_VAL;
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
A_h, 0));
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&result_d),
result, 0));
struct {
TestType* p;
TestType* result;
} args_f{A_d, result_d};
auto size = sizeof(args_f);
void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f,
HIP_LAUNCH_PARAM_BUFFER_SIZE,
&size, HIP_LAUNCH_PARAM_END};
hipModuleLaunchKernel(f_kernel, 1, 1, 1, 1, 1, 1, 0,
nullptr, nullptr, config_d);
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(A_h[0] == INITIAL_VAL + INCREMENT_VAL);
REQUIRE(*result == INITIAL_VAL);
HIP_CHECK(hipHostFree(A_h));
HIP_CHECK(hipHostFree(result));
}
HIP_CHECK(hipModuleUnload(module));
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
"skipping the testcase for this GPU " << device);
}
}
+9
Ver fichero
@@ -42,6 +42,15 @@ set(TEST_SRC
hipGraphAddMemcpyNode1D.cc
hipGraphAddChildGraphNode.cc
hipGraphNodeGetType.cc
hipGraphExecMemcpyNodeSetParams1D.cc
hipGraphGetEdges.cc
hipGraphRemoveDependencies.cc
hipGraphInstantiate.cc
hipGraphExecUpdate.cc
hipGraphExecEventRecordNodeSetEvent.cc
hipGraphMemsetNodeGetParams.cc
hipGraphMemsetNodeSetParams.cc
hipGraphExecMemcpyNodeSetParamsFromSymbol.cc
)
hip_add_exe_to_target(NAME GraphsTest
@@ -20,6 +20,7 @@ THE SOFTWARE.
/**
Testcase Scenarios :
1) Create and add empty node to graph and verify addition is successful.
2) Negative Scenarios
*/
#include <hip_test_common.hh>
@@ -55,3 +56,45 @@ TEST_CASE("Unit_hipGraphAddEmptyNode_Functional") {
HIP_CHECK(hipFree(pOutBuff_d));
HIP_CHECK(hipGraphDestroy(graph));
}
/**
* Negative Scenarios hipGraphAddEmptyNode
*/
TEST_CASE("Unit_hipGraphAddEmptyNode_NegTest") {
char *pOutBuff_d{};
constexpr size_t size = 1024;
hipGraph_t graph;
hipGraphNode_t memsetNode{}, emptyNode{};
std::vector<hipGraphNode_t> dependencies;
HIP_CHECK(hipMalloc(&pOutBuff_d, size));
hipMemsetParams memsetParams{};
memsetParams.dst = reinterpret_cast<void*>(pOutBuff_d);
memsetParams.value = 0;
memsetParams.pitch = 0;
memsetParams.elementSize = sizeof(char);
memsetParams.width = size * sizeof(char);
memsetParams.height = 1;
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
&memsetParams));
dependencies.push_back(memsetNode);
// pGraphNode is nullptr
SECTION("Null Empty Graph Node") {
REQUIRE(hipErrorInvalidValue == hipGraphAddEmptyNode(nullptr, graph,
dependencies.data(), dependencies.size()));
}
// graph is nullptr
SECTION("Null Graph") {
REQUIRE(hipErrorInvalidValue == hipGraphAddEmptyNode(&emptyNode, nullptr,
dependencies.data(), dependencies.size()));
}
// pDependencies is nullptr
SECTION("Null Graph") {
REQUIRE(hipErrorInvalidValue == hipGraphAddEmptyNode(&emptyNode, graph,
nullptr, dependencies.size()));
}
HIP_CHECK(hipFree(pOutBuff_d));
HIP_CHECK(hipGraphDestroy(graph));
}
@@ -0,0 +1,289 @@
/*
Copyright (c) 2022 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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/**
Testcase Scenarios :
1) Create a graph with event record nodes as follows:
event_record_start_node(event1) --> MemcpyH2DNode --> kernel -->
MemcpyD2HNode --> event_record_stop_node(event2).Instantiate the graph.
Set a different event 'event3' in event_record_stop_node using
hipGraphExecEventRecordNodeSetEvent. Launch the graph. Verify the
hipGraphExecEventRecordNodeSetEvent functionality by measuring the
time difference between event2 & event1 and between event3 and
event1.
2) Scenario to verify that hipGraphExecEventRecordNodeSetEvent does not
impact the graph and changes only the executable graph.
Create an event record node with event1 and add it to graph. Instantiate
the graph to create an executable graph. Change the event in the
executable graph to event2. Verify that the event record node still
contains event1.
3) Negative Scenarios
- Input executable graph is a nullptr.
- Input node is a nullptr.
- Input event to set is a nullptr.
- Input executable graph is uninitialized.
- Input node is uninitialized.
- Input event is uninitialized.
- Event record node does not exist in graph.
- Input node is a memset node.
- Input node is a event wait node.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip_test_kernels.hh>
#define GRID_DIM 512
#define BLK_DIM 512
#define LEN (GRID_DIM * BLK_DIM)
/**
* Kernel Functions to copy.
*/
static __global__ void copy_ker_func(int* a, int* b) {
int tx = hipBlockIdx_x*hipBlockDim_x + hipThreadIdx_x;
if (tx < LEN) b[tx] = a[tx];
}
/**
* Scenario 1: Functional scenario (See description Above)
*/
TEST_CASE("Unit_hipGraphExecEventRecordNodeSetEvent_Functional") {
size_t memsize = LEN*sizeof(int);
hipGraph_t graph;
HIP_CHECK(hipGraphCreate(&graph, 0));
// Create events
hipEvent_t event_start, event1_end, event2_end;
HIP_CHECK(hipEventCreate(&event_start));
HIP_CHECK(hipEventCreate(&event1_end));
HIP_CHECK(hipEventCreate(&event2_end));
// Create nodes with event_start and event1_end
hipGraphNode_t event_start_rec, event_end_rec;
HIP_CHECK(hipGraphAddEventRecordNode(&event_start_rec, graph, nullptr, 0,
event_start));
HIP_CHECK(hipGraphAddEventRecordNode(&event_end_rec, graph, nullptr, 0,
event1_end));
int *inp_h, *inp_d, *out_h, *out_d;
// Allocate host buffers
inp_h = reinterpret_cast<int*>(malloc(memsize));
REQUIRE(inp_h != nullptr);
out_h = reinterpret_cast<int*>(malloc(memsize));
REQUIRE(out_h != nullptr);
// Allocate device buffers
HIP_CHECK(hipMalloc(&inp_d, memsize));
HIP_CHECK(hipMalloc(&out_d, memsize));
// Initialize host buffer
for (uint32_t i = 0; i < LEN; i++) {
inp_h[i] = i;
out_h[i] = 0;
}
// graph creation ...........
// Create memcpy and kernel nodes for graph
hipGraphNode_t memcpyH2D, memcpyD2H, kernelnode;
hipKernelNodeParams kernelNodeParams{};
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D, graph, nullptr, 0, inp_d,
inp_h, memsize, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H, graph, nullptr, 0,
out_h, out_d, memsize, hipMemcpyDeviceToHost));
void* kernelArgs1[] = {&inp_d, &out_d};
kernelNodeParams.func = reinterpret_cast<void *>(copy_ker_func);
kernelNodeParams.gridDim = dim3(GRID_DIM);
kernelNodeParams.blockDim = dim3(BLK_DIM);
kernelNodeParams.sharedMemBytes = 0;
kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs1);
kernelNodeParams.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kernelnode, graph, nullptr, 0,
&kernelNodeParams));
// Create dependencies for graph
HIP_CHECK(hipGraphAddDependencies(graph, &event_start_rec,
&memcpyH2D, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D,
&kernelnode, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &kernelnode,
&memcpyD2H, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &memcpyD2H,
&event_end_rec, 1));
// Instantiate and launch the graph
hipStream_t streamForGraph;
hipGraphExec_t graphExec;
HIP_CHECK(hipStreamCreate(&streamForGraph));
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
// Change the event at event_end_rec node to event2_end
HIP_CHECK(hipGraphExecEventRecordNodeSetEvent(graphExec,
event_end_rec, event2_end));
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
// Wait for graph to complete
HIP_CHECK(hipStreamSynchronize(streamForGraph));
// Validate output
bool btestPassed = true;
for (uint32_t i = 0; i < LEN; i++) {
if (out_h[i] != inp_h[i]) {
btestPassed = false;
break;
}
}
REQUIRE(btestPassed == true);
// Validate the changed events
float t = 0.0f;
HIP_CHECK(hipEventElapsedTime(&t, event_start, event2_end));
REQUIRE(t > 0.0f);
// Since event1_end is never recorded, hipEventElapsedTime
// should return error code.
REQUIRE(hipErrorInvalidResourceHandle ==
hipEventElapsedTime(&t, event_start, event1_end));
// Free resources
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipStreamDestroy(streamForGraph));
HIP_CHECK(hipFree(inp_d));
HIP_CHECK(hipFree(out_d));
free(inp_h);
free(out_h);
HIP_CHECK(hipEventDestroy(event_start));
HIP_CHECK(hipEventDestroy(event1_end));
HIP_CHECK(hipEventDestroy(event2_end));
HIP_CHECK(hipGraphDestroy(graph));
}
/**
* Scenario 2: This test verifies that changes to executable graph does
* not impact the original graph.
*/
TEST_CASE("Unit_hipGraphExecEventRecordNodeSetEvent_VerifyEventNotChanged") {
hipGraph_t graph;
HIP_CHECK(hipGraphCreate(&graph, 0));
hipEvent_t event1, event2, event_out;
HIP_CHECK(hipEventCreate(&event1));
HIP_CHECK(hipEventCreate(&event2));
hipGraphNode_t eventrec;
HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0,
event1));
hipGraphExec_t graphExec;
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphExecEventRecordNodeSetEvent(graphExec,
eventrec, event2));
HIP_CHECK(hipGraphEventRecordNodeGetEvent(eventrec, &event_out));
// validate set event and get event are same
REQUIRE(event1 == event_out);
// Instantiate and launch the graph
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipEventDestroy(event2));
HIP_CHECK(hipEventDestroy(event1));
}
/**
* Scenario 3: Negative Tests
*/
TEST_CASE("Unit_hipGraphExecEventRecordNodeSetEvent_Negative") {
hipGraph_t graph;
HIP_CHECK(hipGraphCreate(&graph, 0));
hipEvent_t event1, event2;
HIP_CHECK(hipEventCreate(&event1));
HIP_CHECK(hipEventCreate(&event2));
hipGraphNode_t eventrec;
HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0,
event1));
// Create memset
constexpr size_t Nbytes = 1024;
char *A_d;
hipGraphNode_t memset_A;
hipMemsetParams memsetParams{};
HIP_CHECK(hipMalloc(&A_d, Nbytes));
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void*>(A_d);
memsetParams.value = 0;
memsetParams.pitch = 0;
memsetParams.elementSize = sizeof(char);
memsetParams.width = Nbytes;
memsetParams.height = 1;
hipGraphExec_t graphExec;
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
SECTION("hGraphExec = nullptr") {
REQUIRE(hipErrorInvalidValue ==
hipGraphExecEventRecordNodeSetEvent(nullptr, eventrec, event2));
}
SECTION("hNode = nullptr") {
REQUIRE(hipErrorInvalidValue ==
hipGraphExecEventRecordNodeSetEvent(graphExec, nullptr, event2));
}
SECTION("event = nullptr") {
REQUIRE(hipErrorInvalidValue ==
hipGraphExecEventRecordNodeSetEvent(graphExec, eventrec, nullptr));
}
SECTION("hGraphExec is uninitialized") {
hipGraphExec_t graphExec1{};
REQUIRE(hipErrorInvalidValue ==
hipGraphExecEventRecordNodeSetEvent(graphExec1, eventrec, event2));
}
SECTION("hNode is uninitialized") {
hipGraphNode_t dummy{};
REQUIRE(hipErrorInvalidValue ==
hipGraphExecEventRecordNodeSetEvent(graphExec, dummy, event2));
}
SECTION("event is uninitialized") {
hipEvent_t event_dummy{};
REQUIRE(hipErrorInvalidValue ==
hipGraphExecEventRecordNodeSetEvent(graphExec, eventrec,
event_dummy));
}
SECTION("event record node does not exist") {
hipGraph_t graph1;
HIP_CHECK(hipGraphCreate(&graph1, 0));
HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph1, nullptr, 0,
&memsetParams));
hipGraphExec_t graphExec1;
HIP_CHECK(hipGraphInstantiate(&graphExec1, graph1, nullptr, nullptr, 0));
REQUIRE(hipErrorInvalidValue ==
hipGraphExecEventRecordNodeSetEvent(graphExec1, eventrec, event2));
HIP_CHECK(hipGraphExecDestroy(graphExec1));
HIP_CHECK(hipGraphDestroy(graph1));
}
SECTION("pass memset node as hNode") {
HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0,
&memsetParams));
REQUIRE(hipErrorInvalidValue ==
hipGraphExecEventRecordNodeSetEvent(graphExec, memset_A, event2));
}
SECTION("pass event wait node as hNode") {
hipGraphNode_t event_wait_node;
HIP_CHECK(hipGraphAddEventWaitNode(&event_wait_node, graph, nullptr, 0,
event1));
REQUIRE(hipErrorInvalidValue ==
hipGraphExecEventRecordNodeSetEvent(graphExec, event_wait_node,
event2));
}
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipEventDestroy(event1));
HIP_CHECK(hipEventDestroy(event2));
}
@@ -0,0 +1,198 @@
/*
Copyright (c) 2022 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.
*/
/*
Testcase Scenarios :
Functional-
1) Instantiate a graph with memcpy node, obtain executable graph and update the
node params with set exec api call. Make sure they are taking effect.
Negative-
1) Pass hGraphExec as nullptr and check if api returns error.
2) Pass GraphNode as nullptr and check if api returns error.
3) Pass destination ptr is nullptr, api expected to return error code.
4) Pass source ptr is nullptr, api expected to return error code.
5) Pass count as zero, api expected to return error code.
6) Pass same pointer as source ptr and destination ptr, api expected to return error code.
7) Pass overlap memory address as source ptr and destination ptr, api expected to return error code.
7) Pass overlap memory as source ptr and destination ptr where source ptr is ahead of destination ptr, api expected to return error code.
8) Pass overlap memory as source ptr and destination ptr where destination ptr is ahead of source ptr, api expected to return error code.
9) If count is more than allocated size for source and destination ptr, api should return error code.
10) If count is less than allocated size for source and destination ptr, api should return error code.
11) Change the hipMemcpyKind from H2D to D2H but allocate pointer memory for H2D, api should return error code.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip_test_kernels.hh>
/* Test verifies hipGraphExecMemcpyNodeSetParams1D API Negative scenarios.
*/
TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParams1D_Negative") {
constexpr size_t N = 1024;
constexpr size_t Nbytes = N * sizeof(int);
int *A_d;
HIP_CHECK(hipMalloc(&A_d, Nbytes));
int *A_h = reinterpret_cast<int*>(malloc(Nbytes));
REQUIRE(A_h != nullptr);
memset(A_h, 0, Nbytes);
hipError_t ret;
hipGraphNode_t memcpyH2D;
hipGraph_t graph;
hipGraphExec_t graphExec;
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D, graph, nullptr, 0, A_d, A_h,
Nbytes, hipMemcpyHostToDevice));
// Instantiate the graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0));
SECTION("Pass hGraphExec as nullptr") {
ret = hipGraphExecMemcpyNodeSetParams1D(nullptr, memcpyH2D, A_d, A_h,
Nbytes, hipMemcpyHostToDevice);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass GraphNode as nullptr") {
ret = hipGraphExecMemcpyNodeSetParams1D(graphExec, nullptr, A_d, A_h,
Nbytes, hipMemcpyHostToDevice);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass destination ptr is nullptr") {
ret = hipGraphExecMemcpyNodeSetParams1D(graphExec, memcpyH2D, nullptr, A_h,
Nbytes, hipMemcpyHostToDevice);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass source ptr is nullptr") {
ret = hipGraphExecMemcpyNodeSetParams1D(graphExec, memcpyH2D, A_d, nullptr,
Nbytes, hipMemcpyHostToDevice);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass count as zero") {
ret = hipGraphExecMemcpyNodeSetParams1D(graphExec, memcpyH2D, A_d, A_h,
0, hipMemcpyHostToDevice);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass same pointer as source ptr and destination ptr") {
ret = hipGraphExecMemcpyNodeSetParams1D(graphExec, memcpyH2D, A_d, A_d,
Nbytes, hipMemcpyDeviceToDevice);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass overlap memory where destination ptr is ahead of source ptr") {
ret = hipGraphExecMemcpyNodeSetParams1D(graphExec, memcpyH2D, A_d, A_d-5,
Nbytes, hipMemcpyDeviceToDevice);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass overlap memory where source ptr is ahead of destination ptr") {
ret = hipGraphExecMemcpyNodeSetParams1D(graphExec, memcpyH2D, A_d+5, A_d,
Nbytes, hipMemcpyDeviceToDevice);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Copy more than allocated memory") {
ret = hipGraphExecMemcpyNodeSetParams1D(graphExec, memcpyH2D, A_d, A_h,
Nbytes+8, hipMemcpyHostToDevice);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Copy less than allocated memory") {
ret = hipGraphExecMemcpyNodeSetParams1D(graphExec, memcpyH2D, A_d, A_h,
Nbytes-8, hipMemcpyHostToDevice);
REQUIRE(hipSuccess == ret);
}
SECTION("Change the hipMemcpyKind from H2D to D2H") {
ret = hipGraphExecMemcpyNodeSetParams1D(graphExec, memcpyH2D, A_d, A_h,
Nbytes, hipMemcpyDeviceToHost);
REQUIRE(hipSuccess != ret);
}
HIP_CHECK(hipFree(A_d));
free(A_h);
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipGraphDestroy(graph));
}
/* Test verifies hipGraphExecMemcpyNodeSetParams1D API Functional scenarios.
*/
TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParams1D_Functional") {
constexpr size_t N = 1024;
constexpr size_t Nbytes = N * sizeof(int);
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
size_t NElem{N};
int *hData = reinterpret_cast<int*>(malloc(Nbytes));
REQUIRE(hData != nullptr);
memset(hData, 0, Nbytes);
hipGraphNode_t memcpyH2D_A, memcpyH2D_B, memcpyD2H_C;
hipGraphNode_t kernel_vecAdd;
hipKernelNodeParams kernelNodeParams{};
hipGraph_t graph;
hipGraphExec_t graphExec;
hipStream_t streamForGraph;
HIP_CHECK(hipStreamCreate(&streamForGraph));
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, C_h, C_d,
Nbytes, hipMemcpyDeviceToHost));
void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
kernelNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
kernelNodeParams.gridDim = dim3(blocks);
kernelNodeParams.blockDim = dim3(threadsPerBlock);
kernelNodeParams.sharedMemBytes = 0;
kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs2);
kernelNodeParams.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0,
&kernelNodeParams));
// Create dependencies
HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &kernel_vecAdd, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &kernel_vecAdd, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2H_C, 1));
// Instantiate the graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphExecMemcpyNodeSetParams1D(graphExec, memcpyD2H_C, hData,
C_d, Nbytes, hipMemcpyDeviceToHost));
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
HIP_CHECK(hipStreamSynchronize(streamForGraph));
// Verify graph execution result
HipTest::checkVectorADD(A_h, B_h, hData, N);
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipStreamDestroy(streamForGraph));
HIP_CHECK(hipGraphDestroy(graph));
free(hData);
}
@@ -0,0 +1,317 @@
/*
Copyright (c) 2022 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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/**
Testcase Scenarios of hipGraphExecMemcpyNodeSetParamsFromSymbol API:
Functional
1) Allocate global symbol memory, Instantiate a graph with memcpy node,
obtain executable graph and update the node params with set exec api call.
Make sure they are taking effect.
2) Allocate const symbol memory, Instantiate a graph with memcpy node,
obtain executable graph and update the node params with set exec api call.
Make sure they are taking effect.
Negative
1) Pass hGraphExec as nullptr and check if api returns error.
2) Pass GraphNode as nullptr and check if api returns error.
3) Pass destination ptr as nullptr, api expected to return error code.
4) Pass symbol ptr as nullptr, api expected to return error code.
5) Pass count as zero, api expected to return error code.
6) Pass offset+count greater than allocated size, api expected to return error code.
7) Pass same symbol pointer as source ptr and destination ptr, api expected to return error code.
8) Pass Pass both dstn ptr and source ptr as 2 different symbol ptr, api expected to return error code.
9) Copy from device ptr to host ptr but pass kind as different, api expected to return error code.
10) Check with other graph node but pass same graphExec, api expected to return error code.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <limits>
#define SIZE 256
__device__ int globalIn[SIZE];
__device__ int globalOut[SIZE];
__device__ __constant__ int globalConst[SIZE];
/* Test verifies hipGraphExecMemcpyNodeSetParamsFromSymbol API Negative scenarios.
*/
TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParamsFromSymbol_Negative") {
constexpr size_t Nbytes = SIZE * sizeof(int);
int *A_d{nullptr}, *B_d{nullptr};
int *A_h{nullptr}, *B_h{nullptr};
HipTest::initArrays<int>(&A_d, &B_d, nullptr,
&A_h, &B_h, nullptr, SIZE, false);
hipError_t ret;
hipGraph_t graph;
hipGraphExec_t graphExec;
hipGraphNode_t memcpyToSymbolNode, memcpyFromSymbolNode, memcpyH2D_A;
std::vector<hipGraphNode_t> dependencies;
HIP_CHECK(hipGraphCreate(&graph, 0));
// Adding MemcpyNode
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h,
Nbytes, hipMemcpyHostToDevice));
dependencies.push_back(memcpyH2D_A);
// Adding MemcpyNodeToSymbol
HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph,
dependencies.data(),
dependencies.size(),
HIP_SYMBOL(globalIn),
A_d, Nbytes, 0,
hipMemcpyDeviceToDevice));
dependencies.clear();
dependencies.push_back(memcpyToSymbolNode);
HIP_CHECK(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph,
dependencies.data(),
dependencies.size(),
B_h,
HIP_SYMBOL(globalConst),
Nbytes, 0,
hipMemcpyDeviceToHost));
HIP_CHECK(hipGraphMemcpyNodeSetParamsFromSymbol(memcpyFromSymbolNode,
B_d,
HIP_SYMBOL(globalIn),
Nbytes, 0,
hipMemcpyDeviceToDevice));
// Instantiate the graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
#if HT_NVIDIA
SECTION("Pass hGraphExec as nullptr") {
ret = hipGraphExecMemcpyNodeSetParamsFromSymbol(nullptr,
memcpyFromSymbolNode, B_d,
HIP_SYMBOL(globalConst),
Nbytes, 0,
hipMemcpyDeviceToDevice);
REQUIRE(hipErrorInvalidValue == ret);
}
#endif
SECTION("Pass GraphNode as nullptr") {
ret = hipGraphExecMemcpyNodeSetParamsFromSymbol(graphExec,
nullptr, B_d,
HIP_SYMBOL(globalConst),
Nbytes, 0,
hipMemcpyDeviceToDevice);
REQUIRE(hipErrorInvalidValue == ret);
}
#if HT_NVIDIA
SECTION("Pass destination ptr as nullptr") {
ret = hipGraphExecMemcpyNodeSetParamsFromSymbol(graphExec,
memcpyFromSymbolNode, nullptr,
HIP_SYMBOL(globalConst),
Nbytes, 0,
hipMemcpyDeviceToDevice);
REQUIRE(hipErrorInvalidValue == ret);
}
#endif
SECTION("Pass symbol ptr as nullptr") {
ret = hipGraphExecMemcpyNodeSetParamsFromSymbol(graphExec,
memcpyFromSymbolNode, B_d,
nullptr,
Nbytes, 0,
hipMemcpyDeviceToDevice);
REQUIRE(hipErrorInvalidSymbol == ret);
}
#if HT_NVIDIA
SECTION("Pass count as zero") {
ret = hipGraphExecMemcpyNodeSetParamsFromSymbol(graphExec,
memcpyFromSymbolNode, B_d,
HIP_SYMBOL(globalConst),
0, 0,
hipMemcpyDeviceToDevice);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass offset+count greater than allocated size") {
ret = hipGraphExecMemcpyNodeSetParamsFromSymbol(graphExec,
memcpyFromSymbolNode, B_d,
HIP_SYMBOL(globalConst),
Nbytes, 10,
hipMemcpyDeviceToDevice);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass same symbol pointer as source and destination ptr") {
ret = hipGraphExecMemcpyNodeSetParamsFromSymbol(graphExec,
memcpyFromSymbolNode,
HIP_SYMBOL(globalIn),
HIP_SYMBOL(globalIn),
Nbytes, 0,
hipMemcpyDeviceToDevice);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass both dstn ptr and source ptr as 2 different symbol ptr") {
ret = hipGraphExecMemcpyNodeSetParamsFromSymbol(graphExec,
memcpyFromSymbolNode,
HIP_SYMBOL(globalIn),
HIP_SYMBOL(globalOut),
Nbytes, 0,
hipMemcpyDeviceToDevice);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Copy from device ptr to host ptr but pass kind as different") {
ret = hipGraphExecMemcpyNodeSetParamsFromSymbol(graphExec,
memcpyFromSymbolNode,
B_h,
HIP_SYMBOL(globalOut),
Nbytes, 0,
hipMemcpyDeviceToDevice);
REQUIRE(hipErrorInvalidValue == ret);
}
#endif
SECTION("Check with other graph node") {
hipGraph_t graph1;
hipGraphNode_t memcpyFromSymbolNode1{};
HIP_CHECK(hipGraphCreate(&graph1, 0));
HIP_CHECK(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode1, graph1,
nullptr,
0,
B_h,
HIP_SYMBOL(globalConst),
Nbytes, 0,
hipMemcpyDeviceToHost));
ret = hipGraphExecMemcpyNodeSetParamsFromSymbol(graphExec,
memcpyFromSymbolNode1,
B_d,
HIP_SYMBOL(globalOut),
Nbytes, 0,
hipMemcpyDeviceToDevice);
REQUIRE(hipErrorInvalidValue == ret);
}
HipTest::freeArrays<int>(A_d, B_d, nullptr,
A_h, B_h, nullptr, false);
HIP_CHECK(hipGraphDestroy(graph));
}
static
void hipGraphExecMemcpyNodeSetParamsFromSymbol_GlobalMem(bool useConstVar) {
constexpr size_t Nbytes = SIZE * sizeof(int);
hipGraphNode_t memcpyD2H_B;
int *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr};
int *A_h{nullptr}, *B_h{nullptr};
HipTest::initArrays<int>(&A_d, &B_d, &C_d,
&A_h, &B_h, nullptr, SIZE, false);
hipGraph_t graph;
hipGraphExec_t graphExec;
hipGraphNode_t memcpyToSymbolNode, memcpyFromSymbolNode, memcpyH2D_A;
std::vector<hipGraphNode_t> dependencies;
HIP_CHECK(hipGraphCreate(&graph, 0));
// Adding MemcpyNode
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h,
Nbytes, hipMemcpyHostToDevice));
dependencies.push_back(memcpyH2D_A);
if (useConstVar) {
HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph,
dependencies.data(),
dependencies.size(),
HIP_SYMBOL(globalConst),
A_d, Nbytes, 0,
hipMemcpyDeviceToDevice));
} else {
HIP_CHECK(hipGraphAddMemcpyNodeToSymbol(&memcpyToSymbolNode, graph,
dependencies.data(),
dependencies.size(),
HIP_SYMBOL(globalIn),
A_d, Nbytes, 0,
hipMemcpyDeviceToDevice));
}
dependencies.clear();
dependencies.push_back(memcpyToSymbolNode);
if (useConstVar) {
HIP_CHECK(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph,
dependencies.data(),
dependencies.size(),
C_d,
HIP_SYMBOL(globalConst),
Nbytes, 0,
hipMemcpyDeviceToDevice));
} else {
HIP_CHECK(hipGraphAddMemcpyNodeFromSymbol(&memcpyFromSymbolNode, graph,
dependencies.data(),
dependencies.size(),
C_d,
HIP_SYMBOL(globalIn),
Nbytes, 0,
hipMemcpyDeviceToDevice));
}
dependencies.clear();
dependencies.push_back(memcpyFromSymbolNode);
// Adding MemcpyNode
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_B, graph, dependencies.data(),
dependencies.size(), B_h, B_d,
Nbytes, hipMemcpyDeviceToHost));
// Instantiate and launch the graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
// Update the node with B_d destination pointer from C_d
if (useConstVar) {
HIP_CHECK(hipGraphExecMemcpyNodeSetParamsFromSymbol(graphExec,
memcpyFromSymbolNode,
B_d,
HIP_SYMBOL(globalConst),
Nbytes, 0,
hipMemcpyDeviceToDevice));
} else {
HIP_CHECK(hipGraphExecMemcpyNodeSetParamsFromSymbol(graphExec,
memcpyFromSymbolNode,
B_d,
HIP_SYMBOL(globalIn),
Nbytes, 0,
hipMemcpyDeviceToDevice));
}
HIP_CHECK(hipGraphLaunch(graphExec, 0));
// Validating the result
for (int i = 0; i < SIZE; i++) {
if (B_h[i] != A_h[i]) {
WARN("Validation failed B_h[i] " << B_h[i] << "A_h[i] " << A_h[i]);
REQUIRE(false);
}
}
HipTest::freeArrays<int>(A_d, B_d, C_d,
A_h, B_h, nullptr, false);
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipGraphDestroy(graph));
}
/* Test verifies hipGraphExecMemcpyNodeSetParamsFromSymbol Functional scenario.
1) Allocate global symbol memory, Instantiate a graph with memcpy node,
obtain executable graph and update the node params with set exec api call.
Make sure they are taking effect.
2) Allocate const symbol memory, Instantiate a graph with memcpy node,
obtain executable graph and update the node params with set exec api call.
Make sure they are taking effect.
*/
TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParamsFromSymbol_Functional") {
SECTION("Check and update with Global Device Symbol Memory") {
hipGraphExecMemcpyNodeSetParamsFromSymbol_GlobalMem(false);
}
SECTION("Check and update with Constant Global Device Symbol Memory") {
hipGraphExecMemcpyNodeSetParamsFromSymbol_GlobalMem(true);
}
}
+341
Ver fichero
@@ -0,0 +1,341 @@
/*
Copyright (c) 2022 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.
*/
/**
Testcase Scenarios :
Functional-
1) Make a clone of the created graph and update the executable-graph from a clone or same graph again.
2) Update the executable-graph from a graph and make sure they are taking effect.
Negative-
1) When Pass hGraphExec as nullptr and verify api returns error code.
2) When Pass hGraph as nullptr and verify api returns error code.
3) When Pass hErrorNode_out as nullptr and verify api returns error code.
4) When Pass updateResult_out as nullptr and verify api returns error code.
5) When the a graphExec was updated with with different type of node and verify api returns error code.
6) When a node is deleted in hGraph but not its pair from hGraphExec and verify api returns error code.
7) When a node is deleted in hGraphExec but not its pair from hGraph and verify api returns error code.
8) When grpah dependencies differ but graph have same node and verify api returns error code.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip_test_kernels.hh>
/* Test verifies hipGraphExecUpdate API Negative nullptr check scenarios.
*/
TEST_CASE("Unit_hipGraphExecUpdate_Negative_Basic") {
hipError_t ret;
hipGraph_t graph{};
hipGraphExec_t graphExec{};
hipGraphNode_t hErrorNode_out{};
hipGraphExecUpdateResult updateResult_out{};
SECTION("Pass hGraphExec as nullptr") {
ret = hipGraphExecUpdate(nullptr, graph, &hErrorNode_out,
&updateResult_out);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass hGraph as nullptr") {
ret = hipGraphExecUpdate(graphExec, nullptr, &hErrorNode_out,
&updateResult_out);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass hErrorNode_out as nullptr") {
ret = hipGraphExecUpdate(graphExec, graph, nullptr, &updateResult_out);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass updateResult_out as nullptr") {
ret = hipGraphExecUpdate(graphExec, graph, &hErrorNode_out, nullptr);
REQUIRE(hipErrorInvalidValue == ret);
}
}
/* Test verifies hipGraphExecUpdate API Negative scenarios.
When the a graphExec was updated with with different type of node
*/
TEST_CASE("Unit_hipGraphExecUpdate_Negative_TypeChange") {
constexpr size_t N = 1024;
constexpr size_t Nbytes = N * sizeof(char);
constexpr size_t val = 0;
char *devData;
int *A_d, *A_h;
HipTest::initArrays<int>(&A_d, nullptr, nullptr,
&A_h, nullptr, nullptr, N, false);
HIP_CHECK(hipMalloc(&devData, Nbytes));
hipGraph_t graph, graph2;
hipGraphExec_t graphExec;
hipStream_t streamForGraph;
hipGraphNode_t memsetNode, memcpy_A, hErrorNode_out;
hipError_t ret;
hipGraphExecUpdateResult updateResult_out;
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipStreamCreate(&streamForGraph));
hipMemsetParams memsetParams{};
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void*>(devData);
memsetParams.value = val;
memsetParams.pitch = 0;
memsetParams.elementSize = sizeof(char);
memsetParams.width = Nbytes;
memsetParams.height = 1;
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
&memsetParams));
std::vector<hipGraphNode_t> dependencies;
dependencies.push_back(memsetNode);
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphCreate(&graph2, 0));
HIP_CHECK(hipStreamCreate(&streamForGraph));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph2, nullptr, 0, A_d, A_h,
Nbytes, hipMemcpyHostToDevice));
// graphExec was created before memcpyTemp was added to graph.
ret = hipGraphExecUpdate(graphExec, graph2, &hErrorNode_out,
&updateResult_out);
REQUIRE(hipGraphExecUpdateErrorNodeTypeChanged == updateResult_out);
REQUIRE(hipErrorGraphExecUpdateFailure == ret);
HIP_CHECK(hipFree(devData));
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipGraphDestroy(graph2));
HIP_CHECK(hipStreamDestroy(streamForGraph));
}
/* Test verifies hipGraphExecUpdate API Negative scenarios.
When the count of nodes differ in hGraphExec and hGraph
*/
TEST_CASE("Unit_hipGraphExecUpdate_Negative_CountDiffer") {
constexpr size_t N = 1024;
constexpr size_t Nbytes = N * sizeof(int);
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
size_t NElem{N};
int *hData = reinterpret_cast<int*>(malloc(Nbytes));
REQUIRE(hData != nullptr);
memset(hData, 0, Nbytes);
hipGraphNode_t memcpy_A, memcpy_B, memcpy_C, memcpyTemp;
hipGraphNode_t kernel_vecAdd;
hipKernelNodeParams kernelNodeParams{};
hipError_t ret;
hipGraph_t graph1, graph2, graph3;
hipGraphExec_t graphExec1, graphExec2;
hipStream_t streamForGraph;
hipGraphNode_t hErrorNode_out;
hipGraphExecUpdateResult updateResult_out;
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HIP_CHECK(hipGraphCreate(&graph1, 0));
HIP_CHECK(hipStreamCreate(&streamForGraph));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph1, nullptr, 0, A_d, A_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph1, nullptr, 0, B_d, B_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph1, nullptr, 0, C_h, C_d,
Nbytes, hipMemcpyDeviceToHost));
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
kernelNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
kernelNodeParams.gridDim = dim3(blocks);
kernelNodeParams.blockDim = dim3(threadsPerBlock);
kernelNodeParams.sharedMemBytes = 0;
kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
kernelNodeParams.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph1, nullptr, 0,
&kernelNodeParams));
// Create dependencies
HIP_CHECK(hipGraphAddDependencies(graph1, &memcpy_A, &kernel_vecAdd, 1));
HIP_CHECK(hipGraphAddDependencies(graph1, &memcpy_B, &kernel_vecAdd, 1));
HIP_CHECK(hipGraphAddDependencies(graph1, &kernel_vecAdd, &memcpy_C, 1));
// Create a cloned graph and added extra node to it
HIP_CHECK(hipGraphClone(&graph2, graph1));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyTemp, graph2, nullptr, 0,
C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
HIP_CHECK(hipGraphInstantiate(&graphExec1, graph1, nullptr, nullptr, 0));
HIP_CHECK(hipGraphInstantiate(&graphExec2, graph2, nullptr, nullptr, 0));
SECTION("When a node deleted from Graph but not from its pair GraphExec") {
ret = hipGraphExecUpdate(graphExec2, graph1, &hErrorNode_out,
&updateResult_out);
REQUIRE(hipErrorGraphExecUpdateFailure == ret);
}
SECTION("When a node deleted from GraphExec but not from its pair Graph") {
ret = hipGraphExecUpdate(graphExec1, graph2, &hErrorNode_out,
&updateResult_out);
REQUIRE(hipErrorGraphExecUpdateFailure == ret);
}
SECTION("When the dependent nodes of a pair differ") {
HIP_CHECK(hipGraphCreate(&graph3, 0));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph3, nullptr, 0, A_d, A_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph3, nullptr, 0, B_d, B_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph3, nullptr, 0, C_h, C_d,
Nbytes, hipMemcpyDeviceToHost));
HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph3, nullptr, 0,
&kernelNodeParams));
// Create dependencies
HIP_CHECK(hipGraphAddDependencies(graph3, &memcpy_A, &kernel_vecAdd, 1));
HIP_CHECK(hipGraphAddDependencies(graph3, &memcpy_B, &kernel_vecAdd, 1));
HIP_CHECK(hipGraphAddDependencies(graph3, &memcpy_C, &kernel_vecAdd, 1));
ret = hipGraphExecUpdate(graphExec1, graph3, &hErrorNode_out,
&updateResult_out);
REQUIRE(hipErrorGraphExecUpdateFailure == ret);
HIP_CHECK(hipGraphDestroy(graph3));
}
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipGraphExecDestroy(graphExec1));
HIP_CHECK(hipGraphExecDestroy(graphExec2));
HIP_CHECK(hipStreamDestroy(streamForGraph));
HIP_CHECK(hipGraphDestroy(graph1));
HIP_CHECK(hipGraphDestroy(graph2));
free(hData);
}
/* Functional Scenario -
1) Make a clone of the created graph and update the executable-graph from a clone graph.
2) Update the executable-graph from a graph and make sure they are taking effect.
*/
TEST_CASE("Unit_hipGraphExecUpdate_Functional") {
constexpr size_t N = 1024;
constexpr size_t Nbytes = N * sizeof(int);
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
size_t NElem{N};
int *hData = reinterpret_cast<int*>(malloc(Nbytes));
REQUIRE(hData != nullptr);
memset(hData, 0, Nbytes);
hipGraphNode_t memcpy_A, memcpy_B, memcpy_C;
hipGraphNode_t kernel_vecAdd, kernel_vecSquare;
hipKernelNodeParams kernelNodeParams{};
hipGraph_t graph, graph2, clonedgraph{};
hipGraphExec_t graphExec;
hipStream_t streamForGraph;
hipGraphNode_t hErrorNode_out;
hipGraphExecUpdateResult updateResult_out;
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipStreamCreate(&streamForGraph));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph, nullptr, 0, A_d, A_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph, nullptr, 0, B_d, B_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph, nullptr, 0, C_h, C_d,
Nbytes, hipMemcpyDeviceToHost));
void* kernelArgs[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
kernelNodeParams.func =
reinterpret_cast<void *>(HipTest::vector_square<int>);
kernelNodeParams.gridDim = dim3(blocks);
kernelNodeParams.blockDim = dim3(threadsPerBlock);
kernelNodeParams.sharedMemBytes = 0;
kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs);
kernelNodeParams.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kernel_vecSquare, graph, nullptr, 0,
&kernelNodeParams));
// Create dependencies
HIP_CHECK(hipGraphAddDependencies(graph, &memcpy_A, &kernel_vecSquare, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &memcpy_B, &kernel_vecSquare, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecSquare, &memcpy_C, 1));
// Instantiate and launch the graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
SECTION("Update graphExec with clone graph") {
HIP_CHECK(hipGraphClone(&clonedgraph, graph));
HIP_CHECK(hipGraphExecUpdate(graphExec, clonedgraph, &hErrorNode_out,
&updateResult_out));
}
// Code for new graph creation with samilar node setup
HIP_CHECK(hipGraphCreate(&graph2, 0));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph2, nullptr, 0, A_d, A_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_B, graph2, nullptr, 0, B_d, B_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_C, graph2, nullptr, 0, C_h, C_d,
Nbytes, hipMemcpyDeviceToHost));
HIP_CHECK(hipGraphMemcpyNodeSetParams1D(memcpy_C, hData, C_d, Nbytes,
hipMemcpyDeviceToHost));
memset(&kernelNodeParams, 0, sizeof(hipKernelNodeParams));
void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
kernelNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
kernelNodeParams.gridDim = dim3(blocks);
kernelNodeParams.blockDim = dim3(threadsPerBlock);
kernelNodeParams.sharedMemBytes = 0;
kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs2);
kernelNodeParams.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph2, nullptr, 0,
&kernelNodeParams));
// Create dependencies
HIP_CHECK(hipGraphAddDependencies(graph2, &memcpy_A, &kernel_vecAdd, 1));
HIP_CHECK(hipGraphAddDependencies(graph2, &memcpy_B, &kernel_vecAdd, 1));
HIP_CHECK(hipGraphAddDependencies(graph2, &kernel_vecAdd, &memcpy_C, 1));
// Update the graphExec graph from graph -> graph2
HIP_CHECK(hipGraphExecUpdate(graphExec, graph2, &hErrorNode_out,
&updateResult_out));
REQUIRE(updateResult_out == hipGraphExecUpdateSuccess);
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
HIP_CHECK(hipStreamSynchronize(streamForGraph));
// Verify graph execution result
HipTest::checkVectorADD(A_h, B_h, hData, N);
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipStreamDestroy(streamForGraph));
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipGraphDestroy(graph2));
HIP_CHECK(hipGraphDestroy(clonedgraph));
free(hData);
}
+251
Ver fichero
@@ -0,0 +1,251 @@
/*
Copyright (c) 2022 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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/**
Testcase Scenarios :
1) Add nodes to graph with dependencies defined. Call api and verify number
of edges and from/to list returned corresponds to the dependencies defined.
2) Pass from and to as nullptr and verify the api returns number of edges.
3) Pass numEdges lesser than actual number and verify the api returns from/to
list with requested number of edges.
4) Pass numEdges greater than actual number and verify the remaining entries
in from/to list are set to null and number of edges actually returned will
be written to numEdges.
5) Validate numEdges when 0 or 1 node is present in graph.
6) Negative Test Cases
- Input graph parameter is a nullptr.
- From node parameter is a nullptr.
- To node parameter is a nullptr.
- numEdges parameter is a nullptr.
- Input graph parameter is uninitialized.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip_test_kernels.hh>
#define EXPECTED_NUM_OF_EDGES 6
/**
* Local Function to validate number of edges.
*/
static void validate_hipGraphGetEdges_fromto(size_t numEdgesToGet,
int testnum,
hipGraphNode_t *nodes_from,
hipGraphNode_t *nodes_to,
hipGraph_t graph) {
int numEdges = static_cast<int>(numEdgesToGet);
hipGraphNode_t *fromnode = new hipGraphNode_t[numEdges]{};
hipGraphNode_t *tonode = new hipGraphNode_t[numEdges]{};
hipGraphNode_t *expected_from_nodes = nodes_from;
hipGraphNode_t *expected_to_nodes = nodes_to;
HIP_CHECK(hipGraphGetEdges(graph, fromnode, tonode, &numEdgesToGet));
bool nodeFound;
int found_count = 0;
for (int idx_from = 0; idx_from < EXPECTED_NUM_OF_EDGES; idx_from++) {
nodeFound = false;
int idx = 0;
for (; idx < EXPECTED_NUM_OF_EDGES; idx++) {
if (expected_from_nodes[idx_from] == fromnode[idx]) {
nodeFound = true;
break;
}
}
if (nodeFound && (tonode[idx] == expected_to_nodes[idx_from])) {
found_count++;
}
}
// Validate
if (testnum == 0) {
REQUIRE(found_count == EXPECTED_NUM_OF_EDGES);
} else if (testnum == 1) {
REQUIRE(found_count == numEdges);
} else if (testnum == 2) {
REQUIRE(found_count == EXPECTED_NUM_OF_EDGES);
for (int idx = (EXPECTED_NUM_OF_EDGES - 1); idx > (numEdges - 1); idx++) {
REQUIRE(fromnode[idx] == nullptr);
REQUIRE(tonode[idx] == nullptr);
}
}
delete[] tonode;
delete[] fromnode;
}
/**
* Scenario 1: Finctionality tests to validate hipGraphGetEdges()
* for different number of edges.
*/
TEST_CASE("Unit_hipGraphGetEdges_Functionality") {
constexpr size_t N = 1024;
constexpr size_t Nbytes = N * sizeof(int);
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
hipGraph_t graph;
hipGraphNode_t memset_A, memset_B, memsetKer_C;
hipGraphNode_t memcpyH2D_A, memcpyH2D_B, memcpyD2H_C;
hipGraphNode_t kernel_vecAdd;
hipKernelNodeParams kernelNodeParams{};
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
hipMemsetParams memsetParams{};
int memsetVal{};
size_t NElem{N};
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HIP_CHECK(hipGraphCreate(&graph, 0));
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void*>(A_d);
memsetParams.value = 0;
memsetParams.pitch = 0;
memsetParams.elementSize = sizeof(char);
memsetParams.width = Nbytes;
memsetParams.height = 1;
HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0,
&memsetParams));
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void*>(B_d);
memsetParams.value = 0;
memsetParams.pitch = 0;
memsetParams.elementSize = sizeof(char);
memsetParams.width = Nbytes;
memsetParams.height = 1;
HIP_CHECK(hipGraphAddMemsetNode(&memset_B, graph, nullptr, 0,
&memsetParams));
void* kernelArgs1[] = {&C_d, &memsetVal, reinterpret_cast<void *>(&NElem)};
kernelNodeParams.func =
reinterpret_cast<void *>(HipTest::memsetReverse<int>);
kernelNodeParams.gridDim = dim3(blocks);
kernelNodeParams.blockDim = dim3(threadsPerBlock);
kernelNodeParams.sharedMemBytes = 0;
kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs1);
kernelNodeParams.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&memsetKer_C, graph, nullptr, 0,
&kernelNodeParams));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, C_h, C_d,
Nbytes, hipMemcpyDeviceToHost));
void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
kernelNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
kernelNodeParams.gridDim = dim3(blocks);
kernelNodeParams.blockDim = dim3(threadsPerBlock);
kernelNodeParams.sharedMemBytes = 0;
kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs2);
kernelNodeParams.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0,
&kernelNodeParams));
// Create dependencies
HIP_CHECK(hipGraphAddDependencies(graph, &memset_A, &memcpyH2D_A, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &memset_B, &memcpyH2D_B, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &kernel_vecAdd, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &kernel_vecAdd, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &memsetKer_C, &kernel_vecAdd, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2H_C, 1));
hipGraphNode_t nodes_from[EXPECTED_NUM_OF_EDGES] = {memset_A, memset_B,
memcpyH2D_A, memcpyH2D_B, memsetKer_C, kernel_vecAdd};
hipGraphNode_t nodes_to[EXPECTED_NUM_OF_EDGES] = {memcpyH2D_A, memcpyH2D_B,
kernel_vecAdd, kernel_vecAdd, kernel_vecAdd, memcpyD2H_C};
// Validate hipGraphGetEdges() API
// Scenario 1
SECTION("Validate number of edges") {
size_t numEdges = 0;
HIP_CHECK(hipGraphGetEdges(graph, nullptr, nullptr, &numEdges));
REQUIRE(numEdges == EXPECTED_NUM_OF_EDGES);
}
// Scenario 2
SECTION("Validate from/to list when numEdges = num of edges") {
validate_hipGraphGetEdges_fromto(EXPECTED_NUM_OF_EDGES, 0,
nodes_from, nodes_to, graph);
}
// Scenario 3
SECTION("Validate from/to list when numEdges = less than num of edges") {
validate_hipGraphGetEdges_fromto(EXPECTED_NUM_OF_EDGES - 1, 1,
nodes_from, nodes_to, graph);
}
// Scenario 4
SECTION("Validate from/to list when numEdges = more than num of edges") {
validate_hipGraphGetEdges_fromto(EXPECTED_NUM_OF_EDGES + 1, 2,
nodes_from, nodes_to, graph);
}
// Scenario 5
SECTION("Validate number of edges when zero or one node in graph") {
size_t numEdges = 0;
hipGraph_t graphempty;
HIP_CHECK(hipGraphCreate(&graphempty, 0));
HIP_CHECK(hipGraphGetEdges(graphempty, nullptr, nullptr, &numEdges));
REQUIRE(numEdges == 0);
// Add an empty node
hipGraphNode_t emptyNode{};
HIP_CHECK(hipGraphAddEmptyNode(&emptyNode, graphempty, nullptr, 0));
HIP_CHECK(hipGraphGetEdges(graphempty, nullptr, nullptr, &numEdges));
REQUIRE(numEdges == 0);
HIP_CHECK(hipGraphDestroy(graphempty));
}
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipGraphDestroy(graph));
}
/**
* Scenario 5: Negative Test Cases
*/
TEST_CASE("Unit_hipGraphGetEdges_Negative") {
hipGraph_t graph{}, graph_uninit{};
HIP_CHECK(hipGraphCreate(&graph, 0));
hipGraphNode_t nodes_from[EXPECTED_NUM_OF_EDGES]{},
nodes_to[EXPECTED_NUM_OF_EDGES]{};
size_t numEdges = 0;
SECTION("graph is nullptr") {
REQUIRE(hipErrorInvalidValue ==
hipGraphGetEdges(nullptr, nodes_from, nodes_to, &numEdges));
}
SECTION("from is nullptr") {
REQUIRE(hipErrorInvalidValue ==
hipGraphGetEdges(graph, nullptr, nodes_to, &numEdges));
}
SECTION("to is nullptr") {
REQUIRE(hipErrorInvalidValue ==
hipGraphGetEdges(graph, nodes_from, nullptr, &numEdges));
}
SECTION("numEdges is nullptr") {
REQUIRE(hipErrorInvalidValue ==
hipGraphGetEdges(graph, nodes_from, nodes_to, nullptr));
}
SECTION("graph is uninitialized") {
REQUIRE(hipErrorInvalidValue ==
hipGraphGetEdges(graph_uninit, nodes_from, nodes_to, &numEdges));
}
HIP_CHECK(hipGraphDestroy(graph));
}
@@ -0,0 +1,75 @@
/*
Copyright (c) 2022 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.
*/
/**
Testcase Scenarios :
Functional -
1) Create a graph and then used it for hipGraphInstantiate without adding any node to graph.
Negative -
1) Pass pGraphExec as null ptr and verify that api returns error code and doesn’t crash.
2) Pass graph as null/invalid ptr and check if api returns error.
3) Pass pGraphExec as un-initilize object and verify that api returns error code and doesn’t crash.
4) Pass Graph as un-initilize and verify that api returns error code and doesn’t crash.
*/
#include <hip_test_common.hh>
/* Test verifies hipGraphInstantiate API Negative scenarios.
*/
TEST_CASE("Unit_hipGraphInstantiate_Negative") {
hipError_t ret;
hipGraphExec_t gExec{};
hipGraph_t graph;
HIP_CHECK(hipGraphCreate(&graph, 0));
SECTION("Pass pGraphExec as nullptr") {
ret = hipGraphInstantiate(nullptr, graph, nullptr, nullptr, 0);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass graph as null/invalid ptr") {
ret = hipGraphInstantiate(&gExec, nullptr, nullptr, nullptr, 0);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass Graph as un-initialize") {
hipGraph_t graph_uninit{};
ret = hipGraphInstantiate(&gExec, graph_uninit, nullptr, nullptr, 0);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass pGraphExec as un-initialize") {
ret = hipGraphInstantiate(&gExec, graph, nullptr, nullptr, 0);
REQUIRE(hipSuccess == ret);
}
HIP_CHECK(hipGraphDestroy(graph));
}
/* Test verifies hipGraphInstantiate Basic scenarios.
Create a graph and then used it for hipGraphInstantiate without adding any node to graph.
*/
TEST_CASE("Unit_hipGraphInstantiate_Basic") {
hipGraph_t graph;
hipGraphExec_t graphExec;
HIP_CHECK(hipGraphCreate(&graph, 0));
REQUIRE(nullptr != graph);
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipGraphDestroy(graph));
}
@@ -0,0 +1,141 @@
/*
Copyright (c) 2022 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.
*/
/**
Testcase Scenarios :
Negative -
1) Pass node as nullptr and verify api returns error code.
2) Pass pNodeParams as nullptr and verify api returns error code.
Functional -
1) Create a graph, add Memset node to graph with desired node params.
Verify api fetches the node params mentioned while adding Memset node.
2) Set Memset node params with hipGraphMemsetNodeSetParams,
now get the params and verify both are same.
*/
#include <hip_test_common.hh>
/* Test verifies hipGraphMemsetNodeGetParams API Negative scenarios.
*/
TEST_CASE("Unit_hipGraphMemsetNodeGetParams_Negative") {
hipError_t ret;
hipGraph_t graph;
hipGraphNode_t memsetNode;
HIP_CHECK(hipGraphCreate(&graph, 0));
char *devData;
HIP_CHECK(hipMalloc(&devData, 1024));
hipMemsetParams memsetParams{};
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void*>(devData);
memsetParams.value = 0;
memsetParams.pitch = 0;
memsetParams.elementSize = sizeof(char);
memsetParams.width = 1024;
memsetParams.height = 1;
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
&memsetParams));
SECTION("Pass node as nullptr") {
ret = hipGraphMemsetNodeGetParams(nullptr, &memsetParams);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass GetNodeParams as nullptr") {
ret = hipGraphMemsetNodeGetParams(memsetNode, nullptr);
REQUIRE(hipErrorInvalidValue == ret);
}
HIP_CHECK(hipFree(devData));
HIP_CHECK(hipGraphDestroy(graph));
}
/* Test verifies hipGraphMemsetNodeGetParams API Functional scenarios.
*/
bool memsetNodeCompare(hipMemsetParams *mNode1, hipMemsetParams *mNode2) {
if (mNode1->dst != mNode2->dst)
return false;
if (mNode1->elementSize != mNode2->elementSize)
return false;
if (mNode1->height != mNode2->height)
return false;
if (mNode1->pitch != mNode2->pitch)
return false;
if (mNode1->value != mNode2->value)
return false;
if (mNode1->width != mNode2->width)
return false;
return true;
}
TEST_CASE("Unit_hipGraphMemsetNodeGetParams_Functional") {
constexpr size_t N = 1024;
constexpr size_t Nbytes = N * sizeof(char);
constexpr size_t val = 0;
char *devData;
HIP_CHECK(hipMalloc(&devData, Nbytes));
hipGraph_t graph;
hipGraphNode_t memsetNode;
HIP_CHECK(hipGraphCreate(&graph, 0));
hipStream_t streamForGraph;
HIP_CHECK(hipStreamCreate(&streamForGraph));
hipMemsetParams memsetParams{};
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void*>(devData);
memsetParams.value = val;
memsetParams.pitch = 0;
memsetParams.elementSize = sizeof(char);
memsetParams.width = Nbytes;
memsetParams.height = 1;
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
&memsetParams));
SECTION("Get Memset Param and verify.") {
hipMemsetParams memsetGetParams;
REQUIRE(hipSuccess == hipGraphMemsetNodeGetParams(memsetNode,
&memsetGetParams));
// Validating the result
REQUIRE(true == memsetNodeCompare(&memsetParams, &memsetGetParams));
}
SECTION("Set memset node params then Get and verify.") {
constexpr size_t updateVal = 2;
char *devData1;
HIP_CHECK(hipMalloc(&devData1, Nbytes));
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void*>(devData1);
memsetParams.value = updateVal;
memsetParams.pitch = 0;
memsetParams.elementSize = sizeof(char);
memsetParams.width = Nbytes;
memsetParams.height = 1;
hipMemsetParams memsetGetParams;
REQUIRE(hipSuccess == hipGraphMemsetNodeSetParams(memsetNode,
&memsetParams));
REQUIRE(hipSuccess == hipGraphMemsetNodeGetParams(memsetNode,
&memsetGetParams));
// Validating the result
REQUIRE(true == memsetNodeCompare(&memsetParams, &memsetGetParams));
HIP_CHECK(hipFree(devData1));
}
HIP_CHECK(hipFree(devData));
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipStreamDestroy(streamForGraph));
}
@@ -0,0 +1,185 @@
/*
Copyright (c) 2022 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.
*/
/**
Testcase Scenarios :
Negative -
1) Pass node as nullptr and verify api returns error code.
2) Pass pNodeParams as nullptr and verify api returns error code.
3) Passing hipMemsetParams::dst as nullptr should return error code.
4) Passing hipMemsetParams::element size other than 1, 2, or 4 and check.
5) Passing hipMemsetParams::height as zero and check api should return error code.
Functional -
1) Add Memset node to graph, update the node params with set and
launch the graph and check the set params are executing properly.
2) Add Memset node to graph, launch graph, then update the Memset node params
with set and launch the graph and check updated params are taking effect.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
/* Test verifies hipGraphMemsetNodeSetParams API invalid params scenarios.
*/
TEST_CASE("Unit_hipGraphMemsetNodeSetParams_InvalidParams") {
hipError_t ret;
hipGraph_t graph;
hipStream_t streamForGraph;
hipGraphNode_t memsetNode;
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipStreamCreate(&streamForGraph));
char *devData;
HIP_CHECK(hipMalloc(&devData, 1024));
hipMemsetParams memsetParams{};
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void*>(devData);
memsetParams.value = 0;
memsetParams.pitch = 0;
memsetParams.elementSize = sizeof(char);
memsetParams.width = 1024;
memsetParams.height = 1;
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
&memsetParams));
SECTION("Pass node as nullptr") {
ret = hipGraphMemsetNodeSetParams(nullptr, &memsetParams);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass GetNodeParams as nullptr") {
ret = hipGraphMemsetNodeSetParams(memsetNode, nullptr);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass dest as nullptr") {
memsetParams.dst = nullptr;
ret = hipGraphMemsetNodeSetParams(memsetNode, &memsetParams);
REQUIRE(hipErrorInvalidValue == ret);
}
#if HT_NVIDIA
SECTION("Pass element size other than 1, 2, or 4") {
memsetParams.dst = reinterpret_cast<void*>(devData);
memsetParams.elementSize = 9;
ret = hipGraphMemsetNodeSetParams(memsetNode, &memsetParams);
REQUIRE(hipErrorInvalidValue == ret);
}
SECTION("Pass height as zero or negative") {
memsetParams.elementSize = 2;
memsetParams.height = 0;
ret = hipGraphMemsetNodeSetParams(memsetNode, &memsetParams);
REQUIRE(hipErrorInvalidValue == ret);
}
#endif
HIP_CHECK(hipFree(devData));
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipStreamDestroy(streamForGraph));
}
static void validate_result(char *hData, size_t size, char val) {
// Validating the result
for (size_t i = 0; i < size; i++) {
if (hData[i] != val) {
WARN("Validation failed at- " << i << " hData[i] " << hData[i]);
REQUIRE(false);
}
}
}
/* Test verifies hipGraphMemsetNodeSetParams API Functional scenarios.
*/
TEST_CASE("Unit_hipGraphMemsetNodeSetParams_Functional") {
constexpr size_t N = 1024;
constexpr size_t Nbytes = N * sizeof(char);
constexpr size_t val = 0;
constexpr size_t updateVal = 1;
constexpr size_t updateVal2 = 2;
char *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr};
char *A_h{nullptr}, *B_h{nullptr};
HipTest::initArrays<char>(&A_d, &B_d, &C_d,
&A_h, &B_h, nullptr, N, false);
hipGraph_t graph;
hipGraphExec_t graphExec;
hipStream_t streamForGraph;
hipGraphNode_t memsetNode;
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipStreamCreate(&streamForGraph));
hipMemsetParams memsetParams{};
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void*>(C_d);
memsetParams.value = val;
memsetParams.pitch = 0;
memsetParams.elementSize = sizeof(char);
memsetParams.width = Nbytes;
memsetParams.height = 1;
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, nullptr, 0,
&memsetParams));
std::vector<hipGraphNode_t> dependencies;
dependencies.push_back(memsetNode);
SECTION("Update the memsetNode and check") {
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void*>(A_d);
memsetParams.value = updateVal;
memsetParams.pitch = 0;
memsetParams.elementSize = sizeof(char);
memsetParams.width = Nbytes;
memsetParams.height = 1;
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, dependencies.data(),
dependencies.size(), &memsetParams));
HIP_CHECK(hipGraphMemsetNodeSetParams(memsetNode, &memsetParams));
dependencies.push_back(memsetNode);
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
HIP_CHECK(hipStreamSynchronize(streamForGraph));
HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
validate_result(A_h, Nbytes, updateVal);
}
SECTION("Update the memsetNode again and check") {
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void*>(B_d);
memsetParams.value = updateVal2;
memsetParams.pitch = 0;
memsetParams.elementSize = sizeof(char);
memsetParams.width = Nbytes;
memsetParams.height = 1;
HIP_CHECK(hipGraphAddMemsetNode(&memsetNode, graph, dependencies.data(),
dependencies.size(), &memsetParams));
HIP_CHECK(hipGraphMemsetNodeSetParams(memsetNode, &memsetParams));
dependencies.push_back(memsetNode);
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
HIP_CHECK(hipStreamSynchronize(streamForGraph));
HIP_CHECK(hipMemcpy(B_h, B_d, Nbytes, hipMemcpyDeviceToHost));
validate_result(B_h, Nbytes, updateVal2);
}
HipTest::freeArrays<char>(A_d, B_d, C_d,
A_h, B_h, nullptr, false);
HIP_CHECK(hipGraphExecDestroy(graphExec));
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipStreamDestroy(streamForGraph));
}
@@ -0,0 +1,462 @@
/*
Copyright (c) 2022 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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/**
Testcase Scenarios :
1) Create a graph and add nodes with dependencies manually. Perform
selective removal of dependencies and make sure they are taking
effect using hipGraphGetEdges() API.
2) Generate graph by capturing stream. Perform selective removal of
dependencies and make sure they are taking effect using
hipGraphGetEdges() API.
3) Pass numDependencies as 0 and verify api returns success but doesn't
remove the depedencies.
4) Create a graph and add nodes with dependencies manually. Perform
selective removal of dependency and add new dependency. Verify the
change by executing the updated graph.
5) Negative Test Cases
- Pass graph parameter as nullptr.
- Pass from node parameter as nullptr.
- Pass to node parameter as nullptr.
- Pass uninitialized graph.
- Node passed in "to" parameter does not exist in graph.
- Remove non existing dependency.
- Remove the same dependency twice.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip_test_kernels.hh>
#define TOTAL_NUM_OF_EDGES 6
/**
* Kernel Functions to perform square and return in the same
* input memory location.
*/
static __global__ void vector_square(int* A_d, size_t N_ELMTS) {
size_t gputhread = (blockIdx.x * blockDim.x + threadIdx.x);
size_t stride = blockDim.x * gridDim.x;
int temp = 0;
for (size_t i = gputhread; i < N_ELMTS; i += stride) {
temp = A_d[i] * A_d[i];
A_d[i] = temp;
}
}
/**
* Scenario 1 and Scenario 3: Validate hipGraphRemoveDependencies
* for manually created graph.
*/
TEST_CASE("Unit_hipGraphRemoveDependencies_Func_Manual") {
constexpr size_t N = 1024;
constexpr size_t Nbytes = N * sizeof(int);
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
hipGraph_t graph;
hipGraphNode_t memset_A, memset_B, memsetKer_C;
hipGraphNode_t memcpyH2D_A, memcpyH2D_B, memcpyD2H_C;
hipGraphNode_t kernel_vecAdd;
hipKernelNodeParams kernelNodeParams{};
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
hipMemsetParams memsetParams{};
int memsetVal{};
size_t NElem{N};
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HIP_CHECK(hipGraphCreate(&graph, 0));
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void*>(A_d);
memsetParams.value = 0;
memsetParams.pitch = 0;
memsetParams.elementSize = sizeof(char);
memsetParams.width = Nbytes;
memsetParams.height = 1;
HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0,
&memsetParams));
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void*>(B_d);
memsetParams.value = 0;
memsetParams.pitch = 0;
memsetParams.elementSize = sizeof(char);
memsetParams.width = Nbytes;
memsetParams.height = 1;
HIP_CHECK(hipGraphAddMemsetNode(&memset_B, graph, nullptr, 0,
&memsetParams));
void* kernelArgs1[] = {&C_d, &memsetVal, reinterpret_cast<void *>(&NElem)};
kernelNodeParams.func =
reinterpret_cast<void *>(HipTest::memsetReverse<int>);
kernelNodeParams.gridDim = dim3(blocks);
kernelNodeParams.blockDim = dim3(threadsPerBlock);
kernelNodeParams.sharedMemBytes = 0;
kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs1);
kernelNodeParams.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&memsetKer_C, graph, nullptr, 0,
&kernelNodeParams));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, C_h, C_d,
Nbytes, hipMemcpyDeviceToHost));
void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
kernelNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
kernelNodeParams.gridDim = dim3(blocks);
kernelNodeParams.blockDim = dim3(threadsPerBlock);
kernelNodeParams.sharedMemBytes = 0;
kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs2);
kernelNodeParams.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0,
&kernelNodeParams));
// Create dependencies
HIP_CHECK(hipGraphAddDependencies(graph, &memset_A, &memcpyH2D_A, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &memset_B, &memcpyH2D_B, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &kernel_vecAdd, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &kernel_vecAdd, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &memsetKer_C, &kernel_vecAdd, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2H_C, 1));
SECTION("scenario 1") {
// Remove some dependencies
constexpr size_t numEdgesRemoved = 3;
HIP_CHECK(hipGraphRemoveDependencies(graph, &memcpyH2D_A,
&kernel_vecAdd, 1));
HIP_CHECK(hipGraphRemoveDependencies(graph, &memcpyH2D_B,
&kernel_vecAdd, 1));
HIP_CHECK(hipGraphRemoveDependencies(graph, &memsetKer_C,
&kernel_vecAdd, 1));
// Validate manually with hipGraphGetEdges() API
hipGraphNode_t fromnode[TOTAL_NUM_OF_EDGES]{};
hipGraphNode_t tonode[TOTAL_NUM_OF_EDGES]{};
size_t numEdges = TOTAL_NUM_OF_EDGES;
HIP_CHECK(hipGraphGetEdges(graph, fromnode, tonode, &numEdges));
hipGraphNode_t expected_from_nodes[numEdgesRemoved] = {memcpyH2D_A,
memcpyH2D_B, memsetKer_C};
hipGraphNode_t expected_to_nodes[numEdgesRemoved] = {kernel_vecAdd,
kernel_vecAdd, kernel_vecAdd};
bool nodeFound;
int found_count = 0;
for (size_t idx_from = 0; idx_from < numEdgesRemoved; idx_from++) {
nodeFound = false;
int idx = 0;
for (; idx < TOTAL_NUM_OF_EDGES; idx++) {
if (expected_from_nodes[idx_from] == fromnode[idx]) {
nodeFound = true;
break;
}
}
if (nodeFound && (tonode[idx] == expected_to_nodes[idx_from])) {
found_count++;
}
}
// Ensure none of the nodes are discovered
REQUIRE(0 == found_count);
// Validate with returned number of edges from hipGraphGetEdges() API
numEdges = 0;
HIP_CHECK(hipGraphGetEdges(graph, nullptr, nullptr, &numEdges));
size_t numEdgesExpected = TOTAL_NUM_OF_EDGES - numEdgesRemoved;
REQUIRE(numEdgesExpected == numEdges);
}
SECTION("scenario 3") {
HIP_CHECK(hipGraphRemoveDependencies(graph, &memcpyH2D_A,
&kernel_vecAdd, 0));
size_t numEdges = 0;
HIP_CHECK(hipGraphGetEdges(graph, nullptr, nullptr, &numEdges));
size_t numEdgesExpected = TOTAL_NUM_OF_EDGES;
REQUIRE(numEdgesExpected == numEdges);
}
// Destroy
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipGraphDestroy(graph));
}
/**
* Scenario 2: Validate hipGraphRemoveDependencies for stream captured graph.
*/
TEST_CASE("Unit_hipGraphRemoveDependencies_Func_StrmCapture") {
hipStream_t stream1, stream2, stream3;
hipEvent_t forkStreamEvent, memsetEvent1, memsetEvent2;
hipGraph_t graph;
constexpr size_t N = 1024;
constexpr size_t Nbytes = N * sizeof(int);
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
size_t NElem{N};
int memsetVal{};
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
// Create streams and events
HIP_CHECK(hipStreamCreate(&stream1));
HIP_CHECK(hipStreamCreate(&stream2));
HIP_CHECK(hipStreamCreate(&stream3));
HIP_CHECK(hipEventCreate(&forkStreamEvent));
HIP_CHECK(hipEventCreate(&memsetEvent1));
HIP_CHECK(hipEventCreate(&memsetEvent2));
// Begin stream capture
HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal));
HIP_CHECK(hipEventRecord(forkStreamEvent, stream1));
HIP_CHECK(hipStreamWaitEvent(stream2, forkStreamEvent, 0));
HIP_CHECK(hipStreamWaitEvent(stream3, forkStreamEvent, 0));
// Add operations to stream3
hipLaunchKernelGGL(HipTest::memsetReverse<int>,
dim3(blocks), dim3(threadsPerBlock), 0, stream3,
C_d, memsetVal, NElem);
HIP_CHECK(hipEventRecord(memsetEvent1, stream3));
// Add operations to stream2
HIP_CHECK(hipMemsetAsync(B_d, 0, Nbytes, stream2));
HIP_CHECK(hipMemcpyAsync(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream2));
HIP_CHECK(hipEventRecord(memsetEvent2, stream2));
// Add operations to stream1
HIP_CHECK(hipMemsetAsync(A_d, 0, Nbytes, stream1));
HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream1));
HIP_CHECK(hipStreamWaitEvent(stream1, memsetEvent2, 0));
HIP_CHECK(hipStreamWaitEvent(stream1, memsetEvent1, 0));
hipLaunchKernelGGL(HipTest::vectorADD<int>,
dim3(blocks), dim3(threadsPerBlock), 0, stream1,
A_d, B_d, C_d, NElem);
HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost,
stream1));
HIP_CHECK(hipStreamEndCapture(stream1, &graph));
hipGraphNode_t* nodes{nullptr};
size_t numNodes = 0, numEdges = 0;
HIP_CHECK(hipGraphGetNodes(graph, nodes, &numNodes));
HIP_CHECK(hipGraphGetEdges(graph, nullptr, nullptr, &numEdges));
REQUIRE(7 == numNodes);
REQUIRE(TOTAL_NUM_OF_EDGES == numEdges);
// Get the edges and remove one edge. Verify edge is removed.
hipGraphNode_t fromnode[TOTAL_NUM_OF_EDGES]{};
hipGraphNode_t tonode[TOTAL_NUM_OF_EDGES]{};
HIP_CHECK(hipGraphGetEdges(graph, fromnode, tonode, &numEdges));
HIP_CHECK(hipGraphRemoveDependencies(graph, &fromnode[0],
&tonode[0], 1));
// Verify
HIP_CHECK(hipGraphGetEdges(graph, nullptr, nullptr, &numEdges));
size_t expected_num_edges = TOTAL_NUM_OF_EDGES - 1;
REQUIRE(expected_num_edges == numEdges);
// Destroy
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipStreamDestroy(stream1));
HIP_CHECK(hipStreamDestroy(stream2));
HIP_CHECK(hipStreamDestroy(stream3));
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
}
/**
* Scenario 4: Dynamically modify dependencies in a graph using
* hipGraphRemoveDependencies and verify the computation.
*/
TEST_CASE("Unit_hipGraphRemoveDependencies_ChangeComputeFunc") {
hipStream_t streamForGraph;
HIP_CHECK(hipStreamCreate(&streamForGraph));
constexpr size_t N = 1024;
constexpr size_t Nbytes = N * sizeof(int);
constexpr auto blocksPerCU = 6; // to hide latency
constexpr auto threadsPerBlock = 256;
hipGraph_t graph;
hipGraphNode_t memcpyH2D_A, memcpyH2D_B, memcpyD2H_C;
hipGraphNode_t kernel_vecAdd, kernel_square;
hipKernelNodeParams kernelNodeParams{};
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
size_t NElem{N};
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HIP_CHECK(hipGraphCreate(&graph, 0));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h,
Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, C_h, C_d,
Nbytes, hipMemcpyDeviceToHost));
void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast<void *>(&NElem)};
kernelNodeParams.func = reinterpret_cast<void *>(HipTest::vectorADD<int>);
kernelNodeParams.gridDim = dim3(blocks);
kernelNodeParams.blockDim = dim3(threadsPerBlock);
kernelNodeParams.sharedMemBytes = 0;
kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs2);
kernelNodeParams.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0,
&kernelNodeParams));
// Create dependencies
HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &kernel_vecAdd, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &kernel_vecAdd, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2H_C, 1));
// Instantiate and execute Graph
hipGraphExec_t graphExec;
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
HIP_CHECK(hipStreamSynchronize(streamForGraph));
// Validate
bool bMismatch = false;
for (size_t idx = 0; idx < NElem; idx++) {
if (C_h[idx] != (A_h[idx] + B_h[idx])) {
bMismatch = true;
break;
}
}
REQUIRE(false == bMismatch);
HIP_CHECK(hipGraphExecDestroy(graphExec));
// Remove dependency memcpyH2D_B -> kernel_vecAdd and
// add new dependencies memcpyH2D_B -> kernel_square -> kernel_vecAdd
// Square kernel
void* kernelArgs1[] = {&B_d, reinterpret_cast<void *>(&NElem)};
kernelNodeParams.func =
reinterpret_cast<void *>(vector_square);
kernelNodeParams.gridDim = dim3(blocks);
kernelNodeParams.blockDim = dim3(threadsPerBlock);
kernelNodeParams.sharedMemBytes = 0;
kernelNodeParams.kernelParams = reinterpret_cast<void**>(kernelArgs1);
kernelNodeParams.extra = nullptr;
HIP_CHECK(hipGraphAddKernelNode(&kernel_square, graph, nullptr, 0,
&kernelNodeParams));
hipGraphRemoveDependencies(graph, &memcpyH2D_B, &kernel_vecAdd, 1);
// Add new dependencies
HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &kernel_square, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &kernel_square,
&kernel_vecAdd, 1));
size_t numEdges = 0, numNodes = 0;
HIP_CHECK(hipGraphGetEdges(graph, nullptr, nullptr, &numEdges));
REQUIRE(4 == numEdges);
HIP_CHECK(hipGraphGetNodes(graph, nullptr, &numNodes));
REQUIRE(5 == numNodes);
// Instantiate and execute graph
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph));
HIP_CHECK(hipStreamSynchronize(streamForGraph));
// Validate
bMismatch = false;
for (size_t idx = 0; idx < NElem; idx++) {
if (C_h[idx] != (A_h[idx] + B_h[idx]*B_h[idx])) {
bMismatch = true;
break;
}
}
REQUIRE(false == bMismatch);
// Destroy
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipStreamDestroy(streamForGraph));
}
/**
* Scenario 5: Negative Tests
*/
TEST_CASE("Unit_hipGraphRemoveDependencies_Negative") {
hipGraph_t graph{};
HIP_CHECK(hipGraphCreate(&graph, 0));
hipEvent_t event_start, event_end;
HIP_CHECK(hipEventCreateWithFlags(&event_start, hipEventDisableTiming));
HIP_CHECK(hipEventCreateWithFlags(&event_end, hipEventDisableTiming));
// memset node
constexpr size_t Nbytes = 1024;
char *A_d;
hipGraphNode_t memset_A;
hipMemsetParams memsetParams{};
HIP_CHECK(hipMalloc(&A_d, Nbytes));
memset(&memsetParams, 0, sizeof(memsetParams));
memsetParams.dst = reinterpret_cast<void*>(A_d);
memsetParams.value = 0;
memsetParams.pitch = 0;
memsetParams.elementSize = sizeof(char);
memsetParams.width = Nbytes;
memsetParams.height = 1;
HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0,
&memsetParams));
// create event record node
hipGraphNode_t event_node_start, event_node_end;
HIP_CHECK(hipGraphAddEventRecordNode(&event_node_start, graph, nullptr, 0,
event_start));
HIP_CHECK(hipGraphAddEventRecordNode(&event_node_end, graph, nullptr, 0,
event_end));
// create empty node
hipGraphNode_t emptyNode{};
HIP_CHECK(hipGraphAddEmptyNode(&emptyNode, graph, nullptr, 0));
// Add dependencies between nodes
HIP_CHECK(hipGraphAddDependencies(graph, &event_node_start, &memset_A, 1));
HIP_CHECK(hipGraphAddDependencies(graph, &memset_A, &event_node_end, 1));
SECTION("graph is nullptr") {
REQUIRE(hipErrorInvalidValue ==
hipGraphRemoveDependencies(nullptr, &event_node_start, &memset_A, 1));
}
SECTION("from is nullptr") {
REQUIRE(hipErrorInvalidValue ==
hipGraphRemoveDependencies(graph, nullptr, &memset_A, 1));
}
SECTION("to is nullptr") {
REQUIRE(hipErrorInvalidValue ==
hipGraphRemoveDependencies(graph, &event_node_start, nullptr, 1));
}
SECTION("graph is uninitialized") {
hipGraph_t graph_uninit{};
REQUIRE(hipErrorInvalidValue ==
hipGraphRemoveDependencies(graph_uninit, &event_node_start,
nullptr, 1));
}
SECTION("non existing node") {
REQUIRE(hipErrorInvalidValue ==
hipGraphRemoveDependencies(graph, &event_node_start,
&emptyNode, 1));
}
SECTION("remove non existing dependency") {
REQUIRE(hipErrorInvalidValue ==
hipGraphRemoveDependencies(graph, &event_node_start,
&event_node_end, 1));
}
SECTION("remove same dependency twice") {
HIP_CHECK(hipGraphRemoveDependencies(graph, &event_node_start,
&memset_A, 1));
REQUIRE(hipErrorInvalidValue ==
hipGraphRemoveDependencies(graph, &event_node_start,
&memset_A, 1));
}
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipGraphDestroy(graph));
HIP_CHECK(hipEventDestroy(event_end));
HIP_CHECK(hipEventDestroy(event_start));
}
+29
Ver fichero
@@ -0,0 +1,29 @@
# Copyright (c) 2022 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.
# Common Tests - Test independent of all platforms
set(TEST_SRC
hipMemFaultStackAllocation.cc
hipLaunchBounds.cc
)
hip_add_exe_to_target(NAME KernelTest
TEST_SRC ${TEST_SRC}
TEST_TARGET_NAME build_tests)
+173
Ver fichero
@@ -0,0 +1,173 @@
/*
Copyright (c) 2022 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.
*/
/**
Testcase Scenarios : hipLaunchBounds_With_maxThreadsPerBlock
1) Passing threadsPerBlock same as kernel launch_bounds.
2) Passing threadsPerBlock less than kernel launch_bounds.
3) Passing threadsPerBlock more than kernel launch_bounds.
4) Passing threadsPerBlock as 0 to kernel launch_bounds.
Testcase Scenarios : hipLaunchBounds_With_maxThreadsPerBlock_blocksPerCU
1) Passing threadsPerBlock same as kernel launch_bounds.
2) Passing threadsPerBlock less than kernel launch_bounds.
3) Passing threadsPerBlock more than kernel launch_bounds.
4) Passing threadsPerBlock as 0 to kernel launch_bounds.
5) Passing blocksPerCU same as kernel launch_bounds.
6) Passing blocksPerCU less than kernel launch_bounds.
7) Passing blocksPerCU more than kernel launch_bounds.
8) Passing blocksPerCU as 0 to kernel launch_bounds.
*/
#include <hip_test_common.hh>
#include <hip_test_kernels.hh>
__global__ void
__launch_bounds__(128, 2)
MyKernel(int N, int *x, int val) {
for (int i = 0; i < N; i++) {
x[i] = val;
}
}
__global__ void
__launch_bounds__(64)
MyKernel_2(int N, int *x, int val) {
for (int i = 0; i < N; i++) {
x[i] = val;
}
}
static bool verify(int N, int *x, int val) {
for (int i = 0; i < N; i++) {
if (x[i] != val) {
return false;
}
}
return true;
}
TEST_CASE("Unit_hipLaunchBounds_With_maxThreadsPerBlock_Check") {
constexpr size_t N = 10000;
hipError_t ret;
int *x;
HIP_CHECK(hipMallocManaged(&x, N*sizeof(int)));
REQUIRE(x != nullptr);
SECTION("Passing threadsPerBlock same as kernel launch_bounds") {
hipLaunchKernelGGL(MyKernel_2, dim3(4), dim3(64), 0, 0, N, x, 2);
ret = hipGetLastError();
REQUIRE(hipSuccess == ret);
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(true == verify(N, x, 2));
}
SECTION("Passing threadsPerBlock less than kernel launch_bounds") {
hipLaunchKernelGGL(MyKernel_2, dim3(4), dim3(32), 0, 0, N, x, 22);
ret = hipGetLastError();
REQUIRE(hipSuccess == ret);
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(true == verify(N, x, 22));
}
SECTION("Passing threadsPerBlock more than kernel launch_bounds") {
hipLaunchKernelGGL(MyKernel_2, dim3(4), dim3(128), 0, 0, N, x, 9);
ret = hipGetLastError();
REQUIRE(hipSuccess != ret);
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(true != verify(N, x, 9));
}
SECTION("Passing threadsPerBlock as 0 to kernel launch_bounds") {
hipLaunchKernelGGL(MyKernel_2, dim3(4), dim3(0), 0, 0, N, x, 19);
ret = hipGetLastError();
REQUIRE(hipSuccess != ret);
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(true != verify(N, x, 19));
}
HIP_CHECK(hipFree(x));
}
TEST_CASE("Unit_hipLaunchBounds_With_maxThreadsPerBlock_blocksPerCU_Check") {
constexpr size_t N = 10000;
hipError_t ret;
int *x;
HIP_CHECK(hipMallocManaged(&x, N*sizeof(int)));
REQUIRE(x != nullptr);
SECTION("Passing threadsPerBlock same as kernel launch_bounds") {
hipLaunchKernelGGL(MyKernel, dim3(1), dim3(128), 0, 0, N, x, 1);
ret = hipGetLastError();
REQUIRE(hipSuccess == ret);
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(true == verify(N, x, 1));
}
SECTION("Passing threadsPerBlock less than kernel launch_bounds") {
hipLaunchKernelGGL(MyKernel, dim3(2), dim3(64), 0, 0, N, x, 11);
ret = hipGetLastError();
REQUIRE(hipSuccess == ret);
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(true == verify(N, x, 11));
}
SECTION("Passing threadsPerBlock more than kernel launch_bounds") {
hipLaunchKernelGGL(MyKernel, dim3(2), dim3(256), 0, 0, N, x, 3);
ret = hipGetLastError();
REQUIRE(hipSuccess != ret);
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(true != verify(N, x, 3));
}
SECTION("Passing threadsPerBlock as 0 to kernel launch_bounds") {
hipLaunchKernelGGL(MyKernel, dim3(2), dim3(0), 0, 0, N, x, 13);
ret = hipGetLastError();
REQUIRE(hipSuccess != ret);
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(true != verify(N, x, 13));
}
SECTION("Passing blocksPerCU same as kernel launch_bounds") {
hipLaunchKernelGGL(MyKernel, dim3(2), dim3(128), 0, 0, N, x, 5);
ret = hipGetLastError();
REQUIRE(hipSuccess == ret);
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(true == verify(N, x, 5));
}
SECTION("Passing blocksPerCU less than kernel launch_bounds") {
hipLaunchKernelGGL(MyKernel, dim3(1), dim3(128), 0, 0, N, x, 25);
ret = hipGetLastError();
REQUIRE(hipSuccess == ret);
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(true == verify(N, x, 25));
}
SECTION("Passing blocksPerCU more than kernel launch_bounds") {
hipLaunchKernelGGL(MyKernel, dim3(4), dim3(128), 0, 0, N, x, 7);
ret = hipGetLastError();
REQUIRE(hipSuccess == ret);
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(true == verify(N, x, 7));
}
SECTION("Passing blocksPerCU as 0 to kernel launch_bounds") {
hipLaunchKernelGGL(MyKernel, dim3(0), dim3(128), 0, 0, N, x, 37);
ret = hipGetLastError();
REQUIRE(hipSuccess != ret);
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(true != verify(N, x, 37));
}
HIP_CHECK(hipFree(x));
}
@@ -0,0 +1,102 @@
/*
Copyright (c) 2022 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.
*/
/**
Testcase Scenarios :
1) Calling Kernel which allocate ConstSize to local array.
2) Calling Kernel which allocate VariableSize to local array.
*/
#include <hip_test_common.hh>
#include <hip_test_kernels.hh>
const size_t N = 100000;
__global__ void MyKernelConstSize(int* C_d, const int* A_d) {
constexpr size_t A1size = 1024;
int A1[A1size];
for (size_t i = 0; i < A1size; ++i) {
A1[i] = i;
}
for (size_t i = 0; i < N; ++i) {
C_d[i] = A_d[i] + A1[i%A1size];
}
}
__global__ void MyKernelVariableSize(int* C_d, const int* A_d) {
constexpr size_t A1size = 1024;
int A1[1024];
for (size_t i = 0; i < A1size; ++i) {
A1[i] = i;
}
for (size_t i = 0; i < N; ++i) {
C_d[i] = A_d[i] + A1[i%A1size];
}
}
static bool verify(const int* C_d, const int* A_d) {
for (size_t i = 0; i < N; i++) {
if (C_d[i] != A_d[i] + i%1024) {
return false;
}
}
return true;
}
TEST_CASE("Unit_hipMemFaultStackAllocation_Check") {
hipError_t ret;
int *A_d, *C_d;
const size_t Nbytes = N * sizeof(int);
const unsigned threadsPerBlock = 256;
const unsigned blocks = (N + threadsPerBlock - 1)/threadsPerBlock;
HIP_CHECK(hipMallocManaged(&A_d, Nbytes));
REQUIRE(A_d != nullptr);
HIP_CHECK(hipMallocManaged(&C_d, Nbytes));
REQUIRE(C_d != nullptr);
for (size_t i = 0; i < N; i++) {
A_d[i] = i%1024;
}
SECTION("Calling Kernel which allocate ConstSize to local array") {
hipLaunchKernelGGL(MyKernelConstSize, dim3(blocks),
dim3(threadsPerBlock), 0, 0, C_d, A_d);
ret = hipGetLastError();
REQUIRE(hipSuccess == ret);
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(true == verify(C_d, A_d));
}
SECTION("Calling Kernel which allocate VariableSize to local array") {
hipLaunchKernelGGL(MyKernelVariableSize, dim3(blocks),
dim3(threadsPerBlock), 0, 0, C_d, A_d);
ret = hipGetLastError();
REQUIRE(hipSuccess == ret);
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(true == verify(C_d, A_d));
}
HIP_CHECK(hipFree(C_d));
HIP_CHECK(hipFree(A_d));
}
@@ -189,6 +189,7 @@ TEST_CASE("Unit_hipMemsetAsync_VerifyExecutionWithKernel") {
UseStrmPerThrd = true;
ret = testhipMemsetAsyncWithKernel(UseStrmPerThrd);
REQUIRE(ret == true);
HIP_CHECK(hipDeviceReset());
}
SECTION("hipMemsetD32Async With Kernel") {
+10 -2
Ver fichero
@@ -1,6 +1,12 @@
# AMD Tests
# Common Tests - Test independent of all platforms
set(TEST_SRC
saxpy.cc
warpsize.cc
)
# AMD only tests
set(AMD_TEST_SRC
customOptions.cc
)
if(HIP_PLATFORM MATCHES "nvidia")
@@ -9,7 +15,9 @@ if(HIP_PLATFORM MATCHES "nvidia")
TEST_TARGET_NAME build_tests
LINKER_LIBS nvrtc)
elseif(HIP_PLATFORM MATCHES "amd")
set(TEST_SRC ${TEST_SRC} ${AMD_TEST_SRC})
hip_add_exe_to_target(NAME RTC
TEST_SRC ${TEST_SRC}
TEST_TARGET_NAME build_tests)
TEST_TARGET_NAME build_tests
LINKER_LIBS hiprtc)
endif()
+98
Ver fichero
@@ -0,0 +1,98 @@
/*
Copyright (c) 2022 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.
*/
#include <hip_test_common.hh>
#include <hip/hiprtc.h>
#include <hip/hip_runtime.h>
#include <cassert>
#include <cstddef>
#include <memory>
#include <iostream>
#include <iterator>
#include <vector>
const char* funcname = "getWarpSize";
static constexpr auto code{
R"(
extern "C"
__global__
void getWarpSize(int* warpSizePtr)
{
if (threadIdx.x == 0 && blockIdx.x == 0) *warpSizePtr = warpSize;
}
)"};
TEST_CASE("Unit_hiprtc_warpsize") {
using namespace std;
hiprtcProgram prog;
HIPRTC_CHECK(hiprtcCreateProgram(&prog, code, "code.cu", 0, nullptr, nullptr));
hipDeviceProp_t props;
int device = 0;
hipGetDeviceProperties(&props, device);
#ifdef __HIP_PLATFORM_AMD__
std::string sarg = std::string("--gpu-architecture=") + props.gcnArchName;
#else
std::string sarg = std::string("--gpu-architecture=compute_")
+ std::to_string(props.major) + std::to_string(props.minor);
#endif
vector<const char*> opts;
opts.push_back(sarg.c_str());
hiprtcResult compileResult{hiprtcCompileProgram(prog, opts.size(), opts.data())};
size_t logSize;
HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize));
if (logSize) {
string log(logSize, '\0');
HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0]));
std::cout << log << '\n';
}
REQUIRE(compileResult == HIPRTC_SUCCESS);
size_t codeSize;
HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize));
vector<char> codec(codeSize);
HIPRTC_CHECK(hiprtcGetCode(prog, codec.data()));
HIPRTC_CHECK(hiprtcDestroyProgram(&prog));
int* d_warpSize;
HIP_CHECK(hipMalloc(&d_warpSize, sizeof(int)));
hipModule_t module;
hipFunction_t function;
HIP_CHECK(hipModuleLoadData(&module, codec.data()));
HIP_CHECK(hipModuleGetFunction(&function, module, funcname));
void* args[] = { &d_warpSize };
HIP_CHECK(hipModuleLaunchKernel(function, 1, 1, 1, 64, 1, 1, 0, 0, args, 0));
HIP_CHECK(hipDeviceSynchronize());
int h_warpSize;
HIP_CHECK(hipMemcpyDtoH(&h_warpSize, reinterpret_cast<hipDeviceptr_t>(d_warpSize), sizeof(int)));
HIP_CHECK(hipFree(d_warpSize));
HIP_CHECK(hipModuleUnload(module));
// Verifies warp size returned by the kernel via hiprtc and runtime to be same
REQUIRE(h_warpSize == props.warpSize);
}
+2
Ver fichero
@@ -9,6 +9,7 @@ set(TEST_SRC
hipStreamCreateWithPriority.cc
hipStreamGetCUMask.cc
hipAPIStreamDisable.cc
streamCommon.cc
)
#skipped in windows - duplicate HipTest::vector_square sym (compiler issue)
@@ -29,6 +30,7 @@ set(TEST_SRC
hipStreamCreateWithFlags.cc
hipStreamCreateWithPriority.cc
hipAPIStreamDisable.cc
streamCommon.cc
)
endif()
+14 -30
Ver fichero
@@ -16,36 +16,20 @@ 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.
*/
#include <hip_test_common.hh>
TEST_CASE("Unit_hipStreamCreate_default") {
hipStream_t stream;
HIP_CHECK(hipStreamCreate(&stream));
}
TEST_CASE("Unit_hipStreamCreateWithFlags_Negative") {
hipStream_t stream;
auto status = hipStreamCreateWithFlags(&stream, 0xFF);
REQUIRE(status == hipErrorInvalidValue);
status = hipStreamCreateWithFlags(nullptr, hipStreamDefault);
REQUIRE(status == hipErrorInvalidValue);
}
TEST_CASE("Unit_hipStreamCreateWithFlags") {
hipStream_t stream;
HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamDefault));
HIP_CHECK(hipStreamDestroy(stream));
HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking));
HIP_CHECK(hipStreamDestroy(stream));
}
TEST_CASE("Unit_hipStreamCreateWithPriority") {
int priority_low = 0;
int priority_high = 0;
HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
hipStream_t stream;
SECTION("Setting high prirority") {
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_high));
}
SECTION("Setting low priority") {
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_low));
}
#include "streamCommon.hh"
TEST_CASE("Unit_hipStreamCreate_default") {
int id = GENERATE(range(0, HipTest::getDeviceCount()));
HIP_CHECK(hipSetDevice(id));
hipStream_t stream{nullptr};
HIP_CHECK(hipStreamCreate(&stream));
REQUIRE(stream != nullptr); // Check if stream has a valid ptr
REQUIRE(hip::checkStream(stream)); // check its flags and priority
HIP_CHECK(hipStreamDestroy(stream));
}
TEST_CASE("Unit_hipStreamCreate_Negative") {
REQUIRE(hipErrorInvalidValue == hipStreamCreate(nullptr));
}
@@ -1,5 +1,5 @@
/*
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2022 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
@@ -43,7 +43,7 @@ kernel tasks on these streams from multiple threads. Validate all the results.
8) Validate stream priorities with event after classifying them as low, medium, high.
*/
#include <hip_test_common.hh>
#include "streamCommon.hh"
#include <hip_test_kernels.hh>
#include <atomic>
#include <vector>
@@ -500,8 +500,6 @@ bool validateStreamPrioritiesWithEvents() {
} // namespace hipStreamCreateWithPriorityTest
/**
Tests following scenarios.
1)Create streams with default flag for all available priority levels and
@@ -565,14 +563,12 @@ TEST_CASE("Unit_hipStreamCreateWithPriority_MulthreadNonblockingflag") {
flag = 0xffffffff.
*/
TEST_CASE("Unit_hipStreamCreateWithPriority_NegTst") {
hipStream_t stream;
int priority_low;
int priority_high;
hipError_t ret;
hipStream_t stream{nullptr};
int priority_low{0};
int priority_high{0};
// Test is to get the Stream Priority Range
HIP_CHECK(
hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
// Check if priorities are indeed supported
if (priority_low == priority_high) {
WARN("Stream priority range not supported. Skipping test.");
@@ -580,26 +576,64 @@ TEST_CASE("Unit_hipStreamCreateWithPriority_NegTst") {
}
SECTION("stream = nullptr test") {
ret = hipStreamCreateWithPriority(nullptr, hipStreamDefault,
priority_low);
REQUIRE(ret != hipSuccess);
REQUIRE(hipErrorInvalidValue ==
hipStreamCreateWithPriority(nullptr, hipStreamDefault, priority_low));
}
SECTION("flag value invalid test") {
ret = hipStreamCreateWithPriority(&stream, 0xffffffff,
priority_low);
REQUIRE(ret != hipSuccess);
REQUIRE(hipErrorInvalidValue == hipStreamCreateWithPriority(&stream, 0xffffffff, priority_low));
}
}
TEST_CASE("Unit_hipStreamCreateWithPriority") {
int id = GENERATE(range(0, HipTest::getDeviceCount()));
HIP_CHECK(hipSetDevice(id));
int priority_low = 0, priority_high = 0;
HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
hipStream_t stream{nullptr};
SECTION("Setting high priority") {
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_high));
REQUIRE(stream != nullptr);
REQUIRE(hip::checkStreamPriorityAndFlags(stream, priority_high));
}
SECTION("Setting low priority") {
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault, priority_low));
REQUIRE(stream != nullptr);
REQUIRE(hip::checkStreamPriorityAndFlags(stream, priority_low));
}
SECTION("Setting lowest possible priority") {
HIP_CHECK(
hipStreamCreateWithPriority(&stream, hipStreamDefault, std::numeric_limits<int>::max()));
REQUIRE(stream != nullptr);
REQUIRE(hip::checkStreamPriorityAndFlags(stream, priority_low));
}
SECTION("Setting highest possible priority") {
HIP_CHECK(
hipStreamCreateWithPriority(&stream, hipStreamDefault, std::numeric_limits<int>::min()));
REQUIRE(stream != nullptr);
REQUIRE(hip::checkStreamPriorityAndFlags(stream, priority_high));
}
SECTION("Setting flags to hipStreamNonBlocking") {
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamNonBlocking, priority_high));
REQUIRE(stream != nullptr);
REQUIRE(hip::checkStreamPriorityAndFlags(stream, priority_high, hipStreamNonBlocking));
}
HIP_CHECK(hipStreamDestroy(stream));
}
/**
* Validate stream priorities with event after classifying them as low, medium and high.
*/
TEST_CASE("Unit_hipStreamCreateWithPriority_ValidateWithEvents") {
bool TestPassed = true;
TestPassed = hipStreamCreateWithPriorityTest::
validateStreamPrioritiesWithEvents<int>();
TestPassed = hipStreamCreateWithPriorityTest::validateStreamPrioritiesWithEvents<int>();
REQUIRE(TestPassed);
}
@@ -35,44 +35,6 @@ TEST_CASE("Unit_hipStreamGetPriority_Negative") {
REQUIRE(hipStreamGetPriority(stream, nullptr) == hipErrorInvalidValue);
}
/**
* Create stream and check default priority of stream is within range.
*/
TEST_CASE("Unit_hipStreamGetPriority_default") {
int priority_low = 0;
int priority_high = 0;
int devID = GENERATE(range(0, HipTest::getDeviceCount()));
HIP_CHECK(hipSetDevice(devID));
HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
hipStream_t stream;
HIP_CHECK(hipStreamCreate(&stream));
int priority = 0;
HIP_CHECK(hipStreamGetPriority(stream, &priority));
// valid priority
// Lower the value higher the priority, higher the value lower the priority
REQUIRE(priority_low >= priority);
REQUIRE(priority >= priority_high);
HIP_CHECK(hipStreamDestroy(stream));
}
/**
* Create stream with high priority and check priority is set as expected.
*/
TEST_CASE("Unit_hipStreamGetPriority_high") {
int priority_low = 0;
int priority_high = 0;
int devID = GENERATE(range(0, HipTest::getDeviceCount()));
HIP_CHECK(hipSetDevice(devID));
HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
hipStream_t stream;
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault,
priority_high));
int priority = 0;
HIP_CHECK(hipStreamGetPriority(stream, &priority));
REQUIRE(priority == priority_high);
HIP_CHECK(hipStreamDestroy(stream));
}
/**
* Create stream with higher priority for the priority range returned.
*/
@@ -91,24 +53,6 @@ TEST_CASE("Unit_hipStreamGetPriority_higher") {
HIP_CHECK(hipStreamDestroy(stream));
}
/**
* Create stream with low priority and check priority is set as expected.
*/
TEST_CASE("Unit_hipStreamGetPriority_low") {
int priority_low = 0;
int priority_high = 0;
int devID = GENERATE(range(0, HipTest::getDeviceCount()));
HIP_CHECK(hipSetDevice(devID));
HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
hipStream_t stream;
HIP_CHECK(hipStreamCreateWithPriority(&stream, hipStreamDefault,
priority_low));
int priority = 0;
HIP_CHECK(hipStreamGetPriority(stream, &priority));
REQUIRE(priority_low == priority);
HIP_CHECK(hipStreamDestroy(stream));
}
/**
* Create stream with lower priority for the priority range returned.
*/
+137
Ver fichero
@@ -0,0 +1,137 @@
/*
Copyright (c) 2022 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.
*/
#include "streamCommon.hh"
namespace hip {
inline namespace internal {
bool checkStreamPriority_(hipStream_t stream, bool checkPriority = false, int priority_ = 0) {
int priority{0};
HIP_CHECK(hipStreamGetPriority(stream, &priority));
if (checkPriority) {
if (priority_ != priority) {
UNSCOPED_INFO("Priority Mismatch, Expected Priority: " << priority_
<< " Actual Priority: " << priority);
return false;
}
} else {
int priority_low{0}, priority_high{0};
HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
if (priority_low < priority || priority_high > priority) {
UNSCOPED_INFO("Priority Mismatch, Expected Priority Range: "
<< priority_low << " - " << priority_high << " Actual Priority: " << priority);
return false;
}
}
return true;
}
bool checkStreamFlags_(hipStream_t stream, bool checkFlags = false, unsigned flags_ = 0) {
unsigned flags{0};
HIP_CHECK(hipStreamGetFlags(stream, &flags));
if (checkFlags) {
if (flags_ != flags) {
UNSCOPED_INFO("Flags Mismatch, Expected Flag: " << flags_ << " Actual Flag: " << flags);
return false;
}
} else {
if (flags != hipStreamDefault && flags != hipStreamNonBlocking) {
UNSCOPED_INFO("Flags Mismatch, Expected Flag: " << hipStreamDefault << " or "
<< hipStreamNonBlocking
<< " Actual Flag: " << flags);
return false;
}
}
return true;
}
} // namespace internal
inline namespace stream {
__device__ int defaultSemaphore = 0;
__global__ void signaling_kernel(int* semaphore) {
size_t tid{blockIdx.x * blockDim.x + threadIdx.x};
if (tid == 0) {
if (semaphore == nullptr) {
atomicAdd(&defaultSemaphore, 1);
} else {
atomicAdd(semaphore, 1);
}
}
}
__global__ void waiting_kernel(int* semaphore) {
size_t tid{blockIdx.x * blockDim.x + threadIdx.x};
if (tid == 0) {
if (semaphore == nullptr) {
while (atomicCAS(&defaultSemaphore, 1, 2) == 0) {
}
} else {
while (atomicCAS(semaphore, 1, 2) == 0) {
}
}
}
}
std::thread startSignalingThread(int* semaphore) {
std::thread signalingThread([semaphore]() {
hipStream_t signalingStream;
HIP_CHECK(hipStreamCreateWithFlags(&signalingStream, hipStreamNonBlocking));
signaling_kernel<<<1, 1, 0, signalingStream>>>(semaphore);
HIP_CHECK(hipStreamSynchronize(signalingStream));
HIP_CHECK(hipStreamDestroy(signalingStream));
});
return signalingThread;
}
bool checkStream(hipStream_t stream) {
{ // Check default flags
auto res = checkStreamFlags_(stream, true, hipStreamDefault);
if (!res) return false;
}
{ // Check default Priority
auto res = checkStreamPriority_(stream);
if (!res) return false;
}
return true;
}
bool checkStreamPriorityAndFlags(hipStream_t stream, int priority, unsigned int flags) {
{ // Check flags
auto res = checkStreamFlags_(stream, true, flags);
if (!res) return false;
}
{ // Check priority
auto res = checkStreamPriority_(stream, true, priority);
if (!res) return false;
}
return true;
}
} // namespace stream
} // namespace hip
+61
Ver fichero
@@ -0,0 +1,61 @@
/*
Copyright (c) 2022 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.
*/
#pragma once
#include <hip_test_common.hh>
namespace hip {
inline namespace stream {
const hipStream_t nullStream = nullptr;
const hipStream_t streamPerThread = hipStreamPerThread;
/**
* @brief Kernel that signals a semaphore to change value from 0 to 1.
*
* @param semaphore the semaphore that needs to be signaled.
*/
__global__ void signaling_kernel(int* semaphore = nullptr);
/**
* @brief Kernel that busy waits until the specified semaphore goes from 0 to 1.
*
* @param semaphore the semaphore to wait for.
*/
__global__ void waiting_kernel(int* semaphore = nullptr);
/**
* @brief Creates a thread that runs a signaling_kernel on a non-blocking stream.
* hipStreamNonBlocking is used here to avoid interfering with tests for the Null Stream.
*
* @param semaphore memory location to signal
* @return std::thread thread that has to be joined after the testing is done.
*/
std::thread startSignalingThread(int* semaphore = nullptr);
// Checks stream for valid values of flags and priority
bool checkStream(hipStream_t stream);
// Checks stream for valid flags and a particular value of priority
bool checkStreamPriorityAndFlags(hipStream_t stream, int priority,
unsigned int flags = hipStreamDefault);
} // namespace stream
} // namespace hip
@@ -77,13 +77,23 @@ TEST_CASE("Unit_hipStreamPerThread_DeviceReset_2") {
for (unsigned int i = 0; i < ele_size; ++i) {
A_h[i] = 123;
}
hipMemcpyAsync(A_d, A_h, ele_size * sizeof(int), hipMemcpyHostToDevice,
status = hipMemcpyAsync(A_d, A_h, ele_size * sizeof(int), hipMemcpyHostToDevice,
hipStreamPerThread);
if (status != hipSuccess) return;
hipStreamSynchronize(hipStreamPerThread);
hipDeviceReset();
hipMemcpyAsync(A_d, A_h, ele_size * sizeof(int), hipMemcpyHostToDevice,
// After reset all memory objects will be destroyed hence allocating them again
// Intension is to use hipStreamPerThread successfully after reset hence not validating
// values after copy
status = hipHostMalloc(&A_h, ele_size*sizeof(int));
if (status != hipSuccess) return;
status = hipMalloc(&A_d, ele_size * sizeof(int));
if (status != hipSuccess) return;
status = hipMemcpyAsync(A_d, A_h, ele_size * sizeof(int), hipMemcpyHostToDevice,
hipStreamPerThread);
if (status != hipSuccess) return;
hipStreamSynchronize(hipStreamPerThread);
}
}
+3
Ver fichero
@@ -31,6 +31,9 @@ set(TEST_SRC
hipTextureRef2D.cc
hipSimpleTexture2DLayered.cc
hipTextureMipmapObj2D.cc
hipGetChanDesc.cc
hipTexObjPitch.cc
hipTextureObj1DFetch.cc
)
hip_add_exe_to_target(NAME TextureTest
+53
Ver fichero
@@ -0,0 +1,53 @@
/*
Copyright (c) 2022 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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#define R 8 // rows, height
#define C 8 // columns, width
TEST_CASE("Unit_hipGetChannelDesc_CreateAndGet") {
hipChannelFormatDesc chan_test, chan_desc;
hipArray *hipArray;
#if HT_AMD
int imageSupport{};
HIP_CHECK(hipDeviceGetAttribute(&imageSupport,
hipDeviceAttributeImageSupport, 0));
if (!imageSupport) {
INFO("Texture is not supported on the device. Test is skipped");
return;
}
#endif
chan_desc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindSigned);
HIP_CHECK(hipMallocArray(&hipArray, &chan_desc, C, R, 0));
HIP_CHECK(hipGetChannelDesc(&chan_test, hipArray));
if ((chan_test.x != 32) || (chan_test.y != 0)
|| (chan_test.z != 0) || (chan_test.f != 0)) {
INFO("Mismatch observed : " << chan_test.x << chan_test.y
<< chan_test.z << chan_test.f);
REQUIRE(false);
}
HIP_CHECK(hipFreeArray(hipArray));
}
+106
Ver fichero
@@ -0,0 +1,106 @@
/*
Copyright (c) 2022 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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#define SIZE_H 20
#define SIZE_W 179
// texture object is a kernel argument
template <typename TYPE_t>
static __global__ void texture2dCopyKernel(hipTextureObject_t texObj,
TYPE_t* dst) {
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
for (int i = 0; i < SIZE_H; i++)
for (int j = 0; j < SIZE_W; j++)
dst[SIZE_W*i+j] = tex2D<TYPE_t>(texObj, j, i);
__syncthreads();
#endif
}
TEMPLATE_TEST_CASE("Unit_hipTexObjPitch_texture2D", "", float, int,
unsigned char, int16_t, char, unsigned int) {
TestType* B;
TestType* A;
TestType* devPtrB;
TestType* devPtrA;
#if HT_AMD
int imageSupport{};
HIP_CHECK(hipDeviceGetAttribute(&imageSupport,
hipDeviceAttributeImageSupport, 0));
if (!imageSupport) {
INFO("Texture is not supported on the device. Test is skipped");
return;
}
#endif
B = new TestType[SIZE_H*SIZE_W];
A = new TestType[SIZE_H*SIZE_W];
for (size_t i=1; i <= (SIZE_H*SIZE_W); i++) {
A[i-1] = i;
}
size_t devPitchA;
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&devPtrA), &devPitchA,
SIZE_W*sizeof(TestType), SIZE_H));
HIP_CHECK(hipMemcpy2D(devPtrA, devPitchA, A, SIZE_W*sizeof(TestType),
SIZE_W*sizeof(TestType), SIZE_H, hipMemcpyHostToDevice));
// Use the texture object
hipResourceDesc texRes;
memset(&texRes, 0, sizeof(texRes));
texRes.resType = hipResourceTypePitch2D;
texRes.res.pitch2D.devPtr = devPtrA;
texRes.res.pitch2D.height = SIZE_H;
texRes.res.pitch2D.width = SIZE_W;
texRes.res.pitch2D.pitchInBytes = devPitchA;
texRes.res.pitch2D.desc = hipCreateChannelDesc<TestType>();
hipTextureDesc texDescr;
memset(&texDescr, 0, sizeof(texDescr));
texDescr.normalizedCoords = false;
texDescr.filterMode = hipFilterModePoint;
texDescr.mipmapFilterMode = hipFilterModePoint;
texDescr.addressMode[0] = hipAddressModeClamp;
texDescr.addressMode[1] = hipAddressModeClamp;
texDescr.addressMode[2] = hipAddressModeClamp;
texDescr.readMode = hipReadModeElementType;
hipTextureObject_t texObj;
HIP_CHECK(hipCreateTextureObject(&texObj, &texRes, &texDescr, NULL));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&devPtrB),
SIZE_W*sizeof(TestType)*SIZE_H));
hipLaunchKernelGGL(texture2dCopyKernel, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0,
texObj, devPtrB);
HIP_CHECK(hipMemcpy2D(B, SIZE_W*sizeof(TestType), devPtrB,
SIZE_W*sizeof(TestType), SIZE_W*sizeof(TestType),
SIZE_H, hipMemcpyDeviceToHost));
HipTest::checkArray(A, B, SIZE_H, SIZE_W);
delete []A;
delete []B;
hipFree(devPtrA);
hipFree(devPtrB);
}
@@ -0,0 +1,84 @@
/*
Copyright (c) 2022 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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#define N 512
static __global__ void tex1dKernel(float *val, hipTextureObject_t obj) {
int k = blockIdx.x * blockDim.x + threadIdx.x;
if (k < N) {
val[k] = tex1Dfetch<float>(obj, k);
}
}
TEST_CASE("Unit_hipCreateTextureObject_tex1DfetchVerification") {
// Allocating the required buffer on gpu device
float *texBuf, *texBufOut;
float val[N], output[N];
for (int i = 0; i < N; i++) {
val[i] = (i + 1) * (i + 1);
output[i] = 0.0;
}
HIP_CHECK(hipMalloc(&texBuf, N * sizeof(float)));
HIP_CHECK(hipMalloc(&texBufOut, N * sizeof(float)));
HIP_CHECK(hipMemcpy(texBuf, val, N * sizeof(float), hipMemcpyHostToDevice));
HIP_CHECK(hipMemset(texBufOut, 0, N * sizeof(float)));
hipResourceDesc resDescLinear;
memset(&resDescLinear, 0, sizeof(resDescLinear));
resDescLinear.resType = hipResourceTypeLinear;
resDescLinear.res.linear.devPtr = texBuf;
resDescLinear.res.linear.desc =
hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindFloat);
resDescLinear.res.linear.sizeInBytes = N * sizeof(float);
hipTextureDesc texDesc;
memset(&texDesc, 0, sizeof(texDesc));
texDesc.readMode = hipReadModeElementType;
// Creating texture object
hipTextureObject_t texObj = 0;
HIP_CHECK(hipCreateTextureObject(&texObj, &resDescLinear, &texDesc, NULL));
dim3 dimBlock(64, 1, 1);
dim3 dimGrid(N / dimBlock.x, 1, 1);
hipLaunchKernelGGL(tex1dKernel, dim3(dimGrid), dim3(dimBlock), 0, 0,
texBufOut, texObj);
HIP_CHECK(hipDeviceSynchronize());
HIP_CHECK(hipMemcpy(output, texBufOut, N * sizeof(float),
hipMemcpyDeviceToHost));
for (int i = 0; i < N; i++) {
if (output[i] != val[i]) {
INFO("Mismatch at index : " << i << ", output[i] " << output[i]
<< ", val[i] " << val[i]);
REQUIRE(false);
}
}
HIP_CHECK(hipDestroyTextureObject(texObj));
HIP_CHECK(hipFree(texBuf));
HIP_CHECK(hipFree(texBufOut));
}
+2 -2
Ver fichero
@@ -1,5 +1,5 @@
/*
Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2015 - 2022 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
@@ -20,7 +20,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s ../test_common.cpp NVCC_OPTIONS -lnvrtc
* BUILD: %t %s ../test_common.cpp HIPCC_OPTIONS -lhiprtc NVCC_OPTIONS -lnvrtc
* TEST: %t
* HIT_END
*/
+2 -2
Ver fichero
@@ -1,5 +1,5 @@
/*
Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2015 - 2022 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
@@ -20,7 +20,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s ../test_common.cpp NVCC_OPTIONS -lnvrtc
* BUILD: %t %s ../test_common.cpp HIPCC_OPTIONS -lhiprtc NVCC_OPTIONS -lnvrtc
* TEST: %t
* HIT_END
*/