AI/rocm-systems
2
0
انشعاب 0
کد مسائل تقاضاهای واکشی Actions Packages پروژه‌ها انتشارها دانشنامه فعالیت

Merge 'master' into 'amd-master'

فهرست منبع 
Change-Id: I81a373b1a9c149e96ede6c72a6c24d99171f8821
This commit is contained in:
Jenkins
2019-06-23 11:47:23 -04:00
والد b3b6ff50cb 00f25833a9
کامیت 7ac69da23e
18فایلهای تغییر یافته به همراه427 افزوده شده و 81 حذف شده
دانلود پرونده وصله دانلود فایل تغییرات diff Expand all files Collapse all files
bin/hipcc
+3
مشاهده پرونده
@@ -118,6 +118,9 @@ if ($HIP_RUNTIME eq "VDI" and !defined $HIP_VDI_HOME) {
if (defined $HIP_VDI_HOME) {
if (!defined $HIP_CLANG_PATH and (-e "$HIP_VDI_HOME/bin/clang" or -e "$HIP_VDI_HOME/bin/clang.exe")) {
$HIP_CLANG_PATH = "$HIP_VDI_HOME/bin";
}
# With HIP_VDI_HOME defined, assume the installation of clang components, including headers.
if (!defined $HIP_CLANG_INCLUDE_PATH) {
$HIP_CLANG_INCLUDE_PATH = "$HIP_VDI_HOME/include/clang";
}
if (!defined $DEVICE_LIB_PATH and -e "$HIP_VDI_HOME/lib/bitcode") {
hip_prof_gen.py
+28 -27
مشاهده پرونده
@@ -13,7 +13,7 @@ line_num = -1
# Verbose message
def message(msg):
if verbose: print >>sys.stdout, msg
if verbose: sys.stdout.write(msg + '\n')
# Fatal error termination
def error(msg):
@@ -24,8 +24,8 @@ def error(msg):
else:
msg = " Warning: " + msg
print >>sys.stdout, msg
print >>sys.stderr, sys.argv[0] + msg
sys.stdout.write(msg + '\n')
sys.stderr.write(sys.argv[0] + msg +'\n')
def fatal(msg):
error(msg)
@@ -301,8 +301,6 @@ def generate_prof_header(f, api_map, opts_map):
f.write('// automatically generated sources\n')
f.write('#ifndef _HIP_PROF_STR_H\n');
f.write('#define _HIP_PROF_STR_H\n');
f.write('#include <sstream>\n');
f.write('#include <string>\n');
# Generating dummy macro for non-public API
f.write('\n// Dummy API primitives\n')
@@ -383,27 +381,30 @@ def generate_prof_header(f, api_map, opts_map):
f.write('#define INIT_CB_ARGS_DATA(cb_id, cb_data) INIT_##cb_id##_CB_ARGS_DATA(cb_data)\n')
# Generating the method for the API string, name and parameters
f.write('\n')
f.write('#if 0\n')
f.write('// HIP API string method, method name and parameters\n')
f.write('const char* hipApiString(hip_api_id_t id, const hip_api_data_t* data) {\n')
f.write(' std::ostringstream oss;\n')
f.write(' switch (id) {\n')
for name, args in api_map.items():
f.write(' case HIP_API_ID_' + name + ':\n')
f.write(' oss << "' + name + '("')
for ind in range(0, len(args)):
arg_tuple = args[ind]
arg_name = arg_tuple[1]
if ind != 0: f.write(' << ","')
f.write('\n << " ' + arg_name + '=" << data->args.' + name + '.' + arg_name)
f.write('\n << ")";\n')
f.write(' break;\n')
f.write(' default: oss << "unknown";\n')
f.write(' };\n')
f.write(' return strdup(oss.str().c_str());\n')
f.write('};\n')
f.write('#endif\n')
if False:
f.write('\n')
f.write('#if 0\n')
f.write('#include <sstream>\n');
f.write('#include <string>\n');
f.write('// HIP API string method, method name and parameters\n')
f.write('const char* hipApiString(hip_api_id_t id, const hip_api_data_t* data) {\n')
f.write(' std::ostringstream oss;\n')
f.write(' switch (id) {\n')
for name, args in api_map.items():
f.write(' case HIP_API_ID_' + name + ':\n')
f.write(' oss << "' + name + '("')
for ind in range(0, len(args)):
arg_tuple = args[ind]
arg_name = arg_tuple[1]
if ind != 0: f.write(' << ","')
f.write('\n << " ' + arg_name + '=" << data->args.' + name + '.' + arg_name)
f.write('\n << ")";\n')
f.write(' break;\n')
f.write(' default: oss << "unknown";\n')
f.write(' };\n')
f.write(' return strdup(oss.str().c_str());\n')
f.write('};\n')
f.write('#endif\n')
f.write('#endif // _HIP_PROF_STR_H\n');
@@ -451,7 +452,7 @@ parse_api(api_hfile, api_map)
parse_src(api_map, src_dir, src_pat, opts_map)
# Checking for non-conformant APIs
for name in opts_map.keys():
for name in list(opts_map.keys()):
m = re.match(r'\.(\S*)', name)
if m:
message("Init missing: " + m.group(1))
hipify-clang/src/main.cpp
+1 -1
مشاهده پرونده
@@ -234,7 +234,7 @@ int main(int argc, const char **argv) {
argc++;
}
llcompat::PrintStackTraceOnErrorSignal();
ct::CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::Optional);
ct::CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::ZeroOrMore);
std::vector<std::string> fileSources = OptionsParser.getSourcePathList();
if (fileSources.empty() && !GeneratePerl && !GeneratePython) {
llvm::errs() << "\n" << sHipify << sError << "Must specify at least 1 positional argument for source file." << "\n";
include/hip/hcc_detail/functional_grid_launch.hpp
+16
مشاهده پرونده
@@ -135,6 +135,22 @@ void hipLaunchKernelGGLImpl(
}
} // Namespace hip_impl.
template <typename F>
inline
void hipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* blockSize,
F kernel, size_t dynSharedMemPerBlk, uint32_t blockSizeLimit) {
using namespace hip_impl;
hip_impl::hip_init();
auto f = get_program_state().kernel_descriptor(reinterpret_cast<std::uintptr_t>(kernel),
target_agent(0));
hipOccupancyMaxPotentialBlockSize(gridSize, blockSize, f,
dynSharedMemPerBlk, blockSizeLimit);
}
template <typename... Args, typename F = void (*)(Args...)>
inline
void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
include/hip/hcc_detail/hip_runtime.h
+5 -5
مشاهده پرونده
@@ -270,14 +270,14 @@ static inline __device__ void printf(const char* format, All... all) {}
#if defined __HCC_CPP__
extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block,
grid_launch_parm* lp, const char* kernelNameStr);
grid_launch_parm* lp, const char* kernelNameStr, bool lockAcquired = 0);
extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block,
grid_launch_parm* lp, const char* kernelNameStr);
grid_launch_parm* lp, const char* kernelNameStr, bool lockAcquired = 0);
extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block,
grid_launch_parm* lp, const char* kernelNameStr);
grid_launch_parm* lp, const char* kernelNameStr, bool lockAcquired = 0);
extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block,
grid_launch_parm* lp, const char* kernelNameStr);
extern void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, grid_launch_parm& lp);
grid_launch_parm* lp, const char* kernelNameStr, bool lockAcquired = 0);
extern void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, grid_launch_parm& lp, bool unlockPostponed = 0);
#if GENERIC_GRID_LAUNCH == 0
//#warning "Original hipLaunchKernel defined"
include/hip/hcc_detail/hip_runtime_api.h
+17 -5
مشاهده پرونده
@@ -2744,6 +2744,21 @@ hipError_t hipLaunchCooperativeKernel(const void* f, dim3 gridDim, dim3 blockDim
hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList,
int numDevices, unsigned int flags);
/**
* @brief determine the grid and block sizes to achieves maximum occupancy for a kernel
*
* @param [out] gridSize minimum grid size for maximum potential occupancy
* @param [out] blockSize block size for maximum potential occupancy
* @param [in] f kernel to launch
* @param [in] dynSharedMemPerBlk dynamic shared memory usage (in bytes) intended for each block
* @param [in] blockSizeLimit the maximum block size for the kernel, use 0 for no limit
*
* @returns hipSuccess, hipInvalidDevice, hipErrorInvalidValue
*/
hipError_t hipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* blockSize,
hipFunction_t f, size_t dynSharedMemPerBlk,
uint32_t blockSizeLimit);
/**
* @brief Returns occupancy for a device function.
*
@@ -2782,7 +2797,6 @@ hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList,
int numDevices, unsigned int flags);
// doxygen end Version Management
/**
* @}
@@ -2973,9 +2987,7 @@ hipError_t hipLaunchByPtr(const void* func);
} /* extern "c" */
#endif
#ifdef __cplusplus
#include <hip/hcc_detail/hip_prof_api.h>
#endif
#include <hip/hcc_detail/hip_prof_str.h>
#ifdef __cplusplus
extern "C" {
@@ -2987,7 +2999,7 @@ hipError_t hipRegisterApiCallback(uint32_t id, void* fun, void* arg);
hipError_t hipRemoveApiCallback(uint32_t id);
hipError_t hipRegisterActivityCallback(uint32_t id, void* fun, void* arg);
hipError_t hipRemoveActivityCallback(uint32_t id);
static inline const char* hipApiName(const uint32_t& id) { return hip_api_name(id); }
const char* hipApiName(uint32_t id);
const char* hipKernelNameRef(const hipFunction_t f);
#ifdef __cplusplus
} /* extern "C" */
include/hip/hcc_detail/program_state.hpp
+1
مشاهده پرونده
@@ -77,6 +77,7 @@ class kernargs_size_align {
public:
std::size_t size(std::size_t n) const;
std::size_t alignment(std::size_t n) const;
const void* getHandle() const {return handle;};
private:
const void* handle;
friend kernargs_size_align program_state::get_kernargs_size_align(std::uintptr_t);
samples/0_Intro/module_api/Makefile
+3 -3
مشاهده پرونده
@@ -5,7 +5,7 @@ endif
HIPCC=$(HIP_PATH)/bin/hipcc
HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --compiler)
all: vcpy_kernel.code runKernel.hip.out launchKernelHcc.hip.out
all: vcpy_kernel.code runKernel.hip.out launchKernelHcc.hip.out defaultDriver.hip.out
runKernel.hip.out: runKernel.cpp
$(HIPCC) $(HIPCC_FLAGS) $< -o $@
@@ -13,8 +13,8 @@ runKernel.hip.out: runKernel.cpp
launchKernelHcc.hip.out: launchKernelHcc.cpp
$(HIPCC) $(HIPCC_FLAGS) $< -o $@
#defaultDriver.hip.out: defaultDriver.cpp
# $(HIPCC) $(HIPCC_FLAGS) $< -o $@
defaultDriver.hip.out: defaultDriver.cpp
$(HIPCC) $(HIPCC_FLAGS) $< -o $@
vcpy_kernel.code: vcpy_kernel.cpp
$(HIPCC) --genco $(GENCO_FLAGS) $^ -o $@
samples/0_Intro/module_api/defaultDriver.cpp
+13 -15
مشاهده پرونده
@@ -21,7 +21,6 @@ THE SOFTWARE.
*/
#include "hip/hip_runtime.h"
#include "hip/hip_runtime_api.h"
#include <iostream>
#include <fstream>
#include <vector>
@@ -29,20 +28,18 @@ THE SOFTWARE.
#define LEN 64
#define SIZE LEN << 2
#define fileName "test.co"
#define kernel_name "vadd"
#define fileName "vcpy_kernel.code"
#define kernel_name "hello_world"
int main() {
float *A, *B, *C;
hipDeviceptr_t Ad, Bd, Cd;
float *A, *B;
hipDeviceptr_t Ad, Bd;
A = new float[LEN];
B = new float[LEN];
C = new float[LEN];
for (uint32_t i = 0; i < LEN; i++) {
A[i] = i * 1.0f;
B[i] = 1.0f;
C[i] = 0.0f;
B[i] = 0.0f;
}
hipInit(0);
@@ -53,28 +50,25 @@ int main() {
hipMalloc((void**)&Ad, SIZE);
hipMalloc((void**)&Bd, SIZE);
hipMalloc((void**)&Cd, SIZE);
hipMemcpyHtoD(Ad, A, SIZE);
hipMemcpyHtoD(Bd, B, SIZE);
hipMemcpyHtoD(Cd, C, SIZE);
hipModule_t Module;
hipFunction_t Function;
hipModuleLoad(&Module, fileName);
hipModuleGetFunction(&Function, Module, kernel_name);
int n = LEN;
void* args[4] = {&Ad, &Bd, &Cd, &n};
void* args[2] = {&Ad, &Bd};
hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, args, nullptr);
hipMemcpyDtoH(C, Cd, SIZE);
hipMemcpyDtoH(B, Bd, SIZE);
int mismatchCount = 0;
for (uint32_t i = 0; i < LEN; i++) {
if (A[i] + B[i] != C[i]) {
if (A[i] != B[i]) {
mismatchCount++;
std::cout << "error: mismatch " << A[i] << " + " << B[i] << " != " << C[i] << std::endl;
std::cout << "error: mismatch " << A[i] << " != " << C[i] << std::endl;
}
}
@@ -84,6 +78,10 @@ int main() {
std::cout << "FAILED!\n";
};
hipFree(Ad);
hipFree(Bd);
delete A;
delete B;
hipCtxDestroy(context);
return 0;
}
samples/0_Intro/module_api/launchKernelHcc.cpp
+4
مشاهده پرونده
@@ -105,6 +105,10 @@ int main() {
std::cout << "FAILED!\n";
};
hipFree(Ad);
hipFree(Bd);
delete A;
delete B;
hipCtxDestroy(context);
return 0;
}
samples/0_Intro/module_api/runKernel.cpp
+4
مشاهده پرونده
@@ -97,6 +97,10 @@ int main() {
std::cout << "FAILED!\n";
};
hipFree(Ad);
hipFree(Bd);
delete A;
delete B;
hipCtxDestroy(context);
return 0;
}
samples/0_Intro/module_api_global/runKernel.cpp
+4
مشاهده پرونده
@@ -145,6 +145,10 @@ int main() {
};
}
hipFree(Ad);
hipFree(Bd);
delete A;
delete B;
hipCtxDestroy(context);
return 0;
}
src/hip_hcc.cpp
+31 -18
مشاهده پرونده
@@ -404,7 +404,7 @@ LockedAccessor_StreamCrit_t ihipStream_t::lockopen_preKernelCommand() {
//---
// Must be called after kernel finishes, this releases the lock on the stream so other commands can
// submit.
void ihipStream_t::lockclose_postKernelCommand(const char* kernelName, hc::accelerator_view* av) {
void ihipStream_t::lockclose_postKernelCommand(const char* kernelName, hc::accelerator_view* av, bool unlockPostponed) {
bool blockThisKernel = false;
if (!g_hipLaunchBlockingKernels.empty()) {
@@ -426,7 +426,10 @@ void ihipStream_t::lockclose_postKernelCommand(const char* kernelName, hc::accel
kernelName);
}
_criticalData.unlock(); // paired with lock from lockopen_preKernelCommand.
// if unlockPostponed is true then this stream will be unlocked later (e.g., see hipExtLaunchMultiKernelMultiDevice for a sample call)
if (!unlockPostponed) {
_criticalData.unlock(); // paired with lock from lockopen_preKernelCommand.
}
};
@@ -1493,7 +1496,7 @@ void ihipStreamCallbackHandler(ihipStreamCallback_t* cb) {
//
// If stream==NULL synchronize appropriately with other streams and return the default av for the
// device. If stream is valid, return the AV to use.
hipStream_t ihipSyncAndResolveStream(hipStream_t stream) {
hipStream_t ihipSyncAndResolveStream(hipStream_t stream, bool lockAcquired) {
if (stream == hipStreamNull) {
// Submitting to NULL stream, call locked_syncDefaultStream to wait for all other streams:
ihipCtx_t* ctx = ihipGetTlsDefaultCtx();
@@ -1535,9 +1538,14 @@ hipStream_t ihipSyncAndResolveStream(hipStream_t stream) {
if (needGatherMarker) {
// ensure any commands sent to this stream wait on the NULL stream before
// continuing
LockedAccessor_StreamCrit_t thisStreamCrit(stream->criticalData());
// TODO - could be "noret" version of create_blocking_marker
thisStreamCrit->_av.create_blocking_marker(dcf, hc::accelerator_scope);
if (!lockAcquired) {
LockedAccessor_StreamCrit_t thisStreamCrit(stream->criticalData());
// TODO - could be "noret" version of create_blocking_marker
thisStreamCrit->_av.create_blocking_marker(dcf, hc::accelerator_scope);
} else {
// this stream is already locked (e.g., call from hipExtLaunchMultiKernelMultiDevice)
stream->criticalData()._av.create_blocking_marker(dcf, hc::accelerator_scope);
}
tprintf(
DB_SYNC,
" %s adding marker to wait for freshly recorded default-stream marker \n",
@@ -1578,8 +1586,8 @@ void ihipPrintKernelLaunch(const char* kernelName, const grid_launch_parm* lp,
// Called just before a kernel is launched from hipLaunchKernel.
// Allows runtime to track some information about the stream.
hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_launch_parm* lp,
const char* kernelNameStr) {
stream = ihipSyncAndResolveStream(stream);
const char* kernelNameStr, bool lockAcquired) {
stream = ihipSyncAndResolveStream(stream, lockAcquired);
lp->grid_dim.x = grid.x;
lp->grid_dim.y = grid.y;
lp->grid_dim.z = grid.z;
@@ -1589,8 +1597,13 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_
lp->barrier_bit = barrier_bit_queue_default;
lp->launch_fence = -1;
auto crit = stream->lockopen_preKernelCommand();
lp->av = &(crit->_av);
if (!lockAcquired) {
auto crit = stream->lockopen_preKernelCommand();
lp->av = &(crit->_av);
} else {
// this stream is already locked (e.g., call from hipExtLaunchMultiKernelMultiDevice)
lp->av = &(stream->criticalData()._av);
}
lp->cf = nullptr;
ihipPrintKernelLaunch(kernelNameStr, lp, stream);
@@ -1599,30 +1612,30 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_
hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, grid_launch_parm* lp,
const char* kernelNameStr) {
return ihipPreLaunchKernel(stream, dim3(grid), block, lp, kernelNameStr);
const char* kernelNameStr, bool lockAcquired) {
return ihipPreLaunchKernel(stream, dim3(grid), block, lp, kernelNameStr, lockAcquired);
}
hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, grid_launch_parm* lp,
const char* kernelNameStr) {
return ihipPreLaunchKernel(stream, grid, dim3(block), lp, kernelNameStr);
const char* kernelNameStr, bool lockAcquired) {
return ihipPreLaunchKernel(stream, grid, dim3(block), lp, kernelNameStr, lockAcquired);
}
hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, grid_launch_parm* lp,
const char* kernelNameStr) {
return ihipPreLaunchKernel(stream, dim3(grid), dim3(block), lp, kernelNameStr);
const char* kernelNameStr, bool lockAcquired) {
return ihipPreLaunchKernel(stream, dim3(grid), dim3(block), lp, kernelNameStr, lockAcquired);
}
//---
// Called after kernel finishes execution.
// This releases the lock on the stream.
void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, grid_launch_parm& lp) {
void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, grid_launch_parm& lp, bool unlockPostponed) {
tprintf(DB_SYNC, "ihipPostLaunchKernel, unlocking stream\n");
stream->lockclose_postKernelCommand(kernelName, lp.av);
stream->lockclose_postKernelCommand(kernelName, lp.av, unlockPostponed);
if (HIP_PROFILE_API) {
MARKER_END();
}
src/hip_hcc_internal.h
+3 -2
مشاهده پرونده
@@ -30,6 +30,7 @@ THE SOFTWARE.
#include "hsa/hsa_ext_amd.h"
#include "hip/hip_runtime.h"
#include "hip_prof_api.h"
#include "hip_util.h"
#include "env.h"
@@ -549,7 +550,7 @@ class ihipStream_t {
// Member functions that begin with locked_ are thread-safe accessors - these acquire / release
// the critical mutex.
LockedAccessor_StreamCrit_t lockopen_preKernelCommand();
void lockclose_postKernelCommand(const char* kernelName, hc::accelerator_view* av);
void lockclose_postKernelCommand(const char* kernelName, hc::accelerator_view* av, bool unlockNotNeeded = 0);
void locked_wait();
@@ -951,7 +952,7 @@ hipError_t hipModuleGetFunctionEx(hipFunction_t* hfunc, hipModule_t hmod,
const char* name, hsa_agent_t *agent);
hipStream_t ihipSyncAndResolveStream(hipStream_t);
hipStream_t ihipSyncAndResolveStream(hipStream_t, bool lockAcquired = 0);
hipError_t ihipStreamSynchronize(hipStream_t stream);
void ihipStreamCallbackHandler(ihipStreamCallback_t* cb);
src/hip_intercept.cpp
+5 -1
مشاهده پرونده
@@ -21,7 +21,7 @@ THE SOFTWARE.
*/
#include "hip/hip_runtime.h"
#include "hip/hcc_detail/hip_prof_api.h"
#include "hip_prof_api.h"
// HIP API callback/activity
@@ -47,3 +47,7 @@ hipError_t hipRegisterActivityCallback(uint32_t id, void* fun, void* arg) {
hipError_t hipRemoveActivityCallback(uint32_t id) {
return callbacks_table.set_activity(id, NULL, NULL) ? hipSuccess : hipErrorInvalidValue;
}
const char* hipApiName(uint32_t id) {
return hip_api_name(id);
}
src/hip_module.cpp
+220 -4
مشاهده پرونده
@@ -119,6 +119,8 @@ string ToString(hipFunction_t v) {
const std::string& FunctionSymbol(const hipFunction_t f) { return f->_name; };
extern hipError_t ihipGetDeviceProperties(hipDeviceProp_t* props, int device);
#define CHECK_HSA(hsaStatus, hipStatus) \
if (hsaStatus != HSA_STATUS_SUCCESS) { \
return hipStatus; \
@@ -148,7 +150,7 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
uint32_t localWorkSizeX, uint32_t localWorkSizeY,
uint32_t localWorkSizeZ, size_t sharedMemBytes,
hipStream_t hStream, void** kernelParams, void** extra,
hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags) {
hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags, bool isStreamLocked = 0) {
using namespace hip_impl;
auto ctx = ihipGetTlsDefaultCtx();
@@ -204,8 +206,7 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
sharedMemBytes; // TODO - this should be part of preLaunchKernel.
hStream = ihipPreLaunchKernel(
hStream, dim3(globalWorkSizeX/localWorkSizeX, globalWorkSizeY/localWorkSizeY, globalWorkSizeZ/localWorkSizeZ),
dim3(localWorkSizeX, localWorkSizeY, localWorkSizeZ), &lp, f->_name.c_str());
dim3(localWorkSizeX, localWorkSizeY, localWorkSizeZ), &lp, f->_name.c_str(), isStreamLocked);
hsa_kernel_dispatch_packet_t aql;
@@ -270,7 +271,9 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
stopEvent->attachToCompletionFuture(&cf, hStream, hipEventTypeStopCommand);
}
ihipPostLaunchKernel(f->_name.c_str(), hStream, lp);
ihipPostLaunchKernel(f->_name.c_str(), hStream, lp, isStreamLocked);
}
return ret;
@@ -313,6 +316,75 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, 0));
}
hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList,
int numDevices, unsigned int flags) {
hipError_t result;
if ((numDevices > g_deviceCnt) || (launchParamsList == nullptr)) {
return hipErrorInvalidValue;
}
hipFunction_t* kds = reinterpret_cast<hipFunction_t*>(malloc(sizeof(hipFunction_t) * numDevices));
if (kds == nullptr) {
return hipErrorNotInitialized;
}
// prepare all kernel descriptors for each device as all streams will be locked in the next loop
for (int i = 0; i < numDevices; ++i) {
const hipLaunchParams& lp = launchParamsList[i];
if (lp.stream == nullptr) {
free(kds);
return hipErrorNotInitialized;
}
kds[i] = hip_impl::get_program_state().kernel_descriptor(reinterpret_cast<std::uintptr_t>(lp.func),
hip_impl::target_agent(lp.stream));
if (kds[i] == nullptr) {
free(kds);
return hipErrorInvalidValue;
}
hip_impl::kernargs_size_align kargs = hip_impl::get_program_state().get_kernargs_size_align(
reinterpret_cast<std::uintptr_t>(lp.func));
kds[i]->_kernarg_layout = *reinterpret_cast<const std::vector<std::pair<std::size_t, std::size_t>>*>(
kargs.getHandle());
}
// lock all streams before launching kernels to each device
for (int i = 0; i < numDevices; ++i) {
LockedAccessor_StreamCrit_t streamCrit(launchParamsList[i].stream->criticalData(), false);
#if (__hcc_workweek__ >= 19213)
streamCrit->_av.acquire_locked_hsa_queue();
#endif
}
// launch kernels for each device
for (int i = 0; i < numDevices; ++i) {
const hipLaunchParams& lp = launchParamsList[i];
result = ihipModuleLaunchKernel(kds[i],
lp.gridDim.x * lp.blockDim.x,
lp.gridDim.y * lp.blockDim.y,
lp.gridDim.z * lp.blockDim.z,
lp.blockDim.x, lp.blockDim.y,
lp.blockDim.z, lp.sharedMem,
lp.stream, lp.args, nullptr, nullptr, nullptr, 0,
true /* stream is already locked above and will be unlocked
in the below code after launching kernels on all devices*/);
}
// unlock all streams
for (int i = 0; i < numDevices; ++i) {
launchParamsList[i].stream->criticalData().unlock();
#if (__hcc_workweek__ >= 19213)
launchParamsList[i].stream->criticalData()._av.release_locked_hsa_queue();
#endif
}
free(kds);
return result;
}
namespace hip_impl {
hsa_executable_t executable_for(hipModule_t hmod) {
return hmod->executable;
@@ -805,3 +877,147 @@ hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const
*texRef = reinterpret_cast<textureReference*>(addr);
return ihipLogStatus(hipSuccess);
}
hipError_t ihipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* blockSize,
hipFunction_t f, size_t dynSharedMemPerBlk,
uint32_t blockSizeLimit)
{
using namespace hip_impl;
auto ctx = ihipGetTlsDefaultCtx();
hipError_t ret = hipSuccess;
if (ctx == nullptr) {
ret = hipErrorInvalidDevice;
}
hipDeviceProp_t prop{};
ihipGetDeviceProperties(&prop, ihipGetTlsDefaultCtx()->getDevice()->_deviceId);
prop.regsPerBlock = prop.regsPerBlock ? prop.regsPerBlock : 64 * 1024;
size_t usedVGPRS = 0;
size_t usedSGPRS = 0;
size_t usedLDS = 0;
bool is_code_object_v3 = f->_name.find(".kd") != std::string::npos;
if (is_code_object_v3) {
const auto header = reinterpret_cast<const amd_kernel_code_v3_t*>(f->_header);
// GRANULATED_WAVEFRONT_VGPR_COUNT is specified in 0:5 bits of COMPUTE_PGM_RSRC1
// the granularity for gfx6-gfx9 is max(0, ceil(vgprs_used / 4) - 1)
usedVGPRS = ((header->compute_pgm_rsrc1 & 0x3F) + 1) << 2;
// GRANULATED_WAVEFRONT_SGPR_COUNT is specified in 6:9 bits of COMPUTE_PGM_RSRC1
// the granularity for gfx9+ is 2 * max(0, ceil(sgprs_used / 16) - 1)
usedSGPRS = ((((header->compute_pgm_rsrc1 & 0x3C0) >> 6) >> 1) + 1) << 4;
usedLDS = header->group_segment_fixed_size;
}
else {
const auto header = f->_header;
// VGPRs granularity is 4
usedVGPRS = ((header->workitem_vgpr_count + 3) >> 2) << 2;
// adding 2 to take into account the 2 VCC registers & handle the granularity of 16
usedSGPRS = header->wavefront_sgpr_count + 2;
usedSGPRS = ((usedSGPRS + 15) >> 4) << 4;
usedLDS = header->workgroup_group_segment_byte_size;
}
// try different workgroup sizes to find the maximum potential occupancy
// based on the usage of VGPRs and LDS
size_t wavefrontSize = prop.warpSize;
size_t maxWavefrontsPerBlock = prop.maxThreadsPerBlock / wavefrontSize;
// Due to SPI and private memory limitations, the max of wavefronts per CU in 32
size_t maxWavefrontsPerCU = min(prop.maxThreadsPerMultiProcessor / wavefrontSize, 32);
const size_t numSIMD = 4;
size_t maxActivWaves = 0;
size_t maxWavefronts = 0;
for (int i = 0; i < maxWavefrontsPerBlock; i++) {
size_t wavefrontsPerWG = i + 1;
// workgroup per CU is 40 for WG size of 1 wavefront; otherwise it is 16
size_t maxWorkgroupPerCU = (wavefrontsPerWG == 1) ? 40 : 16;
size_t maxWavesWGLimited = min(wavefrontsPerWG * maxWorkgroupPerCU, maxWavefrontsPerCU);
// Compute VGPR limited wavefronts per block
size_t wavefrontsVGPRS;
if (usedVGPRS == 0) {
wavefrontsVGPRS = maxWavesWGLimited;
}
else {
// find how many VGPRs are available for each SIMD
size_t numVGPRsPerSIMD = (prop.regsPerBlock / wavefrontSize / numSIMD);
wavefrontsVGPRS = (numVGPRsPerSIMD / usedVGPRS) * numSIMD;
}
size_t maxWavesVGPRSLimited = 0;
if (wavefrontsVGPRS > maxWavesWGLimited) {
maxWavesVGPRSLimited = maxWavesWGLimited;
}
else {
maxWavesVGPRSLimited = (wavefrontsVGPRS / wavefrontsPerWG) * wavefrontsPerWG;
}
// Compute SGPR limited wavefronts per block
size_t wavefrontsSGPRS;
if (usedSGPRS == 0) {
wavefrontsSGPRS = maxWavesWGLimited;
}
else {
const size_t numSGPRsPerSIMD = (prop.gcnArch < 900) ? 512 : 800;
wavefrontsSGPRS = (numSGPRsPerSIMD / usedSGPRS) * numSIMD;
}
size_t maxWavesSGPRSLimited = 0;
if (wavefrontsSGPRS > maxWavesWGLimited) {
maxWavesSGPRSLimited = maxWavesWGLimited;
}
else {
maxWavesSGPRSLimited = (wavefrontsSGPRS / wavefrontsPerWG) * wavefrontsPerWG;
}
// Compute LDS limited wavefronts per block
size_t wavefrontsLDS;
if (usedLDS == 0) {
wavefrontsLDS = maxWorkgroupPerCU * wavefrontsPerWG;
}
else {
size_t availableSharedMemPerCU = prop.maxSharedMemoryPerMultiProcessor;
size_t workgroupPerCU = availableSharedMemPerCU / (usedLDS + dynSharedMemPerBlk);
wavefrontsLDS = min(workgroupPerCU, maxWorkgroupPerCU) * wavefrontsPerWG;
}
size_t maxWavesLDSLimited = min(wavefrontsLDS, maxWavefrontsPerCU);
size_t activeWavefronts = 0;
size_t tmp_min = (size_t)min(maxWavesLDSLimited, maxWavesWGLimited);
tmp_min = min(maxWavesSGPRSLimited, tmp_min);
activeWavefronts = min(maxWavesVGPRSLimited, tmp_min);
if (maxActivWaves < activeWavefronts) {
maxActivWaves = activeWavefronts;
maxWavefronts = wavefrontsPerWG;
}
}
// determine the grid and block sizes for maximum potential occupancy
size_t maxThreadsCnt = prop.maxThreadsPerMultiProcessor*prop.multiProcessorCount;
if (blockSizeLimit > 0) {
maxThreadsCnt = min(maxThreadsCnt, blockSizeLimit);
}
*blockSize = maxWavefronts * wavefrontSize;
*gridSize = min((maxThreadsCnt + *blockSize - 1) / *blockSize, prop.multiProcessorCount);
return ret;
}
hipError_t hipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* blockSize,
hipFunction_t f, size_t dynSharedMemPerBlk,
uint32_t blockSizeLimit)
{
HIP_INIT_API(hipOccupancyMaxPotentialBlockSize, gridSize, blockSize, f, dynSharedMemPerBlk, blockSizeLimit);
return ihipLogStatus(ihipOccupancyMaxPotentialBlockSize(
gridSize, blockSize, f, dynSharedMemPerBlk, blockSizeLimit));
}
include/hip/hcc_detail/hip_prof_api.h → src/hip_prof_api.h
مشاهده پرونده
tests/src/runtimeApi/module/hipOccupancyMaxPotentialBlockSize.cpp
+69
مشاهده پرونده
@@ -0,0 +1,69 @@
/*
Copyright (c) 2019 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.
*/
// Test the Grid_Launch syntax.
/* HIT_START
* BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
* TEST: %t
* HIT_END
*/
#include "hip/hip_runtime.h"
#include "test_common.h"
#define fileName "vcpy_kernel.code"
#define kernel_name "hello_world"
__global__ void f1(float *a) { *a = 1.0; }
template <typename T>
__global__ void f2(T *a) { *a = 1; }
int main(int argc, char* argv[]) {
// test case for using kernel function pointer
uint32_t gridSize = 0;
uint32_t blockSize = 0;
hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f1, 0, 0);
assert(gridSize != 0 && blockSize != 0);
// test case for using kernel function pointer with template
gridSize = 0;
blockSize = 0;
hipOccupancyMaxPotentialBlockSize<void(*)(int *)>(&gridSize, &blockSize, f2, 0, 0);
assert(gridSize != 0 && blockSize != 0);
// test case for using kernel with hipFunction_t type
gridSize = 0;
blockSize = 0;
hipModule_t Module;
hipFunction_t Function;
HIPCHECK(hipModuleLoad(&Module, fileName));
HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name));
HIPCHECK(hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, Function, 0, 0));
assert(gridSize != 0 && blockSize != 0);
passed();
}