AI/rocm-systems
2
0
Rozštěpit 0
Zdrojový kód Úkoly Pull requesty Akce Balíčky Projekty Vydání Wiki Aktivita

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

Procházet zdrojové kódy 
Change-Id: Ic4b3f1042c69f9810a7134a999c4779642852f59
Tento commit je obsažen v:
Jenkins
2021-09-06 07:11:25 -04:00
rodič ff00b75617 6a09344dba
revize 757b01e534
17 změnil soubory, kde provedl 371 přidání a 23 odebrání
Stáhněte soubor záplaty Stáhněte rozdílový soubor Rozbalit všechny soubory Sbalit všechny soubory
include/hip/hip_runtime_api.h
+237 -23
Zobrazit soubor
@@ -1441,6 +1441,7 @@ const char* hipGetErrorString(hipError_t hipError);
* The following Stream APIs are not (yet) supported in HIP:
* - hipStreamAttachMemAsync is a nop
*/
/**
* @brief Create an asynchronous stream.
*
@@ -2187,7 +2188,7 @@ hipError_t hipMemAdvise(const void* dev_ptr,
/**
* @brief Query an attribute of a given memory range in HIP.
*
* @param [in/out] data a pointer to a memory location where the result of each
* @param [in,out] data a pointer to a memory location where the result of each
* attribute query will be written to
* @param [in] data_size the size of data
* @param [in] attribute the attribute to query
@@ -2204,7 +2205,7 @@ hipError_t hipMemRangeGetAttribute(void* data,
/**
* @brief Query attributes of a given memory range in HIP.
*
* @param [in/out] data a two-dimensional array containing pointers to memory locations
* @param [in,out] data a two-dimensional array containing pointers to memory locations
* where the result of each attribute query will be written to
* @param [in] data_sizes an array, containing the sizes of each result
* @param [in] attributes the attribute to query
@@ -3652,7 +3653,7 @@ hipError_t hipLaunchCooperativeKernel(const void* f, dim3 gridDim, dim3 blockDim
* @brief Launches kernels on multiple devices where thread blocks can cooperate and
* synchronize as they execute.
*
* @param [in] hipLaunchParams List of launch parameters, one per device.
* @param [in] launchParamsList List of launch parameters, one per device.
* @param [in] numDevices Size of the launchParamsList array.
* @param [in] flags Flags to control launch behavior.
*
@@ -4138,6 +4139,7 @@ const char* hipApiName(uint32_t id);
const char* hipKernelNameRef(const hipFunction_t f);
const char* hipKernelNameRefByPtr(const void* hostFunction, hipStream_t stream);
int hipGetStreamDeviceId(hipStream_t stream);
/**
* An opaque value that represents a hip graph
*/
@@ -4150,6 +4152,12 @@ 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
@@ -4164,6 +4172,7 @@ typedef enum hipGraphNodeType {
hipGraphNodeTypeMemcpyToSymbol = 11, ///< MemcpyToSymbol node
hipGraphNodeTypeCount
} hipGraphNodeType;
typedef void (*hipHostFn_t)(void* userData);
typedef struct hipHostNodeParams {
hipHostFn_t fn;
@@ -4185,13 +4194,19 @@ typedef struct hipMemsetParams {
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 =
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
@@ -4199,6 +4214,7 @@ typedef enum hipGraphExecUpdateResult {
0x6, ///< The update failed because something about the node is not supported
hipGraphExecUpdateErrorUnsupportedFunctionChange = 0x7
} hipGraphExecUpdateResult;
typedef enum hipStreamCaptureMode {
hipStreamCaptureModeGlobal = 0,
hipStreamCaptureModeThreadLocal,
@@ -4212,63 +4228,261 @@ typedef enum hipStreamCaptureStatus {
} hipStreamCaptureStatus;
hipError_t hipStreamBeginCapture(hipStream_t stream, hipStreamCaptureMode mode);
hipError_t hipStreamEndCapture(hipStream_t stream, hipGraph_t* pGraph);
// Creates a graph.
/**
*-------------------------------------------------------------------------------------------------
*-------------------------------------------------------------------------------------------------
* @defgroup Graph Graph Management
* @{
* This section describes the graph management functions of HIP runtime API.
*/
/**
* @brief Creates a graph
*
* @param [out] pGraph - pointer to graph to create.
* @param [in] flags - flags for graph creation, must be 0.
*
* @returns #hipSuccess.
*
*/
hipError_t hipGraphCreate(hipGraph_t* pGraph, unsigned int flags);
// Destroys a graph.
/**
* @brief Destroys a graph
*
* @param [in] graph - instance of graph to destroy.
*
* @returns #hipSuccess.
*
*/
hipError_t hipGraphDestroy(hipGraph_t graph);
// Destroys an executable graph.
/**
* @brief Destroys an executable graph
*
* @param [in] pGraphExec - instance of executable graph to destry.
*
* @returns #hipSuccess.
*
*/
hipError_t hipGraphExecDestroy(hipGraphExec_t pGraphExec);
// Creates an executable graph from a graph.
/**
* @brief Creates an executable graph from a graph
*
* @param [out] pGraphExec - pointer to instantiated executable graph to create.
* @param [in] graph - instance of graph to instantiate.
* @param [out] pErrorNode - pointer to error node in case error occured in graph instantiation,
* it could modify the correponding node.
* @param [out] pLogBuffer - pointer to log buffer.
* @param [in] bufferSize - the size of log buffer.
*
* @returns #hipSuccess, #hipErrorOutOfMemory.
*
*/
hipError_t hipGraphInstantiate(hipGraphExec_t* pGraphExec, hipGraph_t graph,
hipGraphNode_t* pErrorNode, char* pLogBuffer, size_t bufferSize);
// Launches an executable graph in a stream.
/**
* @brief launches an executable graph in a stream
*
* @param [in] graphExec - instance of executable graph to launch.
* @param [in] stream - instance of stream in which to launch executable graph.
* @returns #hipSuccess, #hipErrorOutOfMemory, #hipErrorInvalidHandle, #hipErrorInvalidValue
*
*/
hipError_t hipGraphLaunch(hipGraphExec_t graphExec, hipStream_t stream);
// Creates a kernel execution node and adds it to a graph.
/**
* @brief Creates a kernel execution node and adds it to a graph.
*
* @param [out] pGraphNode - pointer to graph node to create.
* @param [in,out] graph - instance of graph to add the created node.
* @param [in] pDependencies - pointer to the dependencies on the kernel execution node.
* @param [in] numDependencies - the number of the dependencies.
* @param [in] pNodeParams - pointer to the parameters to the kernel execution node on the GPU.
* @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidDeviceFunction
*
*/
hipError_t hipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
const hipGraphNode_t* pDependencies, size_t numDependencies,
const hipKernelNodeParams* pNodeParams);
// Creates a memcpy node and adds it to a graph.
/**
* @brief Creates a memcpy node and adds it to a graph.
*
* @param [out] pGraphNode - pointer to graph node to create.
* @param [in,out] graph - instance of graph to add the created node.
* @param [in] pDependencies - const pointer to the dependencies on the kernel execution node.
* @param [in] numDependencies - the number of the dependencies.
* @param [in] pCopyParams - const pointer to the parameters for the memory copy.
* @returns #hipSuccess, #hipErrorInvalidValue
*
*/
hipError_t hipGraphAddMemcpyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
const hipGraphNode_t* pDependencies, size_t numDependencies,
const hipMemcpy3DParms* pCopyParams);
// Creates a 1D memcpy node and adds it to a graph.
/**
* @brief Creates a 1D memcpy node and adds it to a graph.
*
* @param [out] pGraphNode - pointer to graph node to create.
* @param [in,out] graph - instance of the graph to add the created node.
* @param [in] pDependencies - const pointer to the dependencies on the kernel execution node.
* @param [in] numDependencies - the number of the dependencies.
* @param [in] dst - pointer to memory address to the destination.
* @param [in] src - pointer to memory address to the source.
* @param [in] count - the size of the memory to copy.
* @param [in] kind - the type of memory copy.
* @returns #hipSuccess, #hipErrorInvalidValue
*
*/
hipError_t hipGraphAddMemcpyNode1D(hipGraphNode_t* pGraphNode, hipGraph_t graph,
const hipGraphNode_t* pDependencies, size_t numDependencies,
void* dst, const void* src, size_t count, hipMemcpyKind kind);
// Creates a memset node and adds it to a graph.
/**
* @brief Creates a memset node and adds it to a graph.
*
* @param [out] pGraphNode - pointer to the graph node to create.
* @param [in,out] graph - instance of the graph to add the created node.
* @param [in] pDependencies - const pointer to the dependencies on the kernel execution node.
* @param [in] numDependencies - the number of the dependencies.
* @param [in] pMemsetParams - const pointer to the parameters for the memory set.
* @returns #hipSuccess, #hipErrorInvalidValue
*
*/
hipError_t hipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
const hipGraphNode_t* pDependencies, size_t numDependencies,
const hipMemsetParams* pMemsetParams);
// Returns graph nodes
/**
* @brief Returns graph nodes.
*
* @param [in] graph - instance of graph to get the nodes.
* @param [out] nodes - pointer to the graph nodes.
* @param [out] numNodes - the number of graph nodes.
* @returns #hipSuccess, #hipErrorInvalidValue
*
*/
hipError_t hipGraphGetNodes(hipGraph_t graph, hipGraphNode_t* nodes, size_t* numNodes);
// Returns graph's root nodes.
/**
* @brief Returns graph's root nodes.
*
* @param [in] graph - instance of the graph to get the nodes.
* @param [out] pRootNodes - pointer to the graph's root nodes.
* @param [out] pNumRootNodes - the number of graph's root nodes.
* @returns #hipSuccess, #hipErrorInvalidValue
*
*/
hipError_t hipGraphGetRootNodes(hipGraph_t graph, hipGraphNode_t* pRootNodes,
size_t* pNumRootNodes);
// Returns a kernel node's parameters.
/**
* @brief Gets kernel node's parameters.
*
* @param [in] node - instance of the node to get parameters from.
* @param [out] pNodeParams - pointer to the parameters
* @returns #hipSuccess, #hipErrorInvalidValue
*
*/
hipError_t hipGraphKernelNodeGetParams(hipGraphNode_t node, hipKernelNodeParams* pNodeParams);
// Sets a kernel node's parameters.
/**
* @brief Sets a kernel node's parameters.
*
* @param [in] node - instance of the node to set parameters to.
* @param [in] pNodeParams - const pointer to the parameters.
* @returns #hipSuccess, #hipErrorInvalidValue
*
*/
hipError_t hipGraphKernelNodeSetParams(hipGraphNode_t node, const hipKernelNodeParams* pNodeParams);
// Returns a memcpy node's parameters.
/**
* @brief Gets a memcpy node's parameters.
*
* @param [in] node - instance of the node to get parameters from.
* @param [out] pNodeParams - pointer to the parameters.
* @returns #hipSuccess, #hipErrorInvalidValue
*
*/
hipError_t hipGraphMemcpyNodeGetParams(hipGraphNode_t node, hipMemcpy3DParms* pNodeParams);
// Sets a memcpy node's parameters.
/**
* @brief Sets a memcpy node's parameters.
*
* @param [in] node - instance of the node to set parameters to.
* @param [in] pNodeParams - const pointer to the parameters.
* @returns #hipSuccess, #hipErrorInvalidValue
*
*/
hipError_t hipGraphMemcpyNodeSetParams(hipGraphNode_t node, const hipMemcpy3DParms* pNodeParams);
// Returns a memset node's parameters.
/**
* @brief Gets a memset node's parameters.
*
* @param [in] node - instane of the node to get parameters from.
* @param [out] pNodeParams - pointer to the parameters.
* @returns #hipSuccess, #hipErrorInvalidValue
*
*/
hipError_t hipGraphMemsetNodeGetParams(hipGraphNode_t node, hipMemsetParams* pNodeParams);
// Sets a memset node's parameters.
/**
* @brief Sets a memset node's parameters.
*
* @param [in] node - instance of the node to set parameters to.
* @param [out] pNodeParams - pointer to the parameters.
* @returns #hipSuccess, #hipErrorInvalidValue
*
*/
hipError_t hipGraphMemsetNodeSetParams(hipGraphNode_t node, const hipMemsetParams* pNodeParams);
// Sets the parameters for a kernel node in the given graphExec.
/**
* @brief Sets the parameters for a kernel node in the given graphExec.
*
* @param [in] hGraphExec - instance of the executable graph with the node.
* @param [in] node - instance of the node to set parameters to.
* @param [in] pNodeParams - const pointer to the kernel node parameters.
* @returns #hipSuccess, #hipErrorInvalidValue
*
*/
hipError_t hipGraphExecKernelNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t node,
const hipKernelNodeParams* pNodeParams);
// Adds dependency edges to a graph.
/**
* @brief Adds dependency edges to a graph.
*
* @param [in] graph - instance of the graph to add dependencies.
* @param [in] from - pointer to the graph nodes with dependenties to add from.
* @param [in] to - pointer to the graph nodes to add dependenties to.
* @param [in] numDependencies - the number of dependencies to add.
* @returns #hipSuccess, #hipErrorInvalidValue
*
*/
hipError_t hipGraphAddDependencies(hipGraph_t graph, const hipGraphNode_t* from,
const hipGraphNode_t* to, size_t numDependencies);
// Creates an empty node and adds it to a graph.
/**
* @brief Creates an empty node and adds it to a graph.
*
* @param [out] pGraphNode - pointer to the graph node to create and add to the graph.
* @param [in,out] graph - instane of the graph the node is add to.
* @param [in] pDependencies - const pointer to the node dependenties.
* @param [in] numDependencies - the number of dependencies.
* @returns #hipSuccess, #hipErrorInvalidValue
*
*/
hipError_t hipGraphAddEmptyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
const hipGraphNode_t* pDependencies, size_t numDependencies);
// doxygen end graph API
/**
* @}
*/
#ifdef __cplusplus
} /* extern "c" */
#endif
tests/catch/unit/deviceLib/CMakeLists.txt
+1
Zobrazit soubor
@@ -17,6 +17,7 @@ set(TEST_SRC
# AMD only tests
set(AMD_TEST_SRC
unsafeAtomicAdd.cc
vectorTypesDevice.cc
mbcnt.cc
bitExtract.cc
tests/catch/unit/deviceLib/unsafeAtomicAdd.cc
+133
Zobrazit soubor
@@ -0,0 +1,133 @@
/*
Copyright (c) 2021 - present 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 <cmath>
static constexpr auto kernel{
R"(
extern "C"
__global__
void unsafeAdd_f(float *p, float v)
{
auto val = unsafeAtomicAdd(p, v);
}
extern "C"
__global__
void unsafeAdd_d(double *p, double v)
{
auto val = unsafeAtomicAdd(p, v);
}
)"};
TEST_CASE("Unit_unsafeAtomicAdd") {
using namespace std;
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;
hiprtcCreateProgram(&prog, // prog
kernel, // 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) {
string log(logSize, '\0');
HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0]));
INFO(log);
}
REQUIRE(compileResult == HIPRTC_SUCCESS);
size_t codeSize;
HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize));
vector<char> code(codeSize);
HIPRTC_CHECK(hiprtcGetCode(prog, code.data()));
HIPRTC_CHECK(hiprtcDestroyProgram(&prog));
float* fX;
double* dX;
HIP_CHECK(hipMalloc(&fX, sizeof(float)));
HIP_CHECK(hipMalloc(&dX, sizeof(double)));
hipModule_t module;
hipFunction_t f_kernel, d_kernel;
HIP_CHECK(hipModuleLoadData(&module, code.data()));
HIP_CHECK(hipModuleGetFunction(&f_kernel, module, "unsafeAdd_f"));
HIP_CHECK(hipModuleGetFunction(&d_kernel, module, "unsafeAdd_d"));
float f_val = 10.1f;
double d_val = 10.1;
HIP_CHECK(hipMemcpy(fX, &f_val, sizeof(float), hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(dX, &d_val, sizeof(double), hipMemcpyHostToDevice));
struct {
float* p;
float val;
} args_f{fX, f_val};
struct {
double* p;
double val;
} args_d{dX, d_val};
auto size_f = sizeof(args_f);
void* config_f[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f, HIP_LAUNCH_PARAM_BUFFER_SIZE,
&size_f, HIP_LAUNCH_PARAM_END};
auto size_d = sizeof(args_d);
void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_d, HIP_LAUNCH_PARAM_BUFFER_SIZE,
&size_d, HIP_LAUNCH_PARAM_END};
hipModuleLaunchKernel(f_kernel, 10, 1, 1, 100, 1, 1, 0, nullptr, nullptr, config_f);
hipModuleLaunchKernel(d_kernel, 10, 1, 1, 100, 1, 1, 0, nullptr, nullptr, config_d);
float res_f = 0.0f;
double res_d = 0.0;
HIP_CHECK(hipMemcpy(&res_f, fX, sizeof(float), hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(&res_d, dX, sizeof(double), hipMemcpyDeviceToHost));
hipFree(dX);
hipFree(fX);
HIP_CHECK(hipModuleUnload(module));
REQUIRE(fabs(res_f - (1000 * f_val)) <= 0.2f);
REQUIRE(fabs(res_d - (1000 * d_val)) <= 0.2);
}
}
tests/catch/unit/stream/CMakeLists.txt
Spustitelný soubor → Normální soubor
Zobrazit soubor
tests/catch/unit/stream/hipMultiStream.cc
Spustitelný soubor → Normální soubor
Zobrazit soubor
tests/catch/unit/stream/hipStreamCreate.cc
Spustitelný soubor → Normální soubor
Zobrazit soubor
tests/catch/unit/stream/hipStreamGetFlags.cc
Spustitelný soubor → Normální soubor
Zobrazit soubor
tests/catch/unit/stream/hipStreamGetPriority.cc
Spustitelný soubor → Normální soubor
Zobrazit soubor
tests/src/p2p/hipPeerToPeer_simple.cpp
Spustitelný soubor → Normální soubor
Zobrazit soubor
tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupType.cpp
Spustitelný soubor → Normální soubor
Zobrazit soubor
tests/src/runtimeApi/module/hipExtLaunchKernelGGL.cpp
Spustitelný soubor → Normální soubor
Zobrazit soubor
tests/src/runtimeApi/module/hipModule.cpp
Spustitelný soubor → Normální soubor
Zobrazit soubor
tests/src/runtimeApi/module/hipModuleOccupancyMaxPotentialActiveBlockSize.cpp
Spustitelný soubor → Normální soubor
Zobrazit soubor
tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp
Spustitelný soubor → Normální soubor
Zobrazit soubor
tests/src/runtimeApi/module/matmul.cpp
Spustitelný soubor → Normální soubor
Zobrazit soubor
tests/src/runtimeApi/module/tex2d_kernel.cpp
Spustitelný soubor → Normální soubor
Zobrazit soubor
tests/src/runtimeApi/p2p/hipDeviceGetP2PAttribute.cpp
Spustitelný soubor → Normální soubor
Zobrazit soubor