AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Add support for creating typed buffers

Browse Source 
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
This commit is contained in:
Vladislav Sytchenko
2020-03-18 15:30:19 -04:00
committed by Vladislav Sytchenko
parent ea7170d33c
commit 78f7954765
5 changed files with 6 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/include/hip/hcc_detail/ockl_image.h
+2
View File
@@ -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);
hipamd/include/hip/hcc_detail/texture_fetch_functions.h
+1 -1
View File
@@ -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);
}
hipamd/include/hip/hcc_detail/texture_indirect_functions.h
+1 -1
View File
@@ -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);
}
hipamd/vdi/hip_conversions.hpp
+1 -1
View File
@@ -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:
hipamd/vdi/hip_memory.cpp
+1 -1
View File
@@ -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,