AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

resolve merge conflict

Просмотр исходного кода
Этот коммит содержится в:
Aryan Salmanpour
2020-02-10 10:30:55 -05:00
родитель d805e98fcb 484d67a485
Коммит dfa565d76f
47 изменённых файлов: 887 добавлений и 1063 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
hipamd/src/hip_clang.cpp
+38
Просмотреть файл
@@ -189,6 +189,44 @@ hipError_t hipConfigureCall(
return hipSuccess;
}
extern "C" hipError_t __hipPushCallConfiguration(
dim3 gridDim,
dim3 blockDim,
size_t sharedMem,
hipStream_t stream)
{
GET_TLS();
auto ctx = ihipGetTlsDefaultCtx();
LockedAccessor_CtxCrit_t crit(ctx->criticalData());
crit->_execStack.push(ihipExec_t{gridDim, blockDim, sharedMem, stream});
return hipSuccess;
}
extern "C" hipError_t __hipPopCallConfiguration(
dim3 *gridDim,
dim3 *blockDim,
size_t *sharedMem,
hipStream_t *stream)
{
GET_TLS();
auto ctx = ihipGetTlsDefaultCtx();
LockedAccessor_CtxCrit_t crit(ctx->criticalData());
ihipExec_t exec;
exec = std::move(crit->_execStack.top());
crit->_execStack.pop();
*gridDim = exec._gridDim;
*blockDim = exec._blockDim;
*sharedMem = exec._sharedMem;
*stream = exec._hStream;
return hipSuccess;
}
hipError_t hipSetupArgument(
const void *arg,
size_t size,
hipamd/src/hip_module.cpp
+49 -52
Просмотреть файл
@@ -134,10 +134,10 @@ extern hipError_t ihipGetDeviceProperties(hipDeviceProp_t* props, int device);
return ihipLogStatus(hipStatus); \
}
hipError_t ihipModuleLaunchKernel(TlsData *tls, hipFunction_t f, uint32_t globalWorkSizeX,
uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
uint32_t localWorkSizeX, uint32_t localWorkSizeY,
uint32_t localWorkSizeZ, size_t sharedMemBytes,
hipError_t ihipModuleLaunchKernel(TlsData *tls, hipFunction_t f, uint32_t gridSizeX,
uint32_t gridSizeY, uint32_t gridSizeZ,
uint32_t blockSizeX, uint32_t blockSizeY,
uint32_t blockSizeZ, size_t sharedMemBytes,
hipStream_t hStream, void** kernelParams, void** extra,
hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags, bool isStreamLocked = 0,
void** impCoopParams = 0) {
@@ -146,6 +146,14 @@ hipError_t ihipModuleLaunchKernel(TlsData *tls, hipFunction_t f, uint32_t global
auto ctx = ihipGetTlsDefaultCtx();
hipError_t ret = hipSuccess;
size_t globalWorkSizeX = (size_t)gridSizeX * (size_t)blockSizeX;
size_t globalWorkSizeY = (size_t)gridSizeY * (size_t)blockSizeY;
size_t globalWorkSizeZ = (size_t)gridSizeZ * (size_t)blockSizeZ;
if(globalWorkSizeX > UINT32_MAX || globalWorkSizeY > UINT32_MAX || globalWorkSizeZ > UINT32_MAX)
{
return hipErrorInvalidConfiguration;
}
if (ctx == nullptr) {
ret = hipErrorInvalidDevice;
@@ -180,7 +188,8 @@ hipError_t ihipModuleLaunchKernel(TlsData *tls, hipFunction_t f, uint32_t global
return hipErrorNotInitialized;
}
} else {
}
else if (f->_kernarg_layout.size() != 0) {
return hipErrorInvalidValue;
}
@@ -202,8 +211,8 @@ hipError_t ihipModuleLaunchKernel(TlsData *tls, hipFunction_t f, uint32_t global
lp.dynamic_group_mem_bytes =
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(), isStreamLocked);
hStream, dim3(globalWorkSizeX/blockSizeX, globalWorkSizeY/blockSizeY, globalWorkSizeZ/blockSizeZ),
dim3(blockSizeX, blockSizeY, blockSizeZ), &lp, f->_name.c_str(), isStreamLocked);
hsa_kernel_dispatch_packet_t aql;
@@ -212,9 +221,9 @@ hipError_t ihipModuleLaunchKernel(TlsData *tls, hipFunction_t f, uint32_t global
// aql.completion_signal._handle = 0;
// aql.kernarg_address = 0;
aql.workgroup_size_x = localWorkSizeX;
aql.workgroup_size_y = localWorkSizeY;
aql.workgroup_size_z = localWorkSizeZ;
aql.workgroup_size_x = blockSizeX;
aql.workgroup_size_y = blockSizeY;
aql.workgroup_size_z = blockSizeZ;
aql.grid_size_x = globalWorkSizeX;
aql.grid_size_y = globalWorkSizeY;
aql.grid_size_z = globalWorkSizeZ;
@@ -275,7 +284,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, uint32_t gridDimX, uint32_t gr
HIP_INIT_API(hipModuleLaunchKernel, f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes,
hStream, kernelParams, extra);
return ihipLogStatus(ihipModuleLaunchKernel(tls,
f, blockDimX * gridDimX, blockDimY * gridDimY, gridDimZ * blockDimZ, blockDimX, blockDimY,
f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY,
blockDimZ, sharedMemBytes, hStream, kernelParams, extra, nullptr, nullptr, 0));
}
@@ -287,8 +296,12 @@ hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags) {
HIP_INIT_API(hipExtModuleLaunchKernel, f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX,
localWorkSizeY, localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra);
if(localWorkSizeX == 0 || localWorkSizeY == 0 || localWorkSizeZ == 0)
return hipErrorInvalidValue;
return ihipLogStatus(ihipModuleLaunchKernel(tls,
f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX, localWorkSizeY,
f, globalWorkSizeX/localWorkSizeX, globalWorkSizeY/localWorkSizeY, globalWorkSizeZ/localWorkSizeZ, localWorkSizeX, localWorkSizeY,
localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, flags));
}
@@ -300,8 +313,12 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
hipEvent_t startEvent, hipEvent_t stopEvent) {
HIP_INIT_API(hipHccModuleLaunchKernel, f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX,
localWorkSizeY, localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra);
if(localWorkSizeX == 0 || localWorkSizeY == 0 || localWorkSizeZ == 0)
return hipErrorInvalidValue;
return ihipLogStatus(ihipModuleLaunchKernel(tls,
f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX, localWorkSizeY,
f, globalWorkSizeX/localWorkSizeX, globalWorkSizeY/localWorkSizeY, globalWorkSizeZ/localWorkSizeZ, localWorkSizeX, localWorkSizeY,
localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, 0));
}
@@ -352,9 +369,9 @@ hipError_t ihipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList
const hipLaunchParams& lp = launchParamsList[i];
result = ihipModuleLaunchKernel(tls, kds[i],
lp.gridDim.x * lp.blockDim.x,
lp.gridDim.y * lp.blockDim.y,
lp.gridDim.z * lp.blockDim.z,
lp.gridDim.x,
lp.gridDim.y,
lp.gridDim.z,
lp.blockDim.x, lp.blockDim.y,
lp.blockDim.z, lp.sharedMem,
lp.stream, lp.args, nullptr, nullptr, nullptr, 0,
@@ -458,9 +475,9 @@ hipError_t ihipLaunchCooperativeKernel(const void* f, dim3 gridDim,
// launch the main kernel
result = ihipModuleLaunchKernel(tls, kd,
gridDim.x * blockDimX.x,
gridDim.y * blockDimX.y,
gridDim.z * blockDimX.z,
gridDim.x,
gridDim.y,
gridDim.z,
blockDimX.x, blockDimX.y, blockDimX.z,
sharedMemBytes, stream, kernelParams, nullptr, nullptr,
nullptr, 0, true, impCoopParams);
@@ -612,9 +629,9 @@ hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsL
impCoopParams[0] = &mg_info_ptr[i];
result = ihipModuleLaunchKernel(tls, kds[i],
lp.gridDim.x * lp.blockDim.x,
lp.gridDim.y * lp.blockDim.y,
lp.gridDim.z * lp.blockDim.z,
lp.gridDim.x,
lp.gridDim.y,
lp.gridDim.z,
lp.blockDim.x, lp.blockDim.y,
lp.blockDim.z, lp.sharedMem,
lp.stream, lp.args, nullptr, nullptr, nullptr, 0,
@@ -1368,7 +1385,7 @@ hipError_t hipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* block
}
hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor(
TlsData *tls, int* numBlocks, hipFunction_t f, int blockSize, size_t dynSharedMemPerBlk)
TlsData *tls, uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk)
{
using namespace hip_impl;
@@ -1408,41 +1425,35 @@ hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor(
: std::min(maxWavesPerSimd, availableSGPRs / usedSGPRS));
// Calculate blocks occupancy per CU based on SGPR usage
*numBlocks = std::min(*numBlocks, (int) (sgprs_alu_occupancy / numWavefronts));
*numBlocks = std::min(*numBlocks, (uint32_t) (sgprs_alu_occupancy / numWavefronts));
size_t total_used_lds = usedLDS + dynSharedMemPerBlk;
if (total_used_lds != 0) {
// Calculate LDS occupacy per CU. lds_per_cu / (static_lsd + dynamic_lds)
size_t lds_occupancy = prop.maxSharedMemoryPerMultiProcessor / total_used_lds;
*numBlocks = std::min(*numBlocks, (int) lds_occupancy);
*numBlocks = std::min(*numBlocks, (uint32_t) lds_occupancy);
}
return hipSuccess;
}
hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(
int* numBlocks, const void* f, int blockSize, size_t dynSharedMemPerBlk)
uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk)
{
HIP_INIT_API(hipOccupancyMaxActiveBlocksPerMultiprocessor, numBlocks, f, blockSize, dynSharedMemPerBlk);
auto F = hip_impl::get_program_state().kernel_descriptor((std::uintptr_t)(f),
hip_impl::target_agent(0));
return ihipLogStatus(ihipOccupancyMaxActiveBlocksPerMultiprocessor(
tls, numBlocks, F, blockSize, dynSharedMemPerBlk));
tls, numBlocks, f, blockSize, dynSharedMemPerBlk));
}
hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
int* numBlocks, const void* f, int blockSize, size_t dynSharedMemPerBlk,
uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk,
unsigned int flags)
{
HIP_INIT_API(hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags, numBlocks, f, blockSize, dynSharedMemPerBlk, flags);
auto F = hip_impl::get_program_state().kernel_descriptor((std::uintptr_t)(f),
hip_impl::target_agent(0));
return ihipLogStatus(ihipOccupancyMaxActiveBlocksPerMultiprocessor(
tls, numBlocks, F, blockSize, dynSharedMemPerBlk));
tls, numBlocks, f, blockSize, dynSharedMemPerBlk));
}
hipError_t hipLaunchKernel(
@@ -1454,21 +1465,7 @@ hipError_t hipLaunchKernel(
hipFunction_t kd = hip_impl::get_program_state().kernel_descriptor((std::uintptr_t)func_addr,
hip_impl::target_agent(stream));
if(kd == nullptr || kd->_header == nullptr)
return ihipLogStatus(hipErrorInvalidValue);
size_t szKernArg = kd->_header->kernarg_segment_byte_size;
if(args == NULL && szKernArg != 0)
return ihipLogStatus(hipErrorInvalidValue);
void* config[]{
HIP_LAUNCH_PARAM_BUFFER_POINTER,
args,
HIP_LAUNCH_PARAM_BUFFER_SIZE,
&szKernArg,
HIP_LAUNCH_PARAM_END};
return ihipLogStatus(ihipModuleLaunchKernel(tls, kd, numBlocks.x * dimBlocks.x, numBlocks.y * dimBlocks.y, numBlocks.z * dimBlocks.z,
dimBlocks.x, dimBlocks.y, dimBlocks.z, sharedMemBytes, stream, nullptr, (void**)&config, nullptr, nullptr, 0));
return hipModuleLaunchKernel(kd, numBlocks.x, numBlocks.y, numBlocks.z,
dimBlocks.x, dimBlocks.y, dimBlocks.z, sharedMemBytes,
stream, args, nullptr);
}
hipamd/src/hip_texture.cpp
+23 -12
Просмотреть файл
@@ -29,24 +29,32 @@ void saveTextureInfo(const hipTexture* pTexture, const hipResourceDesc* pResDesc
}
}
void getDrvChannelOrderAndType(const enum hipArray_Format Format, unsigned int NumChannels,
void getDrvChannelOrderAndType(const enum hipArray_Format Format, enum hipTextureReadMode readMode, unsigned int NumChannels,
hsa_ext_image_channel_order_t* channelOrder,
hsa_ext_image_channel_type_t* channelType) {
switch (Format) {
case HIP_AD_FORMAT_UNSIGNED_INT8:
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8;
*channelType = readMode == hipReadModeNormalizedFloat
? HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT8
: HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8;
break;
case HIP_AD_FORMAT_UNSIGNED_INT16:
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16;
*channelType = readMode == hipReadModeNormalizedFloat
? HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT16
: HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16;
break;
case HIP_AD_FORMAT_UNSIGNED_INT32:
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32;
break;
case HIP_AD_FORMAT_SIGNED_INT8:
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8;
*channelType = readMode == hipReadModeNormalizedFloat
? HSA_EXT_IMAGE_CHANNEL_TYPE_SNORM_INT8
: HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8;
break;
case HIP_AD_FORMAT_SIGNED_INT16:
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT16;
*channelType = readMode == hipReadModeNormalizedFloat
? HSA_EXT_IMAGE_CHANNEL_TYPE_SNORM_INT16
: HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT16;
break;
case HIP_AD_FORMAT_SIGNED_INT32:
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32;
@@ -422,7 +430,7 @@ hipError_t ihipBindTextureImpl(TlsData *tls_, int dim, enum hipTextureReadMode r
hsa_ext_image_channel_order_t channelOrder;
hsa_ext_image_channel_type_t channelType;
if (NULL == desc) {
getDrvChannelOrderAndType(tex->format, tex->numChannels, &channelOrder, &channelType);
getDrvChannelOrderAndType(tex->format, readMode, tex->numChannels, &channelOrder, &channelType);
} else {
getChannelOrderAndType(*desc, readMode, &channelOrder, &channelType);
}
@@ -497,7 +505,7 @@ hipError_t ihipBindTexture2DImpl(TlsData *tls, int dim, enum hipTextureReadMode
hsa_ext_image_channel_type_t channelType;
if (NULL == desc) {
getDrvChannelOrderAndType(tex->format, tex->numChannels, &channelOrder, &channelType);
getDrvChannelOrderAndType(tex->format, readMode, tex->numChannels, &channelOrder, &channelType);
} else {
getChannelOrderAndType(*desc, readMode, &channelOrder, &channelType);
}
@@ -602,7 +610,7 @@ hipError_t ihipBindTextureToArrayImpl(TlsData *tls_, int dim, enum hipTextureRea
hsa_ext_image_channel_order_t channelOrder;
hsa_ext_image_channel_type_t channelType;
if (array->isDrv) {
getDrvChannelOrderAndType(array->Format, array->NumChannels,
getDrvChannelOrderAndType(array->Format, readMode, array->NumChannels,
&channelOrder, &channelType);
} else {
getChannelOrderAndType(desc, readMode, &channelOrder, &channelType);
@@ -724,7 +732,10 @@ hipError_t hipTexRefSetFormat(textureReference* tex, hipArray_Format fmt, int Nu
hipError_t hipTexRefSetFlags(textureReference* tex, unsigned int flags) {
HIP_INIT_API(hipTexRefSetFlags, tex, flags);
hipError_t hip_status = hipSuccess;
tex->normalized = flags;
if(flags == HIP_TRSF_READ_AS_INTEGER)
tex->readMode = hipReadModeElementType;
else if(flags == HIP_TRSF_NORMALIZED_COORDINATES)
tex->normalized = flags;
return ihipLogStatus(hip_status);
}
@@ -757,7 +768,7 @@ hipError_t hipTexRefSetArray(textureReference* tex, hipArray_const_t array, unsi
HIP_INIT_API(hipTexRefSetArray, tex, array, flags);
hipError_t hip_status = hipSuccess;
hip_status = ihipBindTextureToArrayImpl(tls, array->textureType, hipReadModeElementType, array,
hip_status = ihipBindTextureToArrayImpl(tls, array->textureType, tex->readMode, array,
array->desc, tex);
return ihipLogStatus(hip_status);
}
@@ -785,7 +796,7 @@ hipError_t hipTexRefSetAddress(size_t* offset, textureReference* tex, hipDevicep
HIP_INIT_API(hipTexRefSetAddress, offset, tex, devPtr, size);
hipError_t hip_status = hipSuccess;
// TODO: hipReadModeElementType is default.
hip_status = ihipBindTextureImpl(tls, hipTextureType1D, hipReadModeElementType, offset, devPtr, NULL,
hip_status = ihipBindTextureImpl(tls, hipTextureType1D, tex->readMode, offset, devPtr, NULL,
size, tex);
return ihipLogStatus(hip_status);
}
@@ -816,7 +827,7 @@ hipError_t hipTexRefSetAddress2D(textureReference* tex, const HIP_ARRAY_DESCRIPT
//TODO: Fix when HSA accepts user defined pitch
if(pitch % 64) pitch =0;
hip_status = ihipBindTexture2DImpl(tls, hipTextureType2D, hipReadModeElementType, &offset, devPtr,
hip_status = ihipBindTexture2DImpl(tls, hipTextureType2D, tex->readMode, &offset, devPtr,
NULL, desc->Width, desc->Height, tex, pitch);
return ihipLogStatus(hip_status);
}
hipamd/src/program_state.inl
+5 -3
Просмотреть файл
@@ -613,7 +613,8 @@ public:
for (auto&& kernel_symbol : it->second) {
functions[aa].second.emplace(
function.first,
Kernel_descriptor{kernel_object(kernel_symbol), it->first});
Kernel_descriptor{kernel_object(kernel_symbol), it->first,
kernargs_size_align(function.first)});
}
}
}, agent);
@@ -672,11 +673,12 @@ public:
auto dx1 = kernels_md.find("CodeProps", dx);
dx = kernels_md.find("Args:", dx);
if (dx1 < dx) {
if (dx1 < dx || dx == std::string::npos) {
dx = dx1;
// create an empty kernarg laybout vector for kernels without any arg
kernargs[fn];
continue;
}
if (dx == std::string::npos) break;
static constexpr decltype(kernels_md.size()) args_sz{5};
dx = parse_args_v2(kernels_md, dx + args_sz, dx1, kernargs[fn]);