Introduce tensor support #70
Conversation
r-abishek
commented
Aug 31, 2021
r-abishek
commented
- Adds all RPP framework headers and changes needed for tensor support.
- Adds all unit test and performance test support for brightness tensor.
… into ar/tensor_support
src/include/hip/rpp_hip_common.hpp
Outdated
| d_half8 dst_h8; | ||
|
|
||
| dst_h8.x.x = __float22half2_rn(*(float2 *)&(dst_f8->x)); | ||
| dst_h8.x.y = __float22half2_rn(*((float2 *)&(dst_f8->x) + 1)); |
There was a problem hiding this comment.
please try to avoid this +1 in your code. Use the x, y, z, and w components of the dst_f8->x.
src/include/hip/rpp_hip_common.hpp
Outdated
|
|
||
| __device__ __forceinline__ void rpp_hip_load24_pkd3_and_unpack_to_float24_pln3(float *srcPtr, uint srcIdx, d_float24 *src_f24) | ||
| { | ||
| src_f24->x.x.x = srcPtr[srcIdx]; |
There was a problem hiding this comment.
this function needs to be rewritten. please don't access the srcPtr memory like the way you've written here by adding some offset. read a float4 each time and then use the x, y, z, w components to fill the src_f24.
src/include/hip/rpp_hip_common.hpp
Outdated
|
|
||
| __device__ __forceinline__ void rpp_hip_load24_pkd3_and_unpack_to_float24_pln3(half *srcPtr, uint srcIdx, d_float24 *src_f24) | ||
| { | ||
| src_f24->x.x.x = __half2float(srcPtr[srcIdx]); |
There was a problem hiding this comment.
use the same comments mentioned above for rpp_hip_load24_pkd3_and_unpack_to_float24_pln3 function to rewrite this function as well.
src/include/hip/rpp_hip_common.hpp
Outdated
| srcPtrG = srcPtrR + increment; | ||
| srcPtrB = srcPtrG + increment; | ||
|
|
||
| src_f24->x.x.x = *srcPtrR; |
There was a problem hiding this comment.
use vector float4 for reading srcPtr memory and use its components when filling the src_f24.
include/rppdefs.h
Outdated
| typedef enum | ||
| { | ||
| rppStatusSuccess = 0, | ||
| rppStatusBadParm = 1, |
There was a problem hiding this comment.
it is better to give -ve values for error codes
|
@asalmanp The I8 images now match exactly with the F16/F32/U8 visually. Please let me know if you have any other thoughts too. We could now proceed for the final testing and merge @paveltc |
|
|
||
| __device__ void brightness_hip_compute(signed char *srcPtr, d_float8 *src_f8, d_float8 *dst_f8, float4 *alpha_f4, float4 *beta_f4) | ||
| { | ||
| dst_f8->x = rpp_hip_pixel_check((src_f8->x + 128) * *alpha_f4 + *beta_f4) - 128; |
There was a problem hiding this comment.
src_f8->x is a float4 so if your intention is to add each of the x/y/z/w with 128 then you need to use (float4)128 here.
similarly, you would need to use (float4)128 to subtract in the end of this line.
There was a problem hiding this comment.
Thats true, but I believe it was doing an auto typecast to float4. In any case, I have added a manual typecast now to be on the safer side.
| { | ||
| int id_x = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x) * 4; | ||
|
|
||
| roiTensorPtrSrc[id_x + 2] -= (roiTensorPtrSrc[id_x] - 1); |
There was a problem hiding this comment.
use int2 to read from the roiTensorPtrSrc instead of using +1
There was a problem hiding this comment.
Thanks! I have changed it now. Also, here roiTensorPtrSrc has 0/1/2/3, so I have used int4.
src/modules/rppi_validate.hpp
Outdated
| return layoutParams; | ||
| } | ||
|
|
||
| // inline void copy_roiTensor(RpptROIPtr roiTensorPtrSrc, rpp::Handle& handle) |
There was a problem hiding this comment.
remove commented out codes if it is not needed.
There was a problem hiding this comment.
Done. Cleaned up this rppi_validate file for just some formatting changes, tabs to spaces, and removal of unwanted commented code.
|
|
||
| __device__ __forceinline__ float rpp_hip_unpack0(int src) | ||
| { | ||
| return (float)(signed char)(src & 0xFF); |
There was a problem hiding this comment.
why the casting to (signed char) is needed here?
There was a problem hiding this comment.
Actually I found this to be the only way to do the unpack correctly. It takes src, gets the first byte by masking with 0xFF, then we explicitly tell the compiler to interpret the MSB as a signed bit, then convert to float. If I directly convert to float, the most significant bit was interpreted not as a sign, and as part of an 8 bit number.
|
@r-abishek please take a look at my latest comments and make the necessary changes. |
|
@paveltc please run rali and rpp unit-tests on this PR so we can merge it |
|
@rrawther @asalmanp This PR passes the unit tests. |
@asalmanp @kiritigowda Could we merge the PR#70 too? The unit tests have been verified. |
