Merge branch 'amd-develop' into amd-master
Change-Id: I77fa88b460be549bfcf9e18d3212e732ffc045f5
Этот коммит содержится в:
Maneesh Gupta
@@ -335,7 +335,7 @@ HIP provides 3 environment variables in the HIP_*_BLOCKING family. These introd
|
||||
|
||||
These options cause HCC to serialize. Useful if you have libraries or code which is calling HCC kernels directly rather than using HIP.
|
||||
- HCC_SERIALZIE_KERNELS : 0x1=pre-serialize before each kernel launch, 0x2=post-serialize after each kernel launch., 0x3= pre- and post- serialize.
|
||||
- HCC_SERIALIZE_COPY : 0x1=pre-serialize before each async copy, 0x2=post-serialize after each async copy., 0x3= pre- and post- serialize.0
|
||||
- HCC_SERIALIZE_COPY : 0x1=pre-serialize before each async copy, 0x2=post-serialize after each async copy., 0x3= pre- and post- serialize.
|
||||
|
||||
- HSA_ENABLE_SDMA=0 : Causes host-to-device and device-to-host copies to use compute shader blit kernels rather than the dedicated DMA copy engines. Compute shader copies have low latency (typically < 5us) and can achieve approximately 80% of the bandwidth of the DMA copy engine. This flag is useful to isolate issues with the hardware copy engines.
|
||||
- HSA_ENABLE_INTERRUPT=0 : Causes completion signals to be detected with memory-based polling rather than interrupts. Can be useful to diagnose interrupt storm issues in the driver.
|
||||
|
||||
@@ -1465,11 +1465,11 @@ public:
|
||||
// to workaround the 'const' MacroArgs passed into this hook.
|
||||
const Token *start = Args->getUnexpArgument(i);
|
||||
size_t len = Args->getArgLength(start) + 1;
|
||||
#if (LLVM_VERSION_MAJOR >= 3) && (LLVM_VERSION_MINOR >= 9)
|
||||
#if (LLVM_VERSION_MAJOR >= 3) && (LLVM_VERSION_MINOR >= 9)
|
||||
_pp->EnterTokenStream(ArrayRef<Token>(start, len), false);
|
||||
#else
|
||||
#else
|
||||
_pp->EnterTokenStream(start, len, false, false);
|
||||
#endif
|
||||
#endif
|
||||
do {
|
||||
toks.push_back(Token());
|
||||
Token &tk = toks.back();
|
||||
@@ -1489,8 +1489,17 @@ public:
|
||||
<< " found as an actual argument in expansion of macro "
|
||||
<< macroName << "\n"
|
||||
<< "will be replaced with: " << repName << "\n");
|
||||
size_t length = name.size();
|
||||
SourceLocation sl = tok.getLocation();
|
||||
Replacement Rep(*_sm, sl, name.size(), repName);
|
||||
if (_sm->isMacroBodyExpansion(sl)) {
|
||||
LangOptions DefaultLangOptions;
|
||||
SourceLocation sl_macro = _sm->getExpansionLoc(sl);
|
||||
SourceLocation sl_end = Lexer::getLocForEndOfToken(sl_macro, 0, *_sm, DefaultLangOptions);
|
||||
length = _sm->getCharacterData(sl_end) - _sm->getCharacterData(sl_macro);
|
||||
name = StringRef(_sm->getCharacterData(sl_macro), length);
|
||||
sl = sl_macro;
|
||||
}
|
||||
Replacement Rep(*_sm, sl, length, repName);
|
||||
Replace->insert(Rep);
|
||||
}
|
||||
} else if (tok.isLiteral()) {
|
||||
|
||||
@@ -25,9 +25,7 @@ THE SOFTWARE.
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
|
||||
#define __CLANG_VERSION__ __clang_major__ * 10 + __clang_minor__
|
||||
|
||||
#ifdef HIP_HALF_HW_SUPPORT
|
||||
#if 0
|
||||
|
||||
typedef __fp16 __half;
|
||||
|
||||
|
||||
@@ -33,7 +33,6 @@ THE SOFTWARE.
|
||||
|
||||
#include <hip/hcc_detail/host_defines.h>
|
||||
#include <hip/hip_runtime_api.h>
|
||||
//#include "hip/hip_hcc.h"
|
||||
|
||||
#if defined (__HCC__) && (__hcc_workweek__ < 16155)
|
||||
#error("This version of HIP requires a newer version of HCC.");
|
||||
@@ -53,7 +52,7 @@ extern "C" {
|
||||
typedef struct ihipCtx_t *hipCtx_t;
|
||||
|
||||
// Note many APIs also use integer deviceIds as an alternative to the device pointer:
|
||||
typedef struct ihipDevice_t *hipDevice_t;
|
||||
typedef int hipDevice_t;
|
||||
|
||||
typedef struct ihipStream_t *hipStream_t;
|
||||
|
||||
@@ -72,7 +71,14 @@ typedef struct ihipIpcEventHandle_t *hipIpcEventHandle_t;
|
||||
|
||||
typedef struct ihipModule_t *hipModule_t;
|
||||
|
||||
typedef struct ihipFunction_t *hipFunction_t;
|
||||
struct ihipModuleSymbol_t{
|
||||
uint64_t _object; // The kernel object.
|
||||
uint32_t _groupSegmentSize;
|
||||
uint32_t _privateSegmentSize;
|
||||
char _name[64]; // TODO - review for performance cost. Name is just used for debug.
|
||||
};
|
||||
|
||||
typedef struct ihipModuleSymbol_t hipFunction_t;
|
||||
|
||||
typedef void* hipDeviceptr_t;
|
||||
|
||||
@@ -1904,18 +1910,6 @@ hipError_t hipIpcCloseMemHandle(void *devPtr);
|
||||
} /* extern "c" */
|
||||
#endif
|
||||
|
||||
#ifdef __cplusplus
|
||||
/**
|
||||
* @brief Returns a PCI Bus Id string for the device.
|
||||
* @param [out] pciBusId
|
||||
* @param [in] len
|
||||
* @param [hipDevice_t] device
|
||||
*
|
||||
* @returns #hipSuccess, #hipErrorInavlidDevice
|
||||
*/
|
||||
hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len,hipDevice_t device);
|
||||
#endif
|
||||
|
||||
/**
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
|
||||
@@ -1,50 +0,0 @@
|
||||
_ Add AQL kernel.
|
||||
_ Fix &*kernel command so the kernel name/type is an argument not a new command.
|
||||
|
||||
_ Add command to parse only.
|
||||
_ Add regression to parse all the hcm files.
|
||||
|
||||
_ Partition HCC, HIP, HSA, OpenCL commands into separate files.
|
||||
|
||||
|
||||
_ Show time for back-to-back copies.
|
||||
_ Add variables.
|
||||
%loopcnt
|
||||
|
||||
./hipCommander %loopcnt=4
|
||||
|
||||
_ Add datasize command.
|
||||
|
||||
|
||||
_ Add ( ) to parsing.
|
||||
_ Add argument parsing and checking.
|
||||
|
||||
_ Add verbose option to print each step of setup.
|
||||
- print deliniater between setup and run. Add run start message.
|
||||
|
||||
- print sizes of all buffers.
|
||||
- print each command before running.
|
||||
- show start/stop of timer routine.
|
||||
|
||||
_
|
||||
_ Clear documentation on what each oepration does.
|
||||
_ Add time instrumentation for each command.
|
||||
_ Add pcie atomic.
|
||||
|
||||
|
||||
_ Add tests for negative cases, ie endloop w/o opening loop.
|
||||
|
||||
|
||||
README tips
|
||||
---
|
||||
- HIP_API_TRACE combined with -v is useful to track the exact commands generates by hipCommander.
|
||||
|
||||
|
||||
Other ideas:
|
||||
---
|
||||
[ ] Perf guide : stream creation very slow on HCC and should be avoided.
|
||||
|
||||
|
||||
Scratch:
|
||||
|
||||
|
||||
@@ -63,6 +63,14 @@ double bytesToGB(size_t s)
|
||||
return (double)s / (1024.0*1024.0*1024.0);
|
||||
}
|
||||
|
||||
#define printLimit(w1, limit, units) \
|
||||
{\
|
||||
size_t val;\
|
||||
cudaDeviceGetLimit(&val, limit);\
|
||||
std::cout << setw(w1) << #limit": " << val << " " << units << std::endl;\
|
||||
}
|
||||
|
||||
|
||||
void printDeviceProp (int deviceId)
|
||||
{
|
||||
using namespace std;
|
||||
@@ -144,6 +152,17 @@ void printDeviceProp (int deviceId)
|
||||
cout << endl;
|
||||
|
||||
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
// Limits:
|
||||
cout << endl;
|
||||
printLimit(w1, cudaLimitStackSize, "bytes/thread");
|
||||
printLimit(w1, cudaLimitPrintfFifoSize, "bytes/device");
|
||||
printLimit(w1, cudaLimitMallocHeapSize, "bytes/device");
|
||||
printLimit(w1, cudaLimitDevRuntimeSyncDepth, "grids");
|
||||
printLimit(w1, cudaLimitDevRuntimePendingLaunchCount, "launches");
|
||||
#endif
|
||||
|
||||
|
||||
|
||||
|
||||
cout << endl;
|
||||
|
||||
@@ -14,7 +14,8 @@ 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,
|
||||
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.
|
||||
*/
|
||||
|
||||
@@ -57,8 +57,8 @@ hipError_t hipCtxCreate(hipCtx_t *ctx, unsigned int flags, hipDevice_t device)
|
||||
{
|
||||
HIP_INIT_API(ctx, flags, device); // FIXME - review if we want to init
|
||||
hipError_t e = hipSuccess;
|
||||
|
||||
*ctx = new ihipCtx_t(device, g_deviceCnt, flags);
|
||||
auto deviceHandle = ihipGetDevice(device);
|
||||
*ctx = new ihipCtx_t(deviceHandle, g_deviceCnt, flags);
|
||||
ihipSetTlsDefaultCtx(*ctx);
|
||||
tls_ctxStack.push(*ctx);
|
||||
|
||||
@@ -69,11 +69,13 @@ hipError_t hipDeviceGet(hipDevice_t *device, int deviceId)
|
||||
{
|
||||
HIP_INIT_API(device, deviceId); // FIXME - review if we want to init
|
||||
|
||||
*device = ihipGetDevice(deviceId);
|
||||
auto deviceHandle = ihipGetDevice(deviceId);
|
||||
|
||||
hipError_t e = hipSuccess;
|
||||
if (*device == NULL) {
|
||||
if (deviceHandle == NULL) {
|
||||
e = hipErrorInvalidDevice;
|
||||
} else {
|
||||
*device = deviceId;
|
||||
}
|
||||
|
||||
return ihipLogStatus(e);
|
||||
@@ -199,9 +201,11 @@ hipError_t hipCtxGetDevice(hipDevice_t *device)
|
||||
|
||||
if(ctx == nullptr) {
|
||||
e = hipErrorInvalidContext;
|
||||
// TODO *device = nullptr;
|
||||
}
|
||||
else {
|
||||
*device = (ihipDevice_t*)ctx->getDevice();
|
||||
auto deviceHandle = ctx->getDevice();
|
||||
*device = deviceHandle->_deviceId;
|
||||
}
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
@@ -321,9 +321,8 @@ hipError_t hipDeviceComputeCapability(int *major, int *minor, hipDevice_t device
|
||||
{
|
||||
HIP_INIT_API(major,minor, device);
|
||||
hipError_t e = hipSuccess;
|
||||
int deviceId= device->_deviceId;
|
||||
e = ihipDeviceGetAttribute(major, hipDeviceAttributeComputeCapabilityMajor, deviceId);
|
||||
e = ihipDeviceGetAttribute(minor, hipDeviceAttributeComputeCapabilityMinor, deviceId);
|
||||
e = ihipDeviceGetAttribute(major, hipDeviceAttributeComputeCapabilityMajor, device);
|
||||
e = ihipDeviceGetAttribute(minor, hipDeviceAttributeComputeCapabilityMinor, device);
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
@@ -331,28 +330,13 @@ hipError_t hipDeviceGetName(char *name,int len,hipDevice_t device)
|
||||
{
|
||||
HIP_INIT_API(name,len, device);
|
||||
hipError_t e = hipSuccess;
|
||||
int nameLen = strlen(device->_props.name);
|
||||
auto deviceHandle = ihipGetDevice(device);
|
||||
int nameLen = strlen(deviceHandle->_props.name);
|
||||
if(nameLen <= len)
|
||||
memcpy(name,device->_props.name,nameLen);
|
||||
memcpy(name,deviceHandle->_props.name,nameLen);
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
#ifdef __cplusplus
|
||||
hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len,hipDevice_t device)
|
||||
{
|
||||
HIP_INIT_API(pciBusId, len, device);
|
||||
hipError_t e = hipSuccess;
|
||||
int deviceId= device->_deviceId;
|
||||
int tempPciBusId = 0;
|
||||
e = ihipDeviceGetAttribute( &tempPciBusId, hipDeviceAttributePciBusId, deviceId);
|
||||
if( e == hipSuccess) {
|
||||
std::string tempPciStr = std::to_string(tempPciBusId);
|
||||
memcpy( pciBusId , tempPciStr.c_str() , tempPciStr.length() );
|
||||
}
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
#endif
|
||||
|
||||
hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len, int device)
|
||||
{
|
||||
HIP_INIT_API(pciBusId, len, device);
|
||||
@@ -365,11 +349,13 @@ hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len, int device)
|
||||
}
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
hipError_t hipDeviceTotalMem (size_t *bytes,hipDevice_t device)
|
||||
{
|
||||
HIP_INIT_API(bytes, device);
|
||||
hipError_t e = hipSuccess;
|
||||
*bytes= device->_props.totalGlobalMem;
|
||||
auto deviceHandle = ihipGetDevice(device);
|
||||
*bytes= deviceHandle->_props.totalGlobalMem;
|
||||
return ihipLogStatus(e);
|
||||
}
|
||||
|
||||
|
||||
@@ -22,8 +22,6 @@ THE SOFTWARE.
|
||||
|
||||
#include"hip/hip_fp16.h"
|
||||
|
||||
#if __CLANG_VERSION__ == 35
|
||||
|
||||
static const unsigned sign_val = 0x8000;
|
||||
static const __half __half_value_one_float = {0x3C00};
|
||||
static const __half __half_value_zero_float = {0x0};
|
||||
@@ -374,4 +372,3 @@ __device__ __half2 __lowhigh2highlow(const __half2 a){
|
||||
__device__ __half2 __low2half2(const __half2 a, const __half2 b){
|
||||
return {a.q, b.q};
|
||||
}
|
||||
#endif
|
||||
|
||||
@@ -387,57 +387,6 @@ void ihipStream_t::lockclose_postKernelCommand(const char * kernelName, hc::acce
|
||||
};
|
||||
|
||||
|
||||
|
||||
|
||||
#if USE_DISPATCH_HSA_KERNEL==0
|
||||
// Precursor: the stream is already locked,specifically so this routine can enqueue work into the specified av.
|
||||
void ihipStream_t::launchModuleKernel(
|
||||
hc::accelerator_view av,
|
||||
hsa_signal_t signal,
|
||||
uint32_t blockDimX,
|
||||
uint32_t blockDimY,
|
||||
uint32_t blockDimZ,
|
||||
uint32_t gridDimX,
|
||||
uint32_t gridDimY,
|
||||
uint32_t gridDimZ,
|
||||
uint32_t groupSegmentSize,
|
||||
uint32_t privateSegmentSize,
|
||||
void *kernarg,
|
||||
size_t kernSize,
|
||||
uint64_t kernel){
|
||||
hsa_status_t status;
|
||||
hsa_queue_t *Queue = (hsa_queue_t*)av.get_hsa_queue();
|
||||
const uint32_t queue_mask = Queue->size-1;
|
||||
uint32_t packet_index = hsa_queue_load_write_index_relaxed(Queue);
|
||||
hsa_kernel_dispatch_packet_t *dispatch_packet = &(((hsa_kernel_dispatch_packet_t*)(Queue->base_address))[packet_index & queue_mask]);
|
||||
|
||||
dispatch_packet->completion_signal = signal;
|
||||
dispatch_packet->workgroup_size_x = blockDimX;
|
||||
dispatch_packet->workgroup_size_y = blockDimY;
|
||||
dispatch_packet->workgroup_size_z = blockDimZ;
|
||||
dispatch_packet->grid_size_x = blockDimX * gridDimX;
|
||||
dispatch_packet->grid_size_y = blockDimY * gridDimY;
|
||||
dispatch_packet->grid_size_z = blockDimZ * gridDimZ;
|
||||
dispatch_packet->group_segment_size = groupSegmentSize;
|
||||
dispatch_packet->private_segment_size = privateSegmentSize;
|
||||
dispatch_packet->kernarg_address = kernarg;
|
||||
dispatch_packet->kernel_object = kernel;
|
||||
uint16_t header = (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) |
|
||||
(1 << HSA_PACKET_HEADER_BARRIER) |
|
||||
(HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) |
|
||||
(HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);
|
||||
|
||||
uint16_t setup = 3 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS;
|
||||
uint32_t header32 = header | (setup << 16);
|
||||
|
||||
__atomic_store_n((uint32_t*)(dispatch_packet), header32, __ATOMIC_RELEASE);
|
||||
|
||||
hsa_queue_store_write_index_relaxed(Queue, packet_index + 1);
|
||||
hsa_signal_store_relaxed(Queue->doorbell_signal, packet_index);
|
||||
}
|
||||
#endif
|
||||
|
||||
|
||||
//=============================================================================
|
||||
// Recompute the peercnt and the packed _peerAgents whenever a peer is added or deleted.
|
||||
// The packed _peerAgents can efficiently be used on each memory allocation.
|
||||
@@ -783,13 +732,10 @@ hipError_t ihipDevice_t::initProperties(hipDeviceProp_t* prop)
|
||||
_isLargeBar = _acc.has_cpu_accessible_am();
|
||||
|
||||
// Get Max Threads Per Multiprocessor
|
||||
|
||||
HsaNodeProperties node_prop = {0};
|
||||
if(HSAKMT_STATUS_SUCCESS == hsaKmtGetNodeProperties(node, &node_prop)) {
|
||||
uint32_t waves_per_cu = node_prop.MaxWavesPerSIMD;
|
||||
uint32_t simd_per_cu = node_prop.NumSIMDPerCU;
|
||||
prop-> maxThreadsPerMultiProcessor = prop->warpSize*waves_per_cu*simd_per_cu;
|
||||
}
|
||||
uint32_t max_waves_per_cu;
|
||||
err = hsa_agent_get_info(_hsaAgent,(hsa_agent_info_t) HSA_AMD_AGENT_INFO_MAX_WAVES_PER_CU, &max_waves_per_cu);
|
||||
DeviceErrorCheck(err);
|
||||
prop-> maxThreadsPerMultiProcessor = prop->warpSize*max_waves_per_cu;
|
||||
|
||||
// Get memory properties
|
||||
err = hsa_agent_iterate_regions(_hsaAgent, get_region_info, prop);
|
||||
|
||||
@@ -33,9 +33,7 @@ THE SOFTWARE.
|
||||
#error("This version of HIP requires a newer version of HCC.");
|
||||
#endif
|
||||
|
||||
#define USE_DISPATCH_HSA_KERNEL 1
|
||||
#define USE_IPC 0
|
||||
//
|
||||
|
||||
|
||||
//---
|
||||
@@ -383,26 +381,6 @@ public:
|
||||
size_t psize;
|
||||
};
|
||||
|
||||
class ihipFunction_t{
|
||||
public:
|
||||
ihipFunction_t(const char *name) {
|
||||
size_t nameSz = strlen(name);
|
||||
char *kernelName = (char*)malloc(nameSz);
|
||||
strncpy(kernelName, name, nameSz);
|
||||
_kernelName = kernelName;
|
||||
};
|
||||
|
||||
~ihipFunction_t() {
|
||||
if (_kernelName) {
|
||||
free((void*)_kernelName);
|
||||
_kernelName = NULL;
|
||||
};
|
||||
};
|
||||
public:
|
||||
const char *_kernelName;
|
||||
hsa_executable_symbol_t _kernelSymbol;
|
||||
uint64_t _kernel;
|
||||
};
|
||||
|
||||
class ihipModule_t {
|
||||
public:
|
||||
@@ -412,20 +390,7 @@ public:
|
||||
void *ptr;
|
||||
size_t size;
|
||||
|
||||
ihipModule_t() : executable(), object(), fileName(), ptr(nullptr), size(0), hipFunctionTable() {}
|
||||
~ihipModule_t() {
|
||||
for (int i = 0; i < hipFunctionTable.size(); ++i) {
|
||||
ihipFunction_t *func = hipFunctionTable[i];
|
||||
delete func;
|
||||
}
|
||||
hipFunctionTable.clear();
|
||||
}
|
||||
|
||||
void registerFunction(ihipFunction_t* func) {
|
||||
hipFunctionTable.push_back(func);
|
||||
}
|
||||
private:
|
||||
std::vector<ihipFunction_t*> hipFunctionTable;
|
||||
ihipModule_t() : executable(), object(), fileName(), ptr(nullptr), size(0) {}
|
||||
};
|
||||
|
||||
template <typename MUTEX_TYPE>
|
||||
|
||||
@@ -12,51 +12,6 @@ define linkonce_odr spir_func void @__threadfence_block() #1 {
|
||||
ret void
|
||||
}
|
||||
|
||||
define linkonce_odr spir_func i32 @__rocm_dp4a(i32 %in1, i32 %in2, i32 %in3) {
|
||||
%val1 = tail call i32 asm "v_mul_u32_u24_sdwa $0, $1, $2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:BYTE_0","=v,v,v"(i32 %in1, i32 %in2)
|
||||
%ret1 = add i32 %val1, %in3
|
||||
%val2 = tail call i32 asm "v_mul_u32_u24_sdwa $0, $1, $2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_1 src1_sel:BYTE_1","=v,v,v"(i32 %in1, i32 %in2)
|
||||
%ret2 = add i32 %ret1, %val2
|
||||
%val3 = tail call i32 asm "v_mul_u32_u24_sdwa $0, $1, $2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_2 src1_sel:BYTE_2","=v,v,v"(i32 %in1, i32 %in2)
|
||||
%ret3 = add i32 %val3, %ret2
|
||||
%val4 = tail call i32 asm "v_mul_u32_u24_sdwa $0, $1, $2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_3 src1_sel:BYTE_3","=v,v,v"(i32 %in1, i32 %in2)
|
||||
%ret4 = add i32 %val4, %ret3
|
||||
ret i32 %ret4
|
||||
}
|
||||
|
||||
define linkonce_odr spir_func i32 @__rocm_hfma(i32 %in1, i32 %in2, i32 %in3) {
|
||||
tail call void asm "v_mac_f16 $0, $1, $2","v,v,v"(i32 %in1, i32 %in2, i32 %in3)
|
||||
ret i32 %in3
|
||||
}
|
||||
|
||||
define linkonce_odr spir_func i32 @__rocm_hadd(i32 %in1, i32 %in2) {
|
||||
%val = tail call i32 asm "v_add_f16 $0, $1, $2","=v,v,v"(i32 %in1, i32 %in2)
|
||||
ret i32 %val
|
||||
}
|
||||
|
||||
define linkonce_odr spir_func half @__hip_hadd_gfx803(half %a, half %b) #1 {
|
||||
%val = tail call half asm "v_add_f16 $0, $1, $2","=v,v,v"(half %a, half %b)
|
||||
ret half %val
|
||||
}
|
||||
|
||||
define linkonce_odr spir_func half @__hip_hfma_gfx803(half %a, half %b, half %c) #1 {
|
||||
%val = tail call half asm "v_fma_f16 $0, $1, $2, $3","=v,v,v,v"(half %a, half %b, half %c)
|
||||
ret half %val
|
||||
}
|
||||
|
||||
define linkonce_odr spir_func half @__hip_hmul_gfx803(half %a, half %b) #1 {
|
||||
%val = tail call half asm "v_mul_f16 $0, $1, $2","=v,v,v"(half %a, half %b)
|
||||
ret half %val
|
||||
}
|
||||
|
||||
define linkonce_odr spir_func half @__hip_hsub_gfx803(half %a, half %b) #1 {
|
||||
%val = tail call half asm "v_sub_f16 $0, $1, $2","=v,v,v"(half %a, half %b)
|
||||
ret half %val
|
||||
}
|
||||
|
||||
define linkonce_odr spir_func i32 @__hip_hadd2_gfx803(i32 %a, i32 %b) #1 {
|
||||
%val = tail call i32 asm "v_add_f16_sdwa $0, $1, $2 dst_sel:WORD_0 dst_unused:UNUSED_PRESERVE src0_sel:WORD_0 src1_sel:WORD_0","=v,v,v"(i32 %a, i32 %b)
|
||||
ret i32 %val
|
||||
}
|
||||
|
||||
attributes #1 = { alwaysinline nounwind }
|
||||
|
||||
@@ -35,6 +35,16 @@ THE SOFTWARE.
|
||||
|
||||
//TODO Use Pool APIs from HCC to get memory regions.
|
||||
|
||||
#define CHECK_HSA(hsaStatus, hipStatus) \
|
||||
if (hsaStatus != HSA_STATUS_SUCCESS) {\
|
||||
return hipStatus;\
|
||||
}
|
||||
|
||||
#define CHECKLOG_HSA(hsaStatus, hipStatus) \
|
||||
if (hsaStatus != HSA_STATUS_SUCCESS) {\
|
||||
return ihipLogStatus(hipStatus);\
|
||||
}
|
||||
|
||||
namespace hipdrv {
|
||||
|
||||
hsa_status_t findSystemRegions(hsa_region_t region, void *data){
|
||||
@@ -103,6 +113,7 @@ uint64_t ElfSize(const void *emi){
|
||||
return total_size;
|
||||
}
|
||||
|
||||
|
||||
hipError_t hipModuleLoad(hipModule_t *module, const char *fname){
|
||||
HIP_INIT_API(module, fname);
|
||||
hipError_t ret = hipSuccess;
|
||||
@@ -154,15 +165,20 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname){
|
||||
}
|
||||
|
||||
status = hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, NULL, &(*module)->executable);
|
||||
if(status != HSA_STATUS_SUCCESS){
|
||||
return ihipLogStatus(hipErrorNotInitialized);
|
||||
}
|
||||
CHECKLOG_HSA(status, hipErrorNotInitialized);
|
||||
|
||||
status = hsa_executable_load_code_object((*module)->executable, agent, (*module)->object, NULL);
|
||||
CHECKLOG_HSA(status, hipErrorNotInitialized);
|
||||
|
||||
status = hsa_executable_freeze((*module)->executable, NULL);
|
||||
CHECKLOG_HSA(status, hipErrorNotInitialized);
|
||||
}
|
||||
}
|
||||
|
||||
return ihipLogStatus(ret);
|
||||
}
|
||||
|
||||
|
||||
hipError_t hipModuleUnload(hipModule_t hmod)
|
||||
{
|
||||
// TODO - improve this synchronization so it is thread-safe.
|
||||
@@ -189,7 +205,7 @@ hipError_t hipModuleUnload(hipModule_t hmod)
|
||||
return ihipLogStatus(ret);
|
||||
}
|
||||
|
||||
hipError_t ihipModuleGetFunction(hipFunction_t *func, hipModule_t hmod, const char *name){
|
||||
hipError_t ihipModuleGetSymbol(hipFunction_t *func, hipModule_t hmod, const char *name){
|
||||
auto ctx = ihipGetTlsDefaultCtx();
|
||||
hipError_t ret = hipSuccess;
|
||||
|
||||
@@ -201,31 +217,33 @@ hipError_t ihipModuleGetFunction(hipFunction_t *func, hipModule_t hmod, const ch
|
||||
ret = hipErrorInvalidContext;
|
||||
|
||||
}else{
|
||||
*func = new ihipFunction_t(name);
|
||||
hmod->registerFunction(*func);
|
||||
int deviceId = ctx->getDevice()->_deviceId;
|
||||
ihipDevice_t *currentDevice = ihipGetDevice(deviceId);
|
||||
hsa_agent_t gpuAgent = (hsa_agent_t)currentDevice->_hsaAgent;
|
||||
|
||||
hsa_status_t status;
|
||||
status = hsa_executable_load_code_object(hmod->executable, gpuAgent, hmod->object, NULL);
|
||||
if(status != HSA_STATUS_SUCCESS){
|
||||
return ihipLogStatus(hipErrorNotInitialized);
|
||||
}
|
||||
|
||||
status = hsa_executable_freeze(hmod->executable, NULL);
|
||||
status = hsa_executable_get_symbol(hmod->executable, NULL, name, gpuAgent, 0, &(*func)->_kernelSymbol);
|
||||
hsa_executable_symbol_t symbol;
|
||||
status = hsa_executable_get_symbol(hmod->executable, NULL, name, gpuAgent, 0, &symbol);
|
||||
if(status != HSA_STATUS_SUCCESS){
|
||||
return ihipLogStatus(hipErrorNotFound);
|
||||
}
|
||||
|
||||
status = hsa_executable_symbol_get_info((*func)->_kernelSymbol,
|
||||
status = hsa_executable_symbol_get_info(symbol,
|
||||
HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT,
|
||||
&(*func)->_kernel);
|
||||
&func->_object);
|
||||
CHECK_HSA(status, hipErrorNotFound);
|
||||
|
||||
if(status != HSA_STATUS_SUCCESS){
|
||||
return ihipLogStatus(hipErrorNotFound);
|
||||
}
|
||||
status = hsa_executable_symbol_get_info(symbol,
|
||||
HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE,
|
||||
&func->_groupSegmentSize);
|
||||
CHECK_HSA(status, hipErrorNotFound);
|
||||
|
||||
status = hsa_executable_symbol_get_info(symbol,
|
||||
HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE,
|
||||
&func->_privateSegmentSize);
|
||||
CHECK_HSA(status, hipErrorNotFound);
|
||||
|
||||
strncpy(func->_name, name, sizeof(func->_name));
|
||||
}
|
||||
return ihipLogStatus(ret);
|
||||
}
|
||||
@@ -234,7 +252,7 @@ hipError_t ihipModuleGetFunction(hipFunction_t *func, hipModule_t hmod, const ch
|
||||
hipError_t hipModuleGetFunction(hipFunction_t *hfunc, hipModule_t hmod,
|
||||
const char *name){
|
||||
HIP_INIT_API(hfunc, hmod, name);
|
||||
return ihipModuleGetFunction(hfunc, hmod, name);
|
||||
return ihipModuleGetSymbol(hfunc, hmod, name);
|
||||
}
|
||||
|
||||
|
||||
@@ -274,31 +292,14 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
|
||||
return ihipLogStatus(hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
uint32_t groupSegmentSize;
|
||||
hsa_status_t status = hsa_executable_symbol_get_info(f->_kernelSymbol,
|
||||
HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE,
|
||||
&groupSegmentSize);
|
||||
if(status != HSA_STATUS_SUCCESS){
|
||||
return ihipLogStatus(hipErrorNotFound);
|
||||
}
|
||||
|
||||
uint32_t privateSegmentSize;
|
||||
status = hsa_executable_symbol_get_info(f->_kernelSymbol,
|
||||
HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE,
|
||||
&privateSegmentSize);
|
||||
if(status != HSA_STATUS_SUCCESS){
|
||||
return ihipLogStatus(hipErrorNotFound);
|
||||
}
|
||||
privateSegmentSize += sharedMemBytes;
|
||||
|
||||
|
||||
/*
|
||||
Kernel argument preparation.
|
||||
*/
|
||||
grid_launch_parm lp;
|
||||
hStream = ihipPreLaunchKernel(hStream, 0, 0, &lp, f->_kernelName);
|
||||
grid_launch_parm lp; // TODO - dummy arg but values are printed during debug.
|
||||
hStream = ihipPreLaunchKernel(hStream, 0, 0, &lp, f._name);
|
||||
|
||||
#if USE_DISPATCH_HSA_KERNEL
|
||||
|
||||
hsa_kernel_dispatch_packet_t aql;
|
||||
|
||||
@@ -313,9 +314,9 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
|
||||
aql.grid_size_x = blockDimX * gridDimX;
|
||||
aql.grid_size_y = blockDimY * gridDimY;
|
||||
aql.grid_size_z = blockDimZ * gridDimZ;
|
||||
aql.group_segment_size = groupSegmentSize;
|
||||
aql.private_segment_size = privateSegmentSize;
|
||||
aql.kernel_object = f->_kernel;
|
||||
aql.group_segment_size = f._groupSegmentSize + sharedMemBytes;
|
||||
aql.private_segment_size = f._privateSegmentSize;
|
||||
aql.kernel_object = f._object;
|
||||
aql.setup = 3 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS;
|
||||
aql.header = (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) |
|
||||
(1 << HSA_PACKET_HEADER_BARRIER) |
|
||||
@@ -323,69 +324,9 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
|
||||
(HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);
|
||||
|
||||
lp.av->dispatch_hsa_kernel(&aql, config[1] /* kernarg*/, kernArgSize);
|
||||
#else
|
||||
|
||||
/*
|
||||
Create signal
|
||||
*/
|
||||
|
||||
hsa_signal_t signal;
|
||||
status = hsa_signal_create(1, 0, NULL, &signal);
|
||||
if(status != HSA_STATUS_SUCCESS){
|
||||
return ihipLogStatus(hipErrorNotFound);
|
||||
}
|
||||
|
||||
/*
|
||||
Allocate kernarg
|
||||
*/
|
||||
void *kern = nullptr;
|
||||
|
||||
hsa_amd_memory_pool_t *pool = reinterpret_cast<hsa_amd_memory_pool_t*>(lp.av->get_hsa_kernarg_region());
|
||||
status = hsa_amd_memory_pool_allocate(*pool, kernArgSize, 0, &kern);
|
||||
if(status != HSA_STATUS_SUCCESS){
|
||||
return ihipLogStatus(hipErrorNotFound);
|
||||
}
|
||||
status = hsa_amd_agents_allow_access(1, (hsa_agent_t*)lp.av->get_hsa_agent(), 0, kern);
|
||||
if(status != HSA_STATUS_SUCCESS){
|
||||
return ihipLogStatus(hipErrorNotFound);
|
||||
}
|
||||
memcpy(kern, config[1], kernArgSize);
|
||||
|
||||
|
||||
/*
|
||||
Launch AQL packet
|
||||
*/
|
||||
hStream->launchModuleKernel(*lp.av, signal, blockDimX, blockDimY, blockDimZ,
|
||||
gridDimX, gridDimY, gridDimZ, groupSegmentSize, privateSegmentSize, kern, kernArgSize, f->_kernel);
|
||||
|
||||
|
||||
/*
|
||||
Wait for signal
|
||||
*/
|
||||
|
||||
hsa_signal_value_t value = hsa_signal_wait_acquire(signal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_BLOCKED);
|
||||
|
||||
/*
|
||||
Destroy kernarg
|
||||
*/
|
||||
status = hsa_amd_memory_pool_free(kern);
|
||||
if(status != HSA_STATUS_SUCCESS){
|
||||
return ihipLogStatus(hipErrorNotFound);
|
||||
}
|
||||
|
||||
/*
|
||||
Destroy the signal
|
||||
*/
|
||||
status = hsa_signal_destroy(signal);
|
||||
if(status != HSA_STATUS_SUCCESS){
|
||||
return ihipLogStatus(hipErrorNotFound);
|
||||
}
|
||||
|
||||
#endif // USE_DISPATCH_HSA_KERNEL
|
||||
|
||||
|
||||
ihipPostLaunchKernel(f->_kernelName, hStream, lp);
|
||||
|
||||
ihipPostLaunchKernel(f._name, hStream, lp);
|
||||
}
|
||||
|
||||
return ihipLogStatus(ret);
|
||||
@@ -405,9 +346,9 @@ hipError_t hipModuleGetGlobal(hipDeviceptr_t *dptr, size_t *bytes,
|
||||
}
|
||||
else{
|
||||
hipFunction_t func;
|
||||
ihipModuleGetFunction(&func, hmod, name);
|
||||
ihipModuleGetSymbol(&func, hmod, name);
|
||||
*bytes = PrintSymbolSizes(hmod->ptr, name) + sizeof(amd_kernel_code_t);
|
||||
*dptr = reinterpret_cast<void*>(func->_kernel);
|
||||
*dptr = reinterpret_cast<void*>(func._object);
|
||||
return ihipLogStatus(ret);
|
||||
}
|
||||
}
|
||||
@@ -419,7 +360,7 @@ hipError_t hipModuleLoadData(hipModule_t *module, const void *image)
|
||||
hipError_t ret = hipSuccess;
|
||||
if(image == NULL || module == NULL){
|
||||
return ihipLogStatus(hipErrorNotInitialized);
|
||||
}else{
|
||||
} else {
|
||||
auto ctx = ihipGetTlsDefaultCtx();
|
||||
*module = new ihipModule_t;
|
||||
int deviceId = ctx->getDevice()->_deviceId;
|
||||
@@ -452,9 +393,13 @@ hipError_t hipModuleLoadData(hipModule_t *module, const void *image)
|
||||
}
|
||||
|
||||
status = hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, NULL, &(*module)->executable);
|
||||
if(status != HSA_STATUS_SUCCESS){
|
||||
return ihipLogStatus(hipErrorNotInitialized);
|
||||
}
|
||||
CHECKLOG_HSA(status, hipErrorNotInitialized);
|
||||
|
||||
status = hsa_executable_load_code_object((*module)->executable, agent, (*module)->object, NULL);
|
||||
CHECKLOG_HSA(status, hipErrorNotInitialized);
|
||||
|
||||
status = hsa_executable_freeze((*module)->executable, NULL);
|
||||
CHECKLOG_HSA(status, hipErrorNotInitialized);
|
||||
}
|
||||
return ihipLogStatus(ret);
|
||||
}
|
||||
|
||||
@@ -72,6 +72,15 @@ inline std::string ToString(hipEvent_t v)
|
||||
return ss.str();
|
||||
};
|
||||
|
||||
// hipEvent_t specialization. TODO - maybe add an event ID for debug?
|
||||
template <>
|
||||
inline std::string ToString(hipFunction_t v)
|
||||
{
|
||||
std::ostringstream ss;
|
||||
ss << "0x" << std::hex << v._object;
|
||||
return ss.str();
|
||||
};
|
||||
|
||||
|
||||
|
||||
// hipStream_t
|
||||
|
||||
@@ -27,7 +27,8 @@ int main(){
|
||||
hipInit(0);
|
||||
hipDevice_t device;
|
||||
hipDeviceGet(&device, 0);
|
||||
char pciBusId[100];
|
||||
char pciBusId[10];
|
||||
memset(pciBusId,0,10);
|
||||
hipDeviceGetPCIBusId(pciBusId,100,device);
|
||||
printf("PCI Bus ID= %s\n",pciBusId);
|
||||
return 0;
|
||||
|
||||
Ссылка в новой задаче
Block a user