25 #include <hip/hip_vector_types.h>
29 #define ADDRESS_SPACE_CONSTANT __attribute__((address_space(4)))
31 __device__ float4::Native_vec_ __ockl_image_load_1D(
unsigned int ADDRESS_SPACE_CONSTANT*i,
int c);
33 __device__ float4::Native_vec_ __ockl_image_load_1Db(
unsigned int ADDRESS_SPACE_CONSTANT*i,
int c);
35 __device__ float4::Native_vec_ __ockl_image_load_1Da(
unsigned int ADDRESS_SPACE_CONSTANT*i, int2::Native_vec_ c);
37 __device__ float4::Native_vec_ __ockl_image_load_2D(
unsigned int ADDRESS_SPACE_CONSTANT*i, int2::Native_vec_ c);
39 __device__ float4::Native_vec_ __ockl_image_load_2Da(
unsigned int ADDRESS_SPACE_CONSTANT*i, int4::Native_vec_ c);
41 __device__ float4::Native_vec_ __ockl_image_load_3D(
unsigned int ADDRESS_SPACE_CONSTANT*i, int4::Native_vec_ c);
43 __device__ float4::Native_vec_ __ockl_image_load_CM(
unsigned int ADDRESS_SPACE_CONSTANT*i, int2::Native_vec_ c,
int f);
45 __device__ float4::Native_vec_ __ockl_image_load_CMa(
unsigned int ADDRESS_SPACE_CONSTANT*i, int4::Native_vec_ c,
int f);
47 __device__ float4::Native_vec_ __ockl_image_load_lod_1D(
unsigned int ADDRESS_SPACE_CONSTANT*i,
int c,
int l);
49 __device__ float4::Native_vec_ __ockl_image_load_lod_1Da(
unsigned int ADDRESS_SPACE_CONSTANT*i, int2::Native_vec_ c,
int l);
51 __device__ float4::Native_vec_ __ockl_image_load_lod_2D(
unsigned int ADDRESS_SPACE_CONSTANT*i, int2::Native_vec_ c,
int l);
53 __device__ float4::Native_vec_ __ockl_image_load_lod_2Da(
unsigned int ADDRESS_SPACE_CONSTANT*i, int4::Native_vec_ c,
int l);
55 __device__ float4::Native_vec_ __ockl_image_load_lod_3D(
unsigned int ADDRESS_SPACE_CONSTANT*i, int4::Native_vec_ c,
int l);
57 __device__ float4::Native_vec_ __ockl_image_load_lod_CM(
unsigned int ADDRESS_SPACE_CONSTANT*i, int2::Native_vec_ c,
int f,
int l);
59 __device__ float4::Native_vec_ __ockl_image_load_lod_CMa(
unsigned int ADDRESS_SPACE_CONSTANT*i, int4::Native_vec_ c,
int f,
int l);
61 __device__
void __ockl_image_store_1D(
unsigned int ADDRESS_SPACE_CONSTANT*i,
int c, float4::Native_vec_ p);
63 __device__
void __ockl_image_store_1Da(
unsigned int ADDRESS_SPACE_CONSTANT*i, int2::Native_vec_ c, float4::Native_vec_ p);
65 __device__
void __ockl_image_store_2D(
unsigned int ADDRESS_SPACE_CONSTANT*i, int2::Native_vec_ c, float4::Native_vec_ p);
67 __device__
void __ockl_image_store_2Da(
unsigned int ADDRESS_SPACE_CONSTANT*i, int4::Native_vec_ c, float4::Native_vec_ p);
69 __device__
void __ockl_image_store_3D(
unsigned int ADDRESS_SPACE_CONSTANT*i, int4::Native_vec_ c, float4::Native_vec_ p);
71 __device__
void __ockl_image_store_CM(
unsigned int ADDRESS_SPACE_CONSTANT*i, int4::Native_vec_ c, float4::Native_vec_ p);
73 __device__
void __ockl_image_store_CMa(
unsigned int ADDRESS_SPACE_CONSTANT*i, int4::Native_vec_ c, float4::Native_vec_ p);
75 __device__
void __ockl_image_store_lod_1D(
unsigned int ADDRESS_SPACE_CONSTANT*i,
int c,
int l, float4::Native_vec_ p);
77 __device__
void __ockl_image_store_lod_1Da(
unsigned int ADDRESS_SPACE_CONSTANT*i, int2::Native_vec_ c,
int l, float4::Native_vec_ p);
79 __device__
void __ockl_image_store_lod_2D(
unsigned int ADDRESS_SPACE_CONSTANT*i, int2::Native_vec_ c,
int l, float4::Native_vec_ p);
81 __device__
void __ockl_image_store_lod_2Da(
unsigned int ADDRESS_SPACE_CONSTANT*i, int4::Native_vec_ c,
int l, float4::Native_vec_ p);
83 __device__
void __ockl_image_store_lod_3D(
unsigned int ADDRESS_SPACE_CONSTANT*i, int4::Native_vec_ c,
int l, float4::Native_vec_ p);
85 __device__
void __ockl_image_store_lod_CM(
unsigned int ADDRESS_SPACE_CONSTANT*i, int4::Native_vec_ c,
int l, float4::Native_vec_ p);
87 __device__
void __ockl_image_store_lod_CMa(
unsigned int ADDRESS_SPACE_CONSTANT*i, int4::Native_vec_ c,
int l, float4::Native_vec_ p);
89 __device__ float4::Native_vec_ __ockl_image_sample_1D(
unsigned int ADDRESS_SPACE_CONSTANT*i,
unsigned int ADDRESS_SPACE_CONSTANT*s,
float c);
91 __device__ float4::Native_vec_ __ockl_image_sample_1Da(
unsigned int ADDRESS_SPACE_CONSTANT*i,
unsigned int ADDRESS_SPACE_CONSTANT*s, float2::Native_vec_ c);
93 __device__ float4::Native_vec_ __ockl_image_sample_2D(
unsigned int ADDRESS_SPACE_CONSTANT*i,
unsigned int ADDRESS_SPACE_CONSTANT*s, float2::Native_vec_ c);
95 __device__ float4::Native_vec_ __ockl_image_sample_2Da(
unsigned int ADDRESS_SPACE_CONSTANT*i,
unsigned int ADDRESS_SPACE_CONSTANT*s, float4::Native_vec_ c);
97 __device__ float4::Native_vec_ __ockl_image_sample_3D(
unsigned int ADDRESS_SPACE_CONSTANT*i,
unsigned int ADDRESS_SPACE_CONSTANT*s, float4::Native_vec_ c);
99 __device__ float4::Native_vec_ __ockl_image_sample_CM(
unsigned int ADDRESS_SPACE_CONSTANT*i,
unsigned int ADDRESS_SPACE_CONSTANT*s, float4::Native_vec_ c);
101 __device__ float4::Native_vec_ __ockl_image_sample_CMa(
unsigned int ADDRESS_SPACE_CONSTANT*i,
unsigned int ADDRESS_SPACE_CONSTANT*s, float4::Native_vec_ c);
103 __device__ float4::Native_vec_ __ockl_image_sample_grad_1D(
unsigned int ADDRESS_SPACE_CONSTANT*i,
unsigned int ADDRESS_SPACE_CONSTANT*s,
float c,
float dx,
float dy);
105 __device__ float4::Native_vec_ __ockl_image_sample_grad_1Da(
unsigned int ADDRESS_SPACE_CONSTANT*i,
unsigned int ADDRESS_SPACE_CONSTANT*s, float2::Native_vec_ c,
float dx,
float dy);
107 __device__ float4::Native_vec_ __ockl_image_sample_grad_2D(
unsigned int ADDRESS_SPACE_CONSTANT*i,
unsigned int ADDRESS_SPACE_CONSTANT*s, float2::Native_vec_ c, float2::Native_vec_ dx, float2::Native_vec_ dy);
109 __device__ float4::Native_vec_ __ockl_image_sample_grad_2Da(
unsigned int ADDRESS_SPACE_CONSTANT*i,
unsigned int ADDRESS_SPACE_CONSTANT*s, float4::Native_vec_ c, float2::Native_vec_ dx, float2::Native_vec_ dy);
111 __device__ float4::Native_vec_ __ockl_image_sample_grad_3D(
unsigned int ADDRESS_SPACE_CONSTANT*i,
unsigned int ADDRESS_SPACE_CONSTANT*s, float4::Native_vec_ c, float4::Native_vec_ dx, float4::Native_vec_ dy);
113 __device__ float4::Native_vec_ __ockl_image_sample_lod_1D(
unsigned int ADDRESS_SPACE_CONSTANT*i,
unsigned int ADDRESS_SPACE_CONSTANT*s,
float c,
float l);
115 __device__ float4::Native_vec_ __ockl_image_sample_lod_1Da(
unsigned int ADDRESS_SPACE_CONSTANT*i,
unsigned int ADDRESS_SPACE_CONSTANT*s, float2::Native_vec_ c,
float l);
117 __device__ float4::Native_vec_ __ockl_image_sample_lod_2D(
unsigned int ADDRESS_SPACE_CONSTANT*i,
unsigned int ADDRESS_SPACE_CONSTANT*s, float2::Native_vec_ c,
float l);
119 __device__ float4::Native_vec_ __ockl_image_sample_lod_2Da(
unsigned int ADDRESS_SPACE_CONSTANT*i,
unsigned int ADDRESS_SPACE_CONSTANT*s, float4::Native_vec_ c,
float l);
121 __device__ float4::Native_vec_ __ockl_image_sample_lod_3D(
unsigned int ADDRESS_SPACE_CONSTANT*i,
unsigned int ADDRESS_SPACE_CONSTANT*s, float4::Native_vec_ c,
float l);
123 __device__ float4::Native_vec_ __ockl_image_sample_lod_CM(
unsigned int ADDRESS_SPACE_CONSTANT*i,
unsigned int ADDRESS_SPACE_CONSTANT*s, float4::Native_vec_ c,
float l);
125 __device__ float4::Native_vec_ __ockl_image_sample_lod_CMa(
unsigned int ADDRESS_SPACE_CONSTANT*i,
unsigned int ADDRESS_SPACE_CONSTANT*s, float4::Native_vec_ c,
float l);
127 __device__ float4::Native_vec_ __ockl_image_gather4r_2D(
unsigned int ADDRESS_SPACE_CONSTANT*i,
unsigned int ADDRESS_SPACE_CONSTANT*s, float2::Native_vec_ c);
129 __device__ float4::Native_vec_ __ockl_image_gather4g_2D(
unsigned int ADDRESS_SPACE_CONSTANT*i,
unsigned int ADDRESS_SPACE_CONSTANT*s, float2::Native_vec_ c);
131 __device__ float4::Native_vec_ __ockl_image_gather4b_2D(
unsigned int ADDRESS_SPACE_CONSTANT*i,
unsigned int ADDRESS_SPACE_CONSTANT*s, float2::Native_vec_ c);
133 __device__ float4::Native_vec_ __ockl_image_gather4a_2D(
unsigned int ADDRESS_SPACE_CONSTANT*i,
unsigned int ADDRESS_SPACE_CONSTANT*s, float2::Native_vec_ c);