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

@@ -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);

include/hip/hcc_detail/texture_fetch_functions.h

@@ -139,7 +139,7 @@ template <typename T, hipTextureReadMode readMode>
 static __forceinline__ __device__ __hip_tex_ret_t<T, readMode> tex1Dfetch(texture<T, hipTextureType1D, readMode> 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<T, readMode>*>(&tmp);
 }
 

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<T*>(&tmp);
 }
 

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:

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,