From af3c38c477b23e98d9277eaa9b6d7766003cd4f6 Mon Sep 17 00:00:00 2001 From: Julia Jiang <56359287+jujiang-del@users.noreply.github.com> Date: Wed, 1 Sep 2021 23:17:38 -0400 Subject: [PATCH 1/3] SWDEV-291894 - Add doxygen for hip graph APIs (#2328) Change-Id: Id878125b34711fc43570dd680c05310d9e2a4015 --- include/hip/hip_runtime_api.h | 260 +++++++++++++++++++++++++++++++--- 1 file changed, 237 insertions(+), 23 deletions(-) diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index 2aa438809b..2e939273ab 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -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 From 56d6d27a6f73470f8ab12f817b886e993ac4d6ae Mon Sep 17 00:00:00 2001 From: Julia Jiang <56359287+jujiang-del@users.noreply.github.com> Date: Wed, 1 Sep 2021 23:17:57 -0400 Subject: [PATCH 2/3] SWDEV-286322 - remove unnecessary executable permission in some HIP files (#2329) Change-Id: If992963761a73d709e08330bc36f2c979d5dcae9 --- tests/catch/unit/stream/CMakeLists.txt | 0 tests/catch/unit/stream/hipMultiStream.cc | 0 tests/catch/unit/stream/hipStreamCreate.cc | 0 tests/catch/unit/stream/hipStreamGetFlags.cc | 0 tests/catch/unit/stream/hipStreamGetPriority.cc | 0 tests/src/p2p/hipPeerToPeer_simple.cpp | 0 tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupType.cpp | 0 tests/src/runtimeApi/module/hipExtLaunchKernelGGL.cpp | 0 tests/src/runtimeApi/module/hipModule.cpp | 0 .../module/hipModuleOccupancyMaxPotentialActiveBlockSize.cpp | 0 tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp | 0 tests/src/runtimeApi/module/matmul.cpp | 0 tests/src/runtimeApi/module/tex2d_kernel.cpp | 0 tests/src/runtimeApi/p2p/hipDeviceGetP2PAttribute.cpp | 0 14 files changed, 0 insertions(+), 0 deletions(-) mode change 100755 => 100644 tests/catch/unit/stream/CMakeLists.txt mode change 100755 => 100644 tests/catch/unit/stream/hipMultiStream.cc mode change 100755 => 100644 tests/catch/unit/stream/hipStreamCreate.cc mode change 100755 => 100644 tests/catch/unit/stream/hipStreamGetFlags.cc mode change 100755 => 100644 tests/catch/unit/stream/hipStreamGetPriority.cc mode change 100755 => 100644 tests/src/p2p/hipPeerToPeer_simple.cpp mode change 100755 => 100644 tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupType.cpp mode change 100755 => 100644 tests/src/runtimeApi/module/hipExtLaunchKernelGGL.cpp mode change 100755 => 100644 tests/src/runtimeApi/module/hipModule.cpp mode change 100755 => 100644 tests/src/runtimeApi/module/hipModuleOccupancyMaxPotentialActiveBlockSize.cpp mode change 100755 => 100644 tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp mode change 100755 => 100644 tests/src/runtimeApi/module/matmul.cpp mode change 100755 => 100644 tests/src/runtimeApi/module/tex2d_kernel.cpp mode change 100755 => 100644 tests/src/runtimeApi/p2p/hipDeviceGetP2PAttribute.cpp diff --git a/tests/catch/unit/stream/CMakeLists.txt b/tests/catch/unit/stream/CMakeLists.txt old mode 100755 new mode 100644 diff --git a/tests/catch/unit/stream/hipMultiStream.cc b/tests/catch/unit/stream/hipMultiStream.cc old mode 100755 new mode 100644 diff --git a/tests/catch/unit/stream/hipStreamCreate.cc b/tests/catch/unit/stream/hipStreamCreate.cc old mode 100755 new mode 100644 diff --git a/tests/catch/unit/stream/hipStreamGetFlags.cc b/tests/catch/unit/stream/hipStreamGetFlags.cc old mode 100755 new mode 100644 diff --git a/tests/catch/unit/stream/hipStreamGetPriority.cc b/tests/catch/unit/stream/hipStreamGetPriority.cc old mode 100755 new mode 100644 diff --git a/tests/src/p2p/hipPeerToPeer_simple.cpp b/tests/src/p2p/hipPeerToPeer_simple.cpp old mode 100755 new mode 100644 diff --git a/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupType.cpp b/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupType.cpp old mode 100755 new mode 100644 diff --git a/tests/src/runtimeApi/module/hipExtLaunchKernelGGL.cpp b/tests/src/runtimeApi/module/hipExtLaunchKernelGGL.cpp old mode 100755 new mode 100644 diff --git a/tests/src/runtimeApi/module/hipModule.cpp b/tests/src/runtimeApi/module/hipModule.cpp old mode 100755 new mode 100644 diff --git a/tests/src/runtimeApi/module/hipModuleOccupancyMaxPotentialActiveBlockSize.cpp b/tests/src/runtimeApi/module/hipModuleOccupancyMaxPotentialActiveBlockSize.cpp old mode 100755 new mode 100644 diff --git a/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp b/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp old mode 100755 new mode 100644 diff --git a/tests/src/runtimeApi/module/matmul.cpp b/tests/src/runtimeApi/module/matmul.cpp old mode 100755 new mode 100644 diff --git a/tests/src/runtimeApi/module/tex2d_kernel.cpp b/tests/src/runtimeApi/module/tex2d_kernel.cpp old mode 100755 new mode 100644 diff --git a/tests/src/runtimeApi/p2p/hipDeviceGetP2PAttribute.cpp b/tests/src/runtimeApi/p2p/hipDeviceGetP2PAttribute.cpp old mode 100755 new mode 100644 From 6a09344dba91a1a9816cb6bcdcc6d8bc6ea564c3 Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary <51944368+cjatin@users.noreply.github.com> Date: Mon, 6 Sep 2021 12:21:38 +0530 Subject: [PATCH 3/3] SWDEV-296918 - Adding unsafe atomics test for gfx90a (#2331) Change-Id: I3936e0b30e15aff154bdab3efe3f6da2ccce3578 --- tests/catch/unit/deviceLib/CMakeLists.txt | 1 + tests/catch/unit/deviceLib/unsafeAtomicAdd.cc | 133 ++++++++++++++++++ 2 files changed, 134 insertions(+) create mode 100644 tests/catch/unit/deviceLib/unsafeAtomicAdd.cc diff --git a/tests/catch/unit/deviceLib/CMakeLists.txt b/tests/catch/unit/deviceLib/CMakeLists.txt index ce11801786..191553ca99 100644 --- a/tests/catch/unit/deviceLib/CMakeLists.txt +++ b/tests/catch/unit/deviceLib/CMakeLists.txt @@ -17,6 +17,7 @@ set(TEST_SRC # AMD only tests set(AMD_TEST_SRC + unsafeAtomicAdd.cc vectorTypesDevice.cc mbcnt.cc bitExtract.cc diff --git a/tests/catch/unit/deviceLib/unsafeAtomicAdd.cc b/tests/catch/unit/deviceLib/unsafeAtomicAdd.cc new file mode 100644 index 0000000000..24e1329c7a --- /dev/null +++ b/tests/catch/unit/deviceLib/unsafeAtomicAdd.cc @@ -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 + +#include +#include + +#include + +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 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); + } +}