diff --git a/include/hip/hcc_detail/ockl_image.h b/include/hip/hcc_detail/ockl_image.h index d9bc296791..b32b23fda0 100644 --- a/include/hip/hcc_detail/ockl_image.h +++ b/include/hip/hcc_detail/ockl_image.h @@ -30,6 +30,8 @@ extern "C" { __device__ float4::Native_vec_ __ockl_image_load_1D(unsigned int ADDRESS_SPACE_CONSTANT*i, int c); +__device__ float4::Native_vec_ __ockl_image_load_1Db(unsigned int ADDRESS_SPACE_CONSTANT*i, int c); + __device__ float4::Native_vec_ __ockl_image_load_1Da(unsigned int ADDRESS_SPACE_CONSTANT*i, int2::Native_vec_ c); __device__ float4::Native_vec_ __ockl_image_load_2D(unsigned int ADDRESS_SPACE_CONSTANT*i, int2::Native_vec_ c); diff --git a/include/hip/hcc_detail/texture_fetch_functions.h b/include/hip/hcc_detail/texture_fetch_functions.h index 23a850c52c..9e8e864520 100644 --- a/include/hip/hcc_detail/texture_fetch_functions.h +++ b/include/hip/hcc_detail/texture_fetch_functions.h @@ -139,7 +139,7 @@ template static __forceinline__ __device__ __hip_tex_ret_t tex1Dfetch(texture t, int x) { TEXTURE_PARAMETERS_INIT; - auto tmp = __ockl_image_sample_1D(i, s, x); + auto tmp = __ockl_image_load_1Db(i, x); return *reinterpret_cast<__hip_tex_ret_t*>(&tmp); } diff --git a/include/hip/hcc_detail/texture_indirect_functions.h b/include/hip/hcc_detail/texture_indirect_functions.h index 4facce436e..2fe33f3ede 100644 --- a/include/hip/hcc_detail/texture_indirect_functions.h +++ b/include/hip/hcc_detail/texture_indirect_functions.h @@ -65,7 +65,7 @@ template < static __device__ T tex1Dfetch(hipTextureObject_t textureObject, int x) { TEXTURE_OBJECT_PARAMETERS_INIT - auto tmp = __ockl_image_sample_1D(i, s, x); + auto tmp = __ockl_image_load_1Db(i, x); return *reinterpret_cast(&tmp); } diff --git a/vdi/hip_conversions.hpp b/vdi/hip_conversions.hpp index 0801c310eb..3b43cae2b1 100644 --- a/vdi/hip_conversions.hpp +++ b/vdi/hip_conversions.hpp @@ -153,7 +153,7 @@ inline cl_mem_object_type getCLMemObjectType(const hipResourceType hipResType) { switch (hipResType) { case hipResourceTypeLinear: - return CL_MEM_OBJECT_IMAGE1D; + return CL_MEM_OBJECT_IMAGE1D_BUFFER; case hipResourceTypePitch2D: return CL_MEM_OBJECT_IMAGE2D; default: diff --git a/vdi/hip_memory.cpp b/vdi/hip_memory.cpp index 55bc11512a..6233c15a38 100644 --- a/vdi/hip_memory.cpp +++ b/vdi/hip_memory.cpp @@ -444,7 +444,7 @@ amd::Image* ihipImageCreate(const cl_channel_order channelOrder, amd::Image* image = nullptr; if (buffer != nullptr) { switch (imageType) { - case CL_MEM_OBJECT_IMAGE1D: + case CL_MEM_OBJECT_IMAGE1D_BUFFER: case CL_MEM_OBJECT_IMAGE2D: image = new (context) amd::Image(*buffer->asBuffer(), imageType,