AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

SWDEV-542815 - Fix unused-parameter compiler warnings (#719)

Просмотр исходного кода 
* SWDEV-542815 - Fix unused-parameter compiler warnings

* SWDEV-542815 - Run clang-format

* SWDEV-542815 - Run clang-format

* SWDEV-542815 - Run clang-format

* SWDEV-542815 - Run clang-format

[ROCm/clr commit: 1385d159ba]
Этот коммит содержится в:
Arandjelovic, Marko
2025-08-15 22:10:23 +02:00
коммит произвёл GitHub
родитель 0d63a5abd3
Коммит e91332340a
3 изменённых файлов: 69 добавлений и 41 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
projects/clr/hipamd/include/hip/amd_detail/amd_surface_functions.h
+9 -8
Просмотреть файл
@@ -114,10 +114,11 @@ template <
typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf1Dread(T* data, hipSurfaceObject_t surfObj, int x,
int boundaryMode = hipBoundaryModeZero) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
auto tmp = __ockl_image_load_1D(i, x);
*data = __hipMapFrom<T>(tmp);
__HIP_SURFACE_OBJECT_PARAMETERS_INIT;
(void)boundaryMode;
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
auto tmp = __ockl_image_load_1D(i, x);
*data = __hipMapFrom<T>(tmp);
}
/** \brief Writes the value data to the one-dimensional surface at coordinate x.
@@ -157,7 +158,7 @@ static __device__ __hip_img_chk__ void surf2Dread(T* data, hipSurfaceObject_t su
*data = __hipMapFrom<T>(tmp);
}
/** \brief Writes the value data to the two-dimensional surface at coordinate
/** \brief Writes the value data to the two-dimensional surface at coordinate
* x, y.
*
* \tparam T The data type of the surface.
@@ -177,7 +178,7 @@ static __device__ __hip_img_chk__ void surf2Dwrite(T data, hipSurfaceObject_t su
__ockl_image_store_2D(i, get_native_vector(coords), tmp);
}
/** \brief Reads the value from the three-dimensional surface at coordinate
/** \brief Reads the value from the three-dimensional surface at coordinate
* x, y, z.
*
* \tparam T The data type of the surface.
@@ -238,7 +239,7 @@ static __device__ __hip_img_chk__ void surf1DLayeredread(T* data, hipSurfaceObje
*data = __hipMapFrom<T>(tmp);
}
/** \brief Writes the value data to the one-dimensional layered surface at
/** \brief Writes the value data to the one-dimensional layered surface at
* coordinate x and layer index.
*
* \tparam T The data type of the surface.
@@ -257,7 +258,7 @@ static __device__ __hip_img_chk__ void surf1DLayeredwrite(T data, hipSurfaceObje
__ockl_image_store_lod_1D(i, x, layer, tmp);
}
/** \brief Reads the value from the two-dimensional layered surface at
/** \brief Reads the value from the two-dimensional layered surface at
* coordinate x, y and layer index.
*
* \tparam T The data type of the surface.
projects/clr/hipamd/include/hip/amd_detail/texture_fetch_functions.h
+15 -3
Просмотреть файл
@@ -31,9 +31,10 @@ THE SOFTWARE.
#include <type_traits>
#endif // !defined(__HIPCC_RTC__)
#define TEXTURE_PARAMETERS_INIT \
unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)t.textureObject; \
unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD;
#define TEXTURE_PARAMETERS_INIT \
unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)t.textureObject; \
unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD; \
(void)s;
template<typename T>
struct __hip_is_tex_surf_scalar_channel_type
@@ -376,6 +377,11 @@ template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> texCubemapGrad(texture<T, hipTextureTypeCubemap, readMode> t, float x, float y, float z, float4 dPdx, float4 dPdy)
{
TEXTURE_PARAMETERS_INIT;
(void)x;
(void)y;
(void)z;
(void)dPdx;
(void)dPdy;
// TODO missing in device libs.
// auto tmp = __ockl_image_sample_grad_CM(i, s, get_native_vector(float4(x, y, z, 0.0f)),
// get_native_vector(float4(dPdx.x, dPdx.y, dPdx.z, 0.0f)), get_native_vector(float4(dPdy.x,
@@ -387,6 +393,12 @@ template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> texCubemapLayeredGrad(texture<T, hipTextureTypeCubemapLayered, readMode> t, float x, float y, float z, int layer, float4 dPdx, float4 dPdy)
{
TEXTURE_PARAMETERS_INIT;
(void)x;
(void)y;
(void)z;
(void)layer;
(void)dPdx;
(void)dPdy;
// TODO missing in device libs.
// auto tmp = __ockl_image_sample_grad_CMa(i, s, get_native_vector(float4(x, y, z, layer)),
// get_native_vector(float4(dPdx.x, dPdx.y, dPdx.z, 0.0f)), get_native_vector(float4(dPdy.x,
projects/clr/hipamd/include/hip/amd_detail/texture_indirect_functions.h
+45 -30
Просмотреть файл
@@ -32,9 +32,10 @@ THE SOFTWARE.
#include <type_traits>
#endif // !defined(__HIPCC_RTC__)
#define TEXTURE_OBJECT_PARAMETERS_INIT \
unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)textureObject; \
unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD;
#define TEXTURE_OBJECT_PARAMETERS_INIT \
unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)textureObject; \
unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD; \
(void)s;
template <
typename T,
@@ -287,10 +288,11 @@ template <
typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex1DLayeredLod(hipTextureObject_t textureObject, float x, int layer, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
float2 coords{x, layer};
auto tmp = __ockl_image_sample_1Da(i, s, get_native_vector(coords));
return __hipMapFrom<T>(tmp);
TEXTURE_OBJECT_PARAMETERS_INIT;
(void)level;
float2 coords{x, layer};
auto tmp = __ockl_image_sample_1Da(i, s, get_native_vector(coords));
return __hipMapFrom<T>(tmp);
}
template <
@@ -306,10 +308,11 @@ template <
typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex2DLayeredLod(hipTextureObject_t textureObject, float x, float y, int layer, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
float4 coords{x, y, layer, 0.0f};
auto tmp = __ockl_image_sample_2Da(i, s, get_native_vector(coords));
return __hipMapFrom<T>(tmp);
TEXTURE_OBJECT_PARAMETERS_INIT;
(void)level;
float4 coords{x, y, layer, 0.0f};
auto tmp = __ockl_image_sample_2Da(i, s, get_native_vector(coords));
return __hipMapFrom<T>(tmp);
}
template <
@@ -344,12 +347,17 @@ template <
typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T texCubemapGrad(hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
// TODO missing in device libs.
// auto tmp = __ockl_image_sample_grad_CM(i, s, get_native_vector(float4(x, y, z, 0.0f)),
// get_native_vector(float4(dPdx.x, dPdx.y, dPdx.z, 0.0f)), get_native_vector(float4(dPdy.x,
// dPdy.y, dPdy.z, 0.0f))); return __hipMapFrom<T>(tmp);
return {};
TEXTURE_OBJECT_PARAMETERS_INIT;
(void)x;
(void)y;
(void)z;
(void)dPdx;
(void)dPdy;
// TODO missing in device libs.
// auto tmp = __ockl_image_sample_grad_CM(i, s, get_native_vector(float4(x, y, z, 0.0f)),
// get_native_vector(float4(dPdx.x, dPdx.y, dPdx.z, 0.0f)), get_native_vector(float4(dPdy.x,
// dPdy.y, dPdy.z, 0.0f))); return __hipMapFrom<T>(tmp);
return {};
}
template <
@@ -422,13 +430,14 @@ template <
typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex3DGrad(hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
float4 coords{x, y, z, 0.0f};
float4 gradx{dPdy.x, dPdy.y, dPdy.z, 0.0f};
float4 grady{dPdy.x, dPdy.y, dPdy.z, 0.0f};
auto tmp = __ockl_image_sample_grad_3D(i, s, get_native_vector(coords),
get_native_vector(gradx), get_native_vector(grady));
return __hipMapFrom<T>(tmp);
TEXTURE_OBJECT_PARAMETERS_INIT;
(void)dPdx;
float4 coords{x, y, z, 0.0f};
float4 gradx{dPdy.x, dPdy.y, dPdy.z, 0.0f};
float4 grady{dPdy.x, dPdy.y, dPdy.z, 0.0f};
auto tmp = __ockl_image_sample_grad_3D(i, s, get_native_vector(coords), get_native_vector(gradx),
get_native_vector(grady));
return __hipMapFrom<T>(tmp);
}
template <
@@ -483,12 +492,18 @@ template <
typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T texCubemapLayeredGrad(hipTextureObject_t textureObject, float x, float y, float z, int layer, float4 dPdx, float4 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
// TODO missing in device libs.
// auto tmp = __ockl_image_sample_grad_CMa(i, s, get_native_vector(float4(x, y, z, layer)),
// get_native_vector(float4(dPdx.x, dPdx.y, dPdx.z, 0.0f)), get_native_vector(float4(dPdy.x,
// dPdy.y, dPdy.z, 0.0f))); return __hipMapFrom<T>(tmp);
return {};
TEXTURE_OBJECT_PARAMETERS_INIT;
(void)x;
(void)y;
(void)z;
(void)layer;
(void)dPdx;
(void)dPdy;
// TODO missing in device libs.
// auto tmp = __ockl_image_sample_grad_CMa(i, s, get_native_vector(float4(x, y, z, layer)),
// get_native_vector(float4(dPdx.x, dPdx.y, dPdx.z, 0.0f)), get_native_vector(float4(dPdy.x,
// dPdy.y, dPdy.z, 0.0f))); return __hipMapFrom<T>(tmp);
return {};
}
template <