AI/rocm-systems
2
0
フォーク 0
コード イシュー プルリクエスト Actions パッケージ プロジェクト リリース Wiki アクティビティ

Merge remote-tracking branch 'origin/master' into feature_use_module_based_dispatch_instead_of_pfe

ソースを参照 
# Conflicts:
#	src/hip_module.cpp
このコミットが含まれているのは:
Alex Voicu
2017-11-28 17:29:11 +00:00
親 265c3b224e 0da0426f94
コミット dc67ca3feb
28個のファイルの変更、5344行の追加、1054行の削除
Patchファイルをダウンロード Diffファイルをダウンロード すべてのファイルを展開 すべてのファイルを折り畳む
src/device_functions.cpp
+9 -4
ファイルの表示
@@ -23,6 +23,11 @@ THE SOFTWARE.
#include <hc_math.hpp>
#include "device_util.h"
extern "C" float __ocml_floor_f32(float);
extern "C" float __ocml_rint_f32(float);
extern "C" float __ocml_ceil_f32(float);
extern "C" float __ocml_trunc_f32(float);
struct holder64Bit{
union{
double d;
@@ -163,19 +168,19 @@ __device__ long long int __double_as_longlong(double x)
__device__ int __float2int_rd(float x)
{
return (int)x;
return (int)__ocml_floor_f32(x);
}
__device__ int __float2int_rn(float x)
{
return (int)x;
return (int)__ocml_rint_f32(x);
}
__device__ int __float2int_ru(float x)
{
return (int)x;
return (int)__ocml_ceil_f32(x);
}
__device__ int __float2int_rz(float x)
{
return (int)x;
return (int)__ocml_trunc_f32(x);
}
__device__ long long int __float2ll_rd(float x)
src/hip_hcc_internal.h
+10 -8
ファイルの表示
@@ -25,8 +25,9 @@ THE SOFTWARE.
#include <hc.hpp>
#include <hsa/hsa.h>
#include "hsa/hsa_ext_amd.h"
#include <unordered_map>
#include "hsa/hsa_ext_amd.h"
#include "hip/hip_runtime.h"
#include "hip_util.h"
#include "env.h"
@@ -373,13 +374,14 @@ public:
class ihipModule_t {
public:
hsa_executable_t executable;
hsa_code_object_t object;
std::string fileName;
void *ptr;
size_t size;
std::list<hipFunction_t> funcTrack;
ihipModule_t() : executable(), object(), fileName(), ptr(nullptr), size(0) {}
hsa_executable_t executable;
hsa_code_object_t object;
std::string fileName;
void *ptr;
size_t size;
std::list<hipFunction_t> funcTrack;
std::unordered_map<std::string, uintptr_t> coGlobals;
ihipModule_t() : executable(), object(), fileName(), ptr(nullptr), size(0) {}
};
src/hip_memory.cpp
+134 -10
ファイルの表示
@@ -407,8 +407,113 @@ hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannel
extern void getChannelOrderAndType(const hipChannelFormatDesc& desc,
enum hipTextureReadMode readMode,
hsa_ext_image_channel_order_t& channelOrder,
hsa_ext_image_channel_type_t& channelType);
hsa_ext_image_channel_order_t* channelOrder,
hsa_ext_image_channel_type_t* channelType);
hipError_t hipArrayCreate ( hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray )
{
HIP_INIT_SPECIAL_API((TRACE_MEM), array, pAllocateArray);
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
if(pAllocateArray->width >0) {
auto ctx = ihipGetTlsDefaultCtx();
*array = (hipArray*)malloc(sizeof(hipArray));
array[0]->drvDesc = *pAllocateArray;
array[0]->width = pAllocateArray->width;
array[0]->height = pAllocateArray->height;
array[0]->isDrv = true;
void ** ptr = &array[0]->data;
if (ctx) {
const unsigned am_flags = 0;
size_t size = pAllocateArray->width;
if(pAllocateArray->height > 0) {
size = size * pAllocateArray->height;
}
hsa_ext_image_channel_type_t channelType;
size_t allocSize = 0;
switch(pAllocateArray->format) {
case HIP_AD_FORMAT_UNSIGNED_INT8:
allocSize = size * sizeof(uint8_t);
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8;
break;
case HIP_AD_FORMAT_UNSIGNED_INT16:
allocSize = size * sizeof(uint16_t);
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16;
break;
case HIP_AD_FORMAT_UNSIGNED_INT32:
allocSize = size * sizeof(uint32_t);
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32;
break;
case HIP_AD_FORMAT_SIGNED_INT8:
allocSize = size * sizeof(int8_t);
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8;
break;
case HIP_AD_FORMAT_SIGNED_INT16:
allocSize = size * sizeof(int16_t);
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT16;
break;
case HIP_AD_FORMAT_SIGNED_INT32:
allocSize = size * sizeof(int32_t);
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32;
break;
case HIP_AD_FORMAT_HALF:
allocSize = size * sizeof(int16_t);
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_HALF_FLOAT;
break;
case HIP_AD_FORMAT_FLOAT:
allocSize = size * sizeof(float);
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT;
break;
default:
hip_status = hipErrorUnknown;
break;
}
hc::accelerator acc = ctx->getDevice()->_acc;
hsa_agent_t* agent =static_cast<hsa_agent_t*>(acc.get_hsa_agent());
size_t allocGranularity = 0;
hsa_amd_memory_pool_t *allocRegion = static_cast<hsa_amd_memory_pool_t*>(acc.get_hsa_am_region());
hsa_amd_memory_pool_get_info(*allocRegion, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, &allocGranularity);
hsa_ext_image_descriptor_t imageDescriptor;
imageDescriptor.width = pAllocateArray->width;
imageDescriptor.height = pAllocateArray->height;
imageDescriptor.depth = 0;
imageDescriptor.array_size = 0;
imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D;
hsa_ext_image_channel_order_t channelOrder;
if (pAllocateArray->numChannels == 4) {
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA;
} else if (pAllocateArray->numChannels == 2) {
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG;
} else if (pAllocateArray->numChannels == 1) {
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R;
}
imageDescriptor.format.channel_order = channelOrder;
imageDescriptor.format.channel_type = channelType;
hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
hsa_ext_image_data_info_t imageInfo;
hsa_status_t status = hsa_ext_image_data_get_info(*agent, &imageDescriptor, permission, &imageInfo);
size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment;
*ptr = hip_internal::allocAndSharePtr("device_array", allocSize, ctx, false/*shareWithAll*/, am_flags, 0, alignment);
if (size && (*ptr == NULL)) {
hip_status = hipErrorMemoryAllocation;
}
} else {
hip_status = hipErrorMemoryAllocation;
}
} else {
hip_status = hipErrorInvalidValue;
}
return ihipLogStatus(hip_status);
}
hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc,
size_t width, size_t height, unsigned int flags)
@@ -425,6 +530,7 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc,
array[0]->height = height;
array[0]->depth = 1;
array[0]->desc = *desc;
array[0]->isDrv = false;
void ** ptr = &array[0]->data;
@@ -480,7 +586,7 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc,
}
hsa_ext_image_channel_order_t channelOrder;
hsa_ext_image_channel_type_t channelType;
getChannelOrderAndType(*desc, hipReadModeElementType, channelOrder, channelType);
getChannelOrderAndType(*desc, hipReadModeElementType, &channelOrder, &channelType);
imageDescriptor.format.channel_order = channelOrder;
imageDescriptor.format.channel_type = channelType;
@@ -577,7 +683,7 @@ hipError_t hipMalloc3DArray(hipArray_t *array,
}
hsa_ext_image_channel_order_t channelOrder;
hsa_ext_image_channel_type_t channelType;
getChannelOrderAndType(*desc, hipReadModeElementType, channelOrder, channelType);
getChannelOrderAndType(*desc, hipReadModeElementType, &channelOrder, &channelType);
imageDescriptor.format.channel_order = channelOrder;
imageDescriptor.format.channel_type = channelType;
@@ -1005,13 +1111,11 @@ hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, h
}
// TODO - review and optimize
hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch,
size_t width, size_t height, hipMemcpyKind kind) {
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind);
hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch,
size_t width, size_t height, hipMemcpyKind kind)
{
if(width > dpitch || width > spitch)
return ihipLogStatus(hipErrorUnknown);
return hipErrorUnknown;
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
@@ -1028,6 +1132,26 @@ hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch,
e = ex._code;
}
return e;
}
hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch,
size_t width, size_t height, hipMemcpyKind kind)
{
HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind);
hipError_t e = hipSuccess;
e = ihipMemcpy2D(dst,dpitch, src, spitch, width, height, kind);
return ihipLogStatus(e);
}
hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy)
{
HIP_INIT_SPECIAL_API((TRACE_MCMD), pCopy);
hipError_t e = hipSuccess;
if(pCopy == nullptr) {
e = hipErrorInvalidValue;
}
e = ihipMemcpy2D(pCopy->dstArray->data, pCopy->widthInBytes, pCopy->srcHost, pCopy->srcPitch, pCopy->widthInBytes, pCopy->height, hipMemcpyDefault);
return ihipLogStatus(e);
}
src/hip_module.cpp
+25 -9
ファイルの表示
@@ -120,18 +120,15 @@ namespace hipdrv {
uint64_t PrintSymbolSizes(const void *emi, const char *name){
using namespace ELFIO;
const ELFIO::Elf64_Ehdr *ehdr = (const ELFIO::Elf64_Ehdr*)emi;
const Elf64_Ehdr *ehdr = (const Elf64_Ehdr*)emi;
if(NULL == ehdr || EV_CURRENT != ehdr->e_version){}
const ELFIO::Elf64_Shdr * shdr =
(const ELFIO::Elf64_Shdr*)((char*)emi + ehdr->e_shoff);
const Elf64_Shdr * shdr = (const Elf64_Shdr*)((char*)emi + ehdr->e_shoff);
for(uint16_t i=0;i<ehdr->e_shnum;++i){
if(shdr[i].sh_type == SHT_SYMTAB){
const ELFIO::Elf64_Sym *syms =
(const ELFIO::Elf64_Sym*)((char*)emi + shdr[i].sh_offset);
const Elf64_Sym *syms = (const Elf64_Sym*)((char*)emi + shdr[i].sh_offset);
assert(syms);
uint64_t numSyms = shdr[i].sh_size/shdr[i].sh_entsize;
const char* strtab =
(const char*)((char*)emi + shdr[shdr[i].sh_link].sh_offset);
const char* strtab = (const char*)((char*)emi + shdr[shdr[i].sh_link].sh_offset);
assert(strtab);
for(uint64_t i=0;i<numSyms;++i){
const char *symname = strtab + syms[i].st_name;
@@ -149,8 +146,8 @@ uint64_t PrintSymbolSizes(const void *emi, const char *name){
uint64_t ElfSize(const void *emi){
using namespace ELFIO;
const ELFIO::Elf64_Ehdr *ehdr = (const ELFIO::Elf64_Ehdr*)emi;
const ELFIO::Elf64_Shdr *shdr = (const ELFIO::Elf64_Shdr*)((char*)emi + ehdr->e_shoff);
const Elf64_Ehdr *ehdr = (const Elf64_Ehdr*)emi;
const Elf64_Shdr *shdr = (const Elf64_Shdr*)((char*)emi + ehdr->e_shoff);
uint64_t max_offset = ehdr->e_shoff;
uint64_t total_size = max_offset + ehdr->e_shentsize * ehdr->e_shnum;
@@ -717,3 +714,22 @@ hipError_t hipModuleLoadDataEx(hipModule_t *module, const void *image, unsigned
{
return hipModuleLoadData(module, image);
}
hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const char* name)
{
HIP_INIT_API(texRef, hmod, name);
hipError_t ret = hipErrorNotFound;
if(texRef == NULL){
ret = hipErrorInvalidValue;
} else {
if(name == NULL || hmod == NULL){
ret = hipErrorNotInitialized;
} else{
const auto it = hmod->coGlobals.find(name);
if (it == hmod->coGlobals.end()) return ihipLogStatus(hipErrorInvalidValue);
*texRef = reinterpret_cast<textureReference*>(it->second);
ret = hipSuccess;
}
}
return ihipLogStatus(ret);
}
src/hip_texture.cpp
+190 -88
ファイルの表示
@@ -32,19 +32,61 @@ void saveTextureInfo(const hipTexture* pTexture,
}
}
void getDrvChannelOrderAndType(const enum hipArray_Format Format,
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;
break;
case HIP_AD_FORMAT_UNSIGNED_INT16:
*channelType = 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;
break;
case HIP_AD_FORMAT_SIGNED_INT16:
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT16;
break;
case HIP_AD_FORMAT_SIGNED_INT32:
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32;
break;
case HIP_AD_FORMAT_HALF:
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_HALF_FLOAT;
break;
case HIP_AD_FORMAT_FLOAT:
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT;
break;
default:
break;
}
if (NumChannels == 4) {
*channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA;
} else if (NumChannels == 2) {
*channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG;
} else if (NumChannels == 1) {
*channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R;
}
}
void getChannelOrderAndType(const hipChannelFormatDesc& desc,
enum hipTextureReadMode readMode,
hsa_ext_image_channel_order_t& channelOrder,
hsa_ext_image_channel_type_t& channelType)
hsa_ext_image_channel_order_t* channelOrder,
hsa_ext_image_channel_type_t* channelType)
{
if (desc.x != 0 && desc.y != 0 && desc.z != 0 && desc.w != 0) {
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA;
*channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA;
} else if (desc.x != 0 && desc.y != 0 && desc.z != 0 && desc.w == 0) {
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGB;
*channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGB;
} else if (desc.x != 0 && desc.y != 0 && desc.z == 0 && desc.w == 0) {
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG;
*channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG;
} else if (desc.x != 0 && desc.y == 0 && desc.z == 0 && desc.w == 0) {
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R;
*channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R;
} else {
}
@@ -52,49 +94,49 @@ void getChannelOrderAndType(const hipChannelFormatDesc& desc,
case hipChannelFormatKindUnsigned:
switch(desc.x) {
case 32:
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32;
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32;
break;
case 16:
channelType = readMode == hipReadModeNormalizedFloat ? HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT16 :
*channelType = readMode == hipReadModeNormalizedFloat ? HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT16 :
HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16;
break;
case 8:
channelType = readMode == hipReadModeNormalizedFloat ? HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT8 :
*channelType = readMode == hipReadModeNormalizedFloat ? HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT8 :
HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8;
break;
default:
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32;
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32;
}
break;
case hipChannelFormatKindSigned:
switch(desc.x) {
case 32:
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32;
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32;
break;
case 16:
channelType = readMode == hipReadModeNormalizedFloat ? HSA_EXT_IMAGE_CHANNEL_TYPE_SNORM_INT16 :
*channelType = readMode == hipReadModeNormalizedFloat ? HSA_EXT_IMAGE_CHANNEL_TYPE_SNORM_INT16 :
HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT16;
break;
case 8:
channelType = readMode == hipReadModeNormalizedFloat ? HSA_EXT_IMAGE_CHANNEL_TYPE_SNORM_INT8 :
*channelType = readMode == hipReadModeNormalizedFloat ? HSA_EXT_IMAGE_CHANNEL_TYPE_SNORM_INT8 :
HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8;
break;
default:
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32;
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32;
}
break;
case hipChannelFormatKindFloat:
switch(desc.x) {
case 32:
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT;
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT;
break;
case 16:
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_HALF_FLOAT;
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_HALF_FLOAT;
break;
case 8:
break;
default:
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT;
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT;
}
break;
case hipChannelFormatKindNone:
@@ -168,8 +210,6 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject,
const hipResourceViewDesc* pResViewDesc)
{
HIP_INIT_API(pTexObject, pResDesc, pTexDesc, pResViewDesc);
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
auto ctx = ihipGetTlsDefaultCtx();
@@ -215,7 +255,7 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject,
imageDescriptor.array_size = 0;
break;
}
getChannelOrderAndType(pResDesc->res.array.array->desc, pTexDesc->readMode, channelOrder, channelType);
getChannelOrderAndType(pResDesc->res.array.array->desc, pTexDesc->readMode, &channelOrder, &channelType);
break;
case hipResourceTypeMipmappedArray:
devPtr = pResDesc->res.mipmap.mipmap->data;
@@ -224,7 +264,7 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject,
imageDescriptor.depth = pResDesc->res.mipmap.mipmap->depth;
imageDescriptor.array_size = 0;
imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D;
getChannelOrderAndType(pResDesc->res.mipmap.mipmap->desc, pTexDesc->readMode, channelOrder, channelType);
getChannelOrderAndType(pResDesc->res.mipmap.mipmap->desc, pTexDesc->readMode, &channelOrder, &channelType);
break;
case hipResourceTypeLinear:
devPtr = pResDesc->res.linear.devPtr;
@@ -233,7 +273,7 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject,
imageDescriptor.depth = 0;
imageDescriptor.array_size = 0;
imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_1D; // ? HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR
getChannelOrderAndType(pResDesc->res.linear.desc, pTexDesc->readMode, channelOrder, channelType);
getChannelOrderAndType(pResDesc->res.linear.desc, pTexDesc->readMode, &channelOrder, &channelType);
break;
case hipResourceTypePitch2D:
devPtr = pResDesc->res.pitch2D.devPtr;
@@ -242,7 +282,7 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject,
imageDescriptor.depth = 0;
imageDescriptor.array_size = 0;
imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D;
getChannelOrderAndType(pResDesc->res.pitch2D.desc, pTexDesc->readMode, channelOrder, channelType);
getChannelOrderAndType(pResDesc->res.pitch2D.desc, pTexDesc->readMode, &channelOrder, &channelType);
break;
default:
break;
@@ -271,7 +311,6 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject,
hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject)
{
HIP_INIT_API(textureObject);
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
@@ -296,8 +335,6 @@ hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject)
hipError_t hipGetTextureObjectResourceDesc(hipResourceDesc* pResDesc, hipTextureObject_t textureObject)
{
HIP_INIT_API(pResDesc, textureObject);
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
auto ctx = ihipGetTlsDefaultCtx();
@@ -313,8 +350,6 @@ hipError_t hipGetTextureObjectResourceDesc(hipResourceDesc* pResDesc, hipTexture
hipError_t hipGetTextureObjectResourceViewDesc(hipResourceViewDesc* pResViewDesc, hipTextureObject_t textureObject)
{
HIP_INIT_API(pResViewDesc, textureObject);
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
auto ctx = ihipGetTlsDefaultCtx();
@@ -330,7 +365,6 @@ hipError_t hipGetTextureObjectResourceViewDesc(hipResourceViewDesc* pResViewDesc
hipError_t hipGetTextureObjectTextureDesc(hipTextureDesc* pTexDesc, hipTextureObject_t textureObject)
{
HIP_INIT_API(pTexDesc, textureObject);
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
@@ -349,18 +383,14 @@ hipError_t ihipBindTextureImpl(int dim,
enum hipTextureReadMode readMode,
size_t *offset,
const void *devPtr,
const struct hipChannelFormatDesc& desc,
size_t size,
enum hipTextureAddressMode addressMode,
enum hipTextureFilterMode filterMode,
int normalizedCoords,
hipTextureObject_t& textureObject)
const struct hipChannelFormatDesc* desc,
size_t size, textureReference* tex )
{
HIP_INIT_API();
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
enum hipTextureAddressMode addressMode = tex->addressMode[0];
enum hipTextureFilterMode filterMode = tex->filterMode;
int normalizedCoords = tex->normalized;
hipTextureObject_t& textureObject = tex->textureObject;
auto ctx = ihipGetTlsDefaultCtx();
if (ctx) {
hc::accelerator acc = ctx->getDevice()->_acc;
@@ -385,7 +415,11 @@ hipError_t ihipBindTextureImpl(int dim,
hsa_ext_image_channel_order_t channelOrder;
hsa_ext_image_channel_type_t channelType;
getChannelOrderAndType(desc, readMode, channelOrder, channelType);
if(NULL == desc) {
getDrvChannelOrderAndType(tex->format, tex->numChannels, &channelOrder, &channelType);
} else {
getChannelOrderAndType(*desc, readMode, &channelOrder, &channelType);
}
imageDescriptor.format.channel_order = channelOrder;
imageDescriptor.format.channel_type = channelType;
@@ -396,13 +430,13 @@ hipError_t ihipBindTextureImpl(int dim,
if (HSA_STATUS_SUCCESS != hsa_ext_image_create_with_layout(*agent, &imageDescriptor, devPtr, permission, HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, 0, 0, &(pTexture->image)) ||
HSA_STATUS_SUCCESS != hsa_ext_sampler_create(*agent, &samplerDescriptor, &(pTexture->sampler))) {
return ihipLogStatus(hipErrorRuntimeOther);
return hipErrorRuntimeOther;
}
getHipTextureObject(&textureObject, pTexture->image, pTexture->sampler);
textureHash[textureObject] = pTexture;
}
return ihipLogStatus(hip_status);
return hip_status;
}
hipError_t hipBindTexture(size_t* offset,
@@ -411,30 +445,28 @@ hipError_t hipBindTexture(size_t* offset,
const hipChannelFormatDesc* desc,
size_t size)
{
HIP_INIT_API(offset, tex, devPtr, desc, size);
hipError_t hip_status = hipSuccess;
// TODO: hipReadModeElementType is default.
return ihipBindTextureImpl(hipTextureType1D, hipReadModeElementType,
offset, devPtr, *desc, size,
tex->addressMode[0], tex->filterMode, tex->normalized,
tex->textureObject);
hip_status = ihipBindTextureImpl(hipTextureType1D, hipReadModeElementType,
offset, devPtr, desc, size, tex);
return ihipLogStatus(hip_status);
}
hipError_t ihipBindTexture2DImpl(int dim,
enum hipTextureReadMode readMode,
size_t *offset,
const void *devPtr,
const struct hipChannelFormatDesc& desc,
const struct hipChannelFormatDesc* desc,
size_t width,
size_t height,
enum hipTextureAddressMode addressMode,
enum hipTextureFilterMode filterMode,
int normalizedCoords,
hipTextureObject_t& textureObject)
textureReference* tex)
{
HIP_INIT_API();
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
enum hipTextureAddressMode addressMode = tex->addressMode[0];
enum hipTextureFilterMode filterMode = tex->filterMode;
int normalizedCoords = tex->normalized;
hipTextureObject_t& textureObject = tex->textureObject;
auto ctx = ihipGetTlsDefaultCtx();
if (ctx) {
hc::accelerator acc = ctx->getDevice()->_acc;
@@ -459,7 +491,12 @@ hipError_t ihipBindTexture2DImpl(int dim,
hsa_ext_image_channel_order_t channelOrder;
hsa_ext_image_channel_type_t channelType;
getChannelOrderAndType(desc, readMode, channelOrder, channelType);
if(NULL == desc) {
getDrvChannelOrderAndType(tex->format, tex->numChannels, &channelOrder, &channelType);
} else {
getChannelOrderAndType(*desc, readMode, &channelOrder, &channelType);
}
imageDescriptor.format.channel_order = channelOrder;
imageDescriptor.format.channel_type = channelType;
@@ -470,13 +507,13 @@ hipError_t ihipBindTexture2DImpl(int dim,
if (HSA_STATUS_SUCCESS != hsa_ext_image_create_with_layout(*agent, &imageDescriptor, devPtr, permission, HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, 0, 0, &(pTexture->image)) ||
HSA_STATUS_SUCCESS != hsa_ext_sampler_create(*agent, &samplerDescriptor, &(pTexture->sampler))) {
return ihipLogStatus(hipErrorRuntimeOther);
return hipErrorRuntimeOther;
}
getHipTextureObject(&textureObject, pTexture->image, pTexture->sampler);
textureHash[textureObject] = pTexture;
}
return ihipLogStatus(hip_status);
return hip_status;
}
hipError_t hipBindTexture2D(size_t* offset,
@@ -487,27 +524,24 @@ hipError_t hipBindTexture2D(size_t* offset,
size_t height,
size_t pitch)
{
// TODO: hipReadModeElementType is default.
return ihipBindTexture2DImpl(hipTextureType2D, hipReadModeElementType,
offset, devPtr, *desc, width, height,
tex->addressMode[0], tex->filterMode, tex->normalized,
tex->textureObject);
HIP_INIT_API(offset, tex, devPtr, desc, width, height, pitch);
hipError_t hip_status = hipSuccess;
hip_status = ihipBindTexture2DImpl(hipTextureType2D, hipReadModeElementType,
offset, devPtr, desc, width, height, tex);
return ihipLogStatus(hip_status);
}
hipError_t ihipBindTextureToArrayImpl(int dim,
enum hipTextureReadMode readMode,
hipArray_const_t array,
const struct hipChannelFormatDesc& desc,
enum hipTextureAddressMode addressMode,
enum hipTextureFilterMode filterMode,
int normalizedCoords,
hipTextureObject_t& textureObject)
textureReference* tex)
{
HIP_INIT_API();
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
enum hipTextureAddressMode addressMode = tex->addressMode[0];
enum hipTextureFilterMode filterMode = tex->filterMode;
int normalizedCoords = tex->normalized;
hipTextureObject_t& textureObject = tex->textureObject;
auto ctx = ihipGetTlsDefaultCtx();
if (ctx) {
hc::accelerator acc = ctx->getDevice()->_acc;
@@ -558,7 +592,11 @@ hipError_t ihipBindTextureToArrayImpl(int dim,
hsa_ext_image_channel_order_t channelOrder;
hsa_ext_image_channel_type_t channelType;
getChannelOrderAndType(desc, readMode, channelOrder, channelType);
if(array->isDrv) {
getDrvChannelOrderAndType(array->drvDesc.format, array->drvDesc.numChannels, &channelOrder, &channelType);
} else {
getChannelOrderAndType(desc, readMode, &channelOrder, &channelType);
}
imageDescriptor.format.channel_order = channelOrder;
imageDescriptor.format.channel_type = channelType;
@@ -569,38 +607,38 @@ hipError_t ihipBindTextureToArrayImpl(int dim,
if (HSA_STATUS_SUCCESS != hsa_ext_image_create_with_layout(*agent, &imageDescriptor, array->data, permission, HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, 0, 0, &(pTexture->image)) ||
HSA_STATUS_SUCCESS != hsa_ext_sampler_create(*agent, &samplerDescriptor, &(pTexture->sampler))) {
return ihipLogStatus(hipErrorRuntimeOther);
return hipErrorRuntimeOther;
}
getHipTextureObject(&textureObject, pTexture->image, pTexture->sampler);
textureHash[textureObject] = pTexture;
}
return ihipLogStatus(hip_status);
return hip_status;
}
hipError_t hipBindTextureToArray(textureReference* tex,
hipArray_const_t array,
const hipChannelFormatDesc* desc)
{
HIP_INIT_API(tex, array, desc);
hipError_t hip_status = hipSuccess;
// TODO: hipReadModeElementType is default.
return ihipBindTextureToArrayImpl(hipTextureType2D, hipReadModeElementType,
array, *desc,
tex->addressMode[0], tex->filterMode, tex->normalized,
tex->textureObject);
hip_status = ihipBindTextureToArrayImpl(hipTextureType2D, hipReadModeElementType,
array, *desc, tex);
return ihipLogStatus(hip_status);
}
hipError_t hipBindTextureToMipmappedArray(textureReference* tex,
hipMipmappedArray_const_t mipmappedArray,
const hipChannelFormatDesc* desc)
{
return hipSuccess;
HIP_INIT_API(tex, mipmappedArray, desc);
hipError_t hip_status = hipSuccess;
return ihipLogStatus(hip_status);
}
hipError_t ihipUnbindTextureImpl(const hipTextureObject_t& textureObject)
{
HIP_INIT_API();
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
auto ctx = ihipGetTlsDefaultCtx();
@@ -619,19 +657,20 @@ hipError_t ihipUnbindTextureImpl(const hipTextureObject_t& textureObject)
}
}
return ihipLogStatus(hip_status);
return hip_status;
}
hipError_t hipUnbindTexture(const textureReference* tex)
{
return ihipUnbindTextureImpl(tex->textureObject);
HIP_INIT_API(tex);
hipError_t hip_status = hipSuccess;
hip_status = ihipUnbindTextureImpl(tex->textureObject);
return ihipLogStatus(hip_status);
}
hipError_t hipGetChannelDesc(hipChannelFormatDesc* desc, hipArray_const_t array)
{
HIP_INIT_API(desc, array);
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
auto ctx = ihipGetTlsDefaultCtx();
@@ -644,7 +683,6 @@ hipError_t hipGetChannelDesc(hipChannelFormatDesc* desc, hipArray_const_t array)
hipError_t hipGetTextureAlignmentOffset(size_t* offset, const textureReference* tex)
{
HIP_INIT_API(offset, tex);
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
@@ -657,7 +695,6 @@ hipError_t hipGetTextureAlignmentOffset(size_t* offset, const textureReference*
hipError_t hipGetTextureReference(const textureReference** tex, const void* symbol)
{
HIP_INIT_API(tex, symbol);
HIP_SET_DEVICE();
hipError_t hip_status = hipSuccess;
@@ -666,3 +703,68 @@ hipError_t hipGetTextureReference(const textureReference** tex, const void* symb
}
return ihipLogStatus(hip_status);
}
hipError_t hipTexRefSetFormat (textureReference* tex, hipArray_Format fmt, int NumPackedComponents )
{
HIP_INIT_API(tex, fmt, NumPackedComponents);
hipError_t hip_status = hipSuccess;
tex->format = fmt;
tex->numChannels = NumPackedComponents;
return ihipLogStatus(hip_status);
}
hipError_t hipTexRefSetFlags ( textureReference* tex, unsigned int flags )
{
HIP_INIT_API(tex, flags);
hipError_t hip_status = hipSuccess;
tex->normalized = flags;
return ihipLogStatus(hip_status);
}
hipError_t hipTexRefSetFilterMode ( textureReference* tex, hipTextureFilterMode fm )
{
HIP_INIT_API(tex, fm);
hipError_t hip_status = hipSuccess;
tex->filterMode = fm;
return ihipLogStatus(hip_status);
}
hipError_t hipTexRefSetAddressMode ( textureReference* tex, int dim, hipTextureAddressMode am )
{
HIP_INIT_API(tex, dim, am);
hipError_t hip_status = hipSuccess;
tex->addressMode[dim] = am;
return ihipLogStatus(hip_status);
}
hipError_t hipTexRefSetArray ( textureReference* tex, hipArray_const_t array, unsigned int flags )
{
HIP_INIT_API(tex, array, flags);
hipError_t hip_status = hipSuccess;
hip_status = ihipBindTextureToArrayImpl(hipTextureType2D, hipReadModeElementType,
array, array->desc,tex );
return ihipLogStatus(hip_status);
}
hipError_t hipTexRefSetAddress( size_t* offset, textureReference* tex, hipDeviceptr_t devPtr, size_t size )
{
HIP_INIT_API(offset, tex, devPtr, size);
hipError_t hip_status = hipSuccess;
// TODO: hipReadModeElementType is default.
hip_status = ihipBindTextureImpl(hipTextureType1D, hipReadModeElementType,
offset, devPtr, NULL, size, tex);
return ihipLogStatus(hip_status);
}
hipError_t hipTexRefSetAddress2D( textureReference* tex, const HIP_ARRAY_DESCRIPTOR* desc, hipDeviceptr_t devPtr, size_t pitch )
{
HIP_INIT_API(tex, desc, devPtr, pitch);
size_t offset;
hipError_t hip_status = hipSuccess;
// TODO: hipReadModeElementType is default.
hip_status = ihipBindTexture2DImpl(hipTextureType2D, hipReadModeElementType,
&offset, devPtr, NULL, desc->width, desc->height, tex);
return ihipLogStatus(hip_status);
}