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

Merge pull request #382 from scchan/fix_texture_addr_space

Просмотр исходного кода 
Change constant address space to 4
Этот коммит содержится в:
Maneesh Gupta
2018-03-27 07:48:50 +05:30
коммит произвёл GitHub
родитель a95ff17549 53d9cce9f9
Коммит c1505becef
1 изменённых файлов: 36 добавлений и 29 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
hipamd/include/hip/hcc_detail/texture_functions.h
+36 -29
Просмотреть файл
@@ -34,14 +34,21 @@ union TData {
};
#define __TEXTURE_FUNCTIONS_DECL__ static __inline__ __device__
#define ADDRESS_SPACE_2 __attribute__((address_space(2)))
#if (__hcc_workweek__ >= 18115)
#define ADDRESS_SPACE_CONSTANT __attribute__((address_space(4)))
#else
#define ADDRESS_SPACE_CONSTANT __attribute__((address_space(2)))
#endif
#define TEXTURE_PARAMETERS_INIT \
unsigned int ADDRESS_SPACE_2* i = (unsigned int ADDRESS_SPACE_2*)textureObject; \
unsigned int ADDRESS_SPACE_2* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD; \
unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)textureObject; \
unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD; \
TData texel;
#define TEXTURE_REF_PARAMETERS_INIT \
unsigned int ADDRESS_SPACE_2* i = (unsigned int ADDRESS_SPACE_2*)texRef.textureObject; \
unsigned int ADDRESS_SPACE_2* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD; \
#define TEXTURE_REF_PARAMETERS_INIT \
unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)texRef.textureObject; \
unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD; \
TData texel;
#define TEXTURE_SET_FLOAT *retVal = texel.f.x;
@@ -146,90 +153,90 @@ union TData {
#define TEXTURE_RETURN_FLOAT_XYZW return float4(texel.f.x, texel.f.y, texel.f.z, texel.f.w);
extern "C" {
hc::short_vector::float4::vector_value_type __ockl_image_sample_1D(unsigned int ADDRESS_SPACE_2* i,
unsigned int ADDRESS_SPACE_2* s,
hc::short_vector::float4::vector_value_type __ockl_image_sample_1D(unsigned int ADDRESS_SPACE_CONSTANT* i,
unsigned int ADDRESS_SPACE_CONSTANT* s,
float c)[[hc]];
hc::short_vector::float4::vector_value_type __ockl_image_sample_1Da(
unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s,
unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s,
hc::short_vector::float2::vector_value_type c)[[hc]];
hc::short_vector::float4::vector_value_type __ockl_image_sample_2D(
unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s,
unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s,
hc::short_vector::float2::vector_value_type c)[[hc]];
hc::short_vector::float4::vector_value_type __ockl_image_sample_2Da(
unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s,
unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s,
hc::short_vector::float4::vector_value_type c)[[hc]];
float __ockl_image_sample_2Dad(unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s,
float __ockl_image_sample_2Dad(unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s,
hc::short_vector::float4::vector_value_type c)[[hc]];
float __ockl_image_sample_2Dd(unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s,
float __ockl_image_sample_2Dd(unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s,
hc::short_vector::float2::vector_value_type c)[[hc]];
hc::short_vector::float4::vector_value_type __ockl_image_sample_3D(
unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s,
unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s,
hc::short_vector::float4::vector_value_type c)[[hc]];
hc::short_vector::float4::vector_value_type __ockl_image_sample_grad_1D(
unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s, float c, float dx,
unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s, float c, float dx,
float dy)[[hc]];
hc::short_vector::float4::vector_value_type __ockl_image_sample_grad_1Da(
unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s,
unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s,
hc::short_vector::float2::vector_value_type c, float dx, float dy)[[hc]];
hc::short_vector::float4::vector_value_type __ockl_image_sample_grad_2D(
unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s,
unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s,
hc::short_vector::float2::vector_value_type c, hc::short_vector::float2::vector_value_type dx,
hc::short_vector::float2::vector_value_type dy)[[hc]];
hc::short_vector::float4::vector_value_type __ockl_image_sample_grad_2Da(
unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s,
unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s,
hc::short_vector::float4::vector_value_type c, hc::short_vector::float2::vector_value_type dx,
hc::short_vector::float2::vector_value_type dy)[[hc]];
float __ockl_image_sample_grad_2Dad(unsigned int ADDRESS_SPACE_2* i,
unsigned int ADDRESS_SPACE_2* s,
float __ockl_image_sample_grad_2Dad(unsigned int ADDRESS_SPACE_CONSTANT* i,
unsigned int ADDRESS_SPACE_CONSTANT* s,
hc::short_vector::float4::vector_value_type c,
hc::short_vector::float2::vector_value_type dx,
hc::short_vector::float2::vector_value_type dy)[[hc]];
float __ockl_image_sample_grad_2Dd(unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s,
float __ockl_image_sample_grad_2Dd(unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s,
hc::short_vector::float2::vector_value_type c,
hc::short_vector::float2::vector_value_type dx,
hc::short_vector::float2::vector_value_type dy)[[hc]];
hc::short_vector::float4::vector_value_type __ockl_image_sample_grad_3D(
unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s,
unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s,
hc::short_vector::float4::vector_value_type c, hc::short_vector::float4::vector_value_type dx,
hc::short_vector::float4::vector_value_type dy)[[hc]];
hc::short_vector::float4::vector_value_type __ockl_image_sample_lod_1D(
unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s, float c, float l)[[hc]];
unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s, float c, float l)[[hc]];
hc::short_vector::float4::vector_value_type __ockl_image_sample_lod_1Da(
unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s,
unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s,
hc::short_vector::float2::vector_value_type c, float l)[[hc]];
hc::short_vector::float4::vector_value_type __ockl_image_sample_lod_2D(
unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s,
unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s,
hc::short_vector::float2::vector_value_type c, float l)[[hc]];
hc::short_vector::float4::vector_value_type __ockl_image_sample_lod_2Da(
unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s,
unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s,
hc::short_vector::float4::vector_value_type c, float l)[[hc]];
float __ockl_image_sample_lod_2Dad(unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s,
float __ockl_image_sample_lod_2Dad(unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s,
hc::short_vector::float4::vector_value_type c, float l)[[hc]];
float __ockl_image_sample_lod_2Dd(unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s,
float __ockl_image_sample_lod_2Dd(unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s,
hc::short_vector::float2::vector_value_type c, float l)[[hc]];
hc::short_vector::float4::vector_value_type __ockl_image_sample_lod_3D(
unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s,
unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s,
hc::short_vector::float4::vector_value_type c, float l)[[hc]];
}