AI/rocm-systems
2
0
Forkuj 0
Kod Zgłoszenia Oczekujące zmiany Actions Packages Projekty Wydania Wiki Aktywność

SWDEV-345188 - hipamd: warnings treated as errors for catch tests

Przeglądaj źródła 
Change-Id: I8981a530da2e0213c60f7a81b09d9c36e27e9025
This commit is contained in:
Ajay
2022-07-06 22:03:12 +00:00
zatwierdzone przez Ajay GunaShekar
rodzic 4274863f5d
commit 9eb45dd1e0
2 zmienionych plików z 18 dodań i 33 usunięć
Ściągnij plik aktualizacji Ściągnij plik porównania Expand all files Collapse all files
hipamd/include/hip/amd_detail/amd_device_functions.h
+5 -9
Wyświetl plik
@@ -101,7 +101,7 @@ __device__ static int32_t __fns64(uint64_t mask, uint32_t base, int32_t offset)
int32_t total = 0;
for (int i = 0x20; i > 0; i >>= 1) {
uint64_t temp_mask_lo = temp_mask & ((1ULL << i) - 1);
uint32_t pcnt = __builtin_popcountll(temp_mask_lo);
int32_t pcnt = __builtin_popcountll(temp_mask_lo);
if (pcnt < temp_offset) {
temp_mask = temp_mask >> i;
temp_offset -= pcnt;
@@ -135,7 +135,7 @@ __device__ static int32_t __fns32(uint64_t mask, uint32_t base, int32_t offset)
int32_t total = 0;
for (int i = 0x20; i > 0; i >>= 1) {
uint64_t temp_mask_lo = temp_mask & ((1ULL << i) - 1);
uint32_t pcnt = __builtin_popcountll(temp_mask_lo);
int32_t pcnt = __builtin_popcountll(temp_mask_lo);
if (pcnt < temp_offset) {
temp_mask = temp_mask >> i;
temp_offset -= pcnt;
@@ -653,7 +653,7 @@ __device__ long long int __clock();
__device__ long long int clock64();
__device__ long long int clock();
// hip.amdgcn.bc - named sync
__device__ void __named_sync(int a, int b);
__device__ void __named_sync();
#ifdef __HIP_DEVICE_COMPILE__
@@ -700,7 +700,7 @@ long long int clock() { return __clock(); }
// hip.amdgcn.bc - named sync
__device__
inline
void __named_sync(int a, int b) { __builtin_amdgcn_s_barrier(); }
void __named_sync() { __builtin_amdgcn_s_barrier(); }
#endif // __HIP_DEVICE_COMPILE__
@@ -871,11 +871,7 @@ void __assert_fail(const char *assertion,
}
extern "C" __device__ __attribute__((noinline)) __attribute__((weak))
void __assertfail(const char *assertion,
const char *file,
unsigned int line,
const char *function,
size_t charsize)
void __assertfail()
{
// ignore all the args for now.
__builtin_trap();
hipamd/include/hip/amd_detail/amd_surface_functions.h
+13 -24
Wyświetl plik
@@ -229,8 +229,7 @@ static __device__ __hip_img_chk__ void surf1Dread(T* data, hipSurfaceObject_t su
template <
typename T,
typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf1Dwrite(T data, hipSurfaceObject_t surfObj, int x,
int boundaryMode = hipBoundaryModeZero) {
static __device__ __hip_img_chk__ void surf1Dwrite(T data, hipSurfaceObject_t surfObj, int x) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
auto tmp = __hipMapToNativeFloat4(data);
@@ -240,8 +239,7 @@ static __device__ __hip_img_chk__ void surf1Dwrite(T data, hipSurfaceObject_t su
template <
typename T,
typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf2Dread(T* data, hipSurfaceObject_t surfObj, int x, int y,
int boundaryMode = hipBoundaryModeZero) {
static __device__ __hip_img_chk__ void surf2Dread(T* data, hipSurfaceObject_t surfObj, int x, int y) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
auto tmp = __ockl_image_load_2D(i, int2(x, y).data);
@@ -251,8 +249,7 @@ static __device__ __hip_img_chk__ void surf2Dread(T* data, hipSurfaceObject_t su
template <
typename T,
typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf2Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y,
int boundaryMode = hipBoundaryModeZero) {
static __device__ __hip_img_chk__ void surf2Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
auto tmp = __hipMapToNativeFloat4(data);
@@ -262,8 +259,7 @@ static __device__ __hip_img_chk__ void surf2Dwrite(T data, hipSurfaceObject_t su
template <
typename T,
typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf3Dread(T* data, hipSurfaceObject_t surfObj, int x, int y, int z,
int boundaryMode = hipBoundaryModeZero) {
static __device__ __hip_img_chk__ void surf3Dread(T* data, hipSurfaceObject_t surfObj, int x, int y, int z) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i));
auto tmp = __ockl_image_load_3D(i, int4(x, y, z, 0).data);
@@ -273,8 +269,7 @@ static __device__ __hip_img_chk__ void surf3Dread(T* data, hipSurfaceObject_t su
template <
typename T,
typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf3Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int z,
int boundaryMode = hipBoundaryModeZero) {
static __device__ __hip_img_chk__ void surf3Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int z) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_3D(i), __ockl_image_channel_order_3D(i));
auto tmp = __hipMapToNativeFloat4(data);
@@ -284,8 +279,7 @@ static __device__ __hip_img_chk__ void surf3Dwrite(T data, hipSurfaceObject_t su
template <
typename T,
typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf1DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int layer,
int boundaryMode = hipBoundaryModeZero) {
static __device__ __hip_img_chk__ void surf1DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int layer) {
__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_lod_1D(i, x, layer);
@@ -295,8 +289,7 @@ static __device__ __hip_img_chk__ void surf1DLayeredread(T* data, hipSurfaceObje
template <
typename T,
typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf1DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int layer,
int boundaryMode = hipBoundaryModeZero) {
static __device__ __hip_img_chk__ void surf1DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int layer) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
auto tmp = __hipMapToNativeFloat4(data);
@@ -306,8 +299,7 @@ static __device__ __hip_img_chk__ void surf1DLayeredwrite(T data, hipSurfaceObje
template <
typename T,
typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf2DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int y, int layer,
int boundaryMode = hipBoundaryModeZero) {
static __device__ __hip_img_chk__ void surf2DLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int y, int layer) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
auto tmp = __ockl_image_load_lod_2D(i, int2(x, y).data, layer);
@@ -317,8 +309,7 @@ static __device__ __hip_img_chk__ void surf2DLayeredread(T* data, hipSurfaceObje
template <
typename T,
typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf2DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int layer,
int boundaryMode = hipBoundaryModeZero) {
static __device__ __hip_img_chk__ void surf2DLayeredwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int layer) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
auto tmp = __hipMapToNativeFloat4(data);
@@ -328,8 +319,7 @@ static __device__ __hip_img_chk__ void surf2DLayeredwrite(T data, hipSurfaceObje
template <
typename T,
typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surfCubemapread(T* data, hipSurfaceObject_t surfObj, int x, int y, int face,
int boundaryMode = hipBoundaryModeZero) {
static __device__ __hip_img_chk__ void surfCubemapread(T* data, hipSurfaceObject_t surfObj, int x, int y, int face) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
auto tmp = __ockl_image_load_CM(i, int2(x, y).data, face);
@@ -339,8 +329,7 @@ static __device__ __hip_img_chk__ void surfCubemapread(T* data, hipSurfaceObject
template <
typename T,
typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surfCubemapwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int face,
int boundaryMode = hipBoundaryModeZero) {
static __device__ __hip_img_chk__ void surfCubemapwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int face) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
auto tmp = __hipMapToNativeFloat4(data);
@@ -351,7 +340,7 @@ template <
typename T,
typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surfCubemapLayeredread(T* data, hipSurfaceObject_t surfObj, int x, int y, int face,
int layer, int boundaryMode = hipBoundaryModeZero) {
int layer) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
auto tmp = __ockl_image_load_lod_CM(i, int2(x, y).data, face, layer);
@@ -362,7 +351,7 @@ template <
typename T,
typename std::enable_if<__hip_is_isurf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surfCubemapLayeredwrite(T* data, hipSurfaceObject_t surfObj, int x, int y, int face,
int layer, int boundaryMode = hipBoundaryModeZero) {
int layer) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_2D(i), __ockl_image_channel_order_2D(i));
auto tmp = __hipMapToNativeFloat4(data);