From 4829a7c2156f5b87d9fb345fe63f4bed4f85f5db Mon Sep 17 00:00:00 2001 From: Vladislav Sytchenko Date: Wed, 18 Mar 2020 15:30:19 -0400 Subject: [PATCH] Add support for creating typed buffers What Cuda refers to "linear texture memory" is the OpenCL equivalent of CL_MEM_OBJECT_IMAGE1D_BUFFER. For these types of allocations we should create a typed buffer instead of an image. Currently there is no check in the texture fetch functions as to what kind of SRD is written into the texture object, so any kind of incorrect programming will cause the TA to hang. Fortunately for us, every one writes correct code :) Change-Id: I80dab85a992f2c0754ebf303d40ac6b5e045c7c1 --- include/hip/hcc_detail/ockl_image.h | 2 ++ include/hip/hcc_detail/texture_fetch_functions.h | 2 +- include/hip/hcc_detail/texture_indirect_functions.h | 2 +- vdi/hip_conversions.hpp | 2 +- vdi/hip_memory.cpp | 2 +- 5 files changed, 6 insertions(+), 4 deletions(-) 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,