Skip to content
This repository has been archived by the owner on May 14, 2024. It is now read-only.

Commit

Permalink
Merge "Add support for accessing cubemaps"
Browse files Browse the repository at this point in the history
  • Loading branch information
@b-sumner
b-sumner authored and Gerrit Code Review committed Jun 22, 2017
2 parents bd357ee + fd43760 commit e807ef7
Show file tree
Hide file tree
Showing 3 changed files with 236 additions and 0 deletions.
5 changes: 5 additions & 0 deletions irif/inc/irif.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -304,6 +304,11 @@ extern __attribute__((const)) ulong __llvm_amdgcn_icmp_i64(ulong, ulong, uint) _
extern __attribute__((const)) ulong __llvm_amdgcn_fcmp_f32(float, float, uint) __asm("llvm.amdgcn.fcmp.f32");
extern __attribute__((const)) ulong __llvm_amdgcn_fcmp_f64(double, double, uint) __asm("llvm.amdgcn.fcmp.f64");

extern __attribute__((const)) float __llvm_amdgcn_cubeid(float, float, float) __asm("llvm.amdgcn.cubeid");
extern __attribute__((const)) float __llvm_amdgcn_cubema(float, float, float) __asm("llvm.amdgcn.cubema");
extern __attribute__((const)) float __llvm_amdgcn_cubesc(float, float, float) __asm("llvm.amdgcn.cubesc");
extern __attribute__((const)) float __llvm_amdgcn_cubetc(float, float, float) __asm("llvm.amdgcn.cubetc");

// Buffer Load/Store

extern __attribute__((pure)) float4 __llvm_amdgcn_buffer_load_format_v4f32(uint4 v, uint i, uint o, bool glc, bool slc) __asm("llvm.amdgcn.buffer.load.format.v4f32");
Expand Down
35 changes: 35 additions & 0 deletions ockl/inc/ockl.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -243,24 +243,32 @@ extern __attribute__((pure)) float4 OCKL_MANGLE_T(image_load,2Da)(TSHARP i, int4
extern __attribute__((pure)) float OCKL_MANGLE_T(image_load,2Dad)(TSHARP i, int4 c);
extern __attribute__((pure)) float OCKL_MANGLE_T(image_load,2Dd)(TSHARP i, int2 c);
extern __attribute__((pure)) float4 OCKL_MANGLE_T(image_load,3D)(TSHARP i, int4 c);
extern __attribute__((pure)) float4 OCKL_MANGLE_T(image_load,CM)(TSHARP i, int2 c, int f);
extern __attribute__((pure)) float4 OCKL_MANGLE_T(image_load,CMa)(TSHARP i, int4 c, int f);
extern __attribute__((pure)) float4 OCKL_MANGLE_T(image_load_mip,1D)(TSHARP i, int c, int l);
extern __attribute__((pure)) float4 OCKL_MANGLE_T(image_load_mip,1Da)(TSHARP i, int2 c, int l);
extern __attribute__((pure)) float4 OCKL_MANGLE_T(image_load_mip,2D)(TSHARP i, int2 c, int l);
extern __attribute__((pure)) float4 OCKL_MANGLE_T(image_load_mip,2Da)(TSHARP i, int4 c, int l);
extern __attribute__((pure)) float OCKL_MANGLE_T(image_load_mip,2Dad)(TSHARP i, int4 c, int l);
extern __attribute__((pure)) float OCKL_MANGLE_T(image_load_mip,2Dd)(TSHARP i, int2 c, int l);
extern __attribute__((pure)) float4 OCKL_MANGLE_T(image_load_mip,3D)(TSHARP i, int4 c, int l);
extern __attribute__((pure)) float4 OCKL_MANGLE_T(image_load_mip,CM)(TSHARP i, int2 c, int f, int l);
extern __attribute__((pure)) float4 OCKL_MANGLE_T(image_load_mip,CMa)(TSHARP i, int4 c, int f, int l);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_loadh,1D)(TSHARP i, int c);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_loadh,1Da)(TSHARP i, int2 c);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_loadh,1Db)(TSHARP i, int c);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_loadh,2D)(TSHARP i, int2 c);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_loadh,2Da)(TSHARP i, int4 c);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_loadh,3D)(TSHARP i, int4 c);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_loadh,CM)(TSHARP i, int2 c, int f);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_loadh,CMa)(TSHARP i, int4 c, int f);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_loadh_mip,1D)(TSHARP i, int c, int l);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_loadh_mip,1Da)(TSHARP i, int2 c, int l);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_loadh_mip,2D)(TSHARP i, int2 c, int l);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_loadh_mip,2Da)(TSHARP i, int4 c, int l);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_loadh_mip,3D)(TSHARP i, int4 c, int l);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_loadh_mip,CM)(TSHARP i, int2 c, int f, int l);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_loadh_mip,CMa)(TSHARP i, int4 c, int f, int l);

extern void OCKL_MANGLE_T(image_store,1D)(TSHARP i, int c, float4 p);
extern void OCKL_MANGLE_T(image_store,1Da)(TSHARP i, int2 c, float4 p);
Expand All @@ -270,24 +278,32 @@ extern void OCKL_MANGLE_T(image_store,2Da)(TSHARP i, int4 c, float4 p);
extern void OCKL_MANGLE_T(image_store,2Dad)(TSHARP i, int4 c, float p);
extern void OCKL_MANGLE_T(image_store,2Dd)(TSHARP i, int2 c, float p);
extern void OCKL_MANGLE_T(image_store,3D)(TSHARP i, int4 c, float4 p);
extern void OCKL_MANGLE_T(image_store,CM)(TSHARP i, int2 c, int f, float4 p);
extern void OCKL_MANGLE_T(image_store,CMa)(TSHARP i, int4 c, int f, float4 p);
extern void OCKL_MANGLE_T(image_store_lod,1D)(TSHARP i, int c, int l, float4 p);
extern void OCKL_MANGLE_T(image_store_lod,1Da)(TSHARP i, int2 c, int l, float4 p);
extern void OCKL_MANGLE_T(image_store_lod,2D)(TSHARP i, int2 c, int l, float4 p);
extern void OCKL_MANGLE_T(image_store_lod,2Da)(TSHARP i, int4 c, int l, float4 p);
extern void OCKL_MANGLE_T(image_store_lod,2Dad)(TSHARP i, int4 c, int l, float p);
extern void OCKL_MANGLE_T(image_store_lod,2Dd)(TSHARP i, int2 c, int l, float p);
extern void OCKL_MANGLE_T(image_store_lod,3D)(TSHARP i, int4 c, int l, float4 p);
extern void OCKL_MANGLE_T(image_store_lod,CM)(TSHARP i, int2 c, int f, int l, float4 p);
extern void OCKL_MANGLE_T(image_store_lod,CMa)(TSHARP i, int4 c, int f, int l, float4 p);
extern void OCKL_MANGLE_T(image_storeh,1D)(TSHARP i, int c, half4 p);
extern void OCKL_MANGLE_T(image_storeh,1Da)(TSHARP i, int2 c, half4 p);
extern void OCKL_MANGLE_T(image_storeh,1Db)(TSHARP i, int c, half4 p);
extern void OCKL_MANGLE_T(image_storeh,2D)(TSHARP i, int2 c, half4 p);
extern void OCKL_MANGLE_T(image_storeh,2Da)(TSHARP i, int4 c, half4 p);
extern void OCKL_MANGLE_T(image_storeh,3D)(TSHARP i, int4 c, half4 p);
extern void OCKL_MANGLE_T(image_storeh,CM)(TSHARP i, int2 c, int f, half4 p);
extern void OCKL_MANGLE_T(image_storeh,CMa)(TSHARP i, int4 c, int f, half4 p);
extern void OCKL_MANGLE_T(image_storeh_lod,1D)(TSHARP i, int c, int l, half4 p);
extern void OCKL_MANGLE_T(image_storeh_lod,1Da)(TSHARP i, int2 c, int l, half4 p);
extern void OCKL_MANGLE_T(image_storeh_lod,2D)(TSHARP i, int2 c, int l, half4 p);
extern void OCKL_MANGLE_T(image_storeh_lod,2Da)(TSHARP i, int4 c, int l, half4 p);
extern void OCKL_MANGLE_T(image_storeh_lod,3D)(TSHARP i, int4 c, int l, half4 p);
extern void OCKL_MANGLE_T(image_storeh_lod,CM)(TSHARP i, int2 c, int f, int l, half4 p);
extern void OCKL_MANGLE_T(image_storeh_lod,CMa)(TSHARP i, int4 c, int f, int l, half4 p);

extern __attribute__((pure)) float4 OCKL_MANGLE_T(image_sample,1D)(TSHARP i, SSHARP s, float c);
extern __attribute__((pure)) float4 OCKL_MANGLE_T(image_sample,1Da)(TSHARP i, SSHARP s, float2 c);
Expand All @@ -296,6 +312,8 @@ extern __attribute__((pure)) float4 OCKL_MANGLE_T(image_sample,2Da)(TSHARP i, SS
extern __attribute__((pure)) float OCKL_MANGLE_T(image_sample,2Dad)(TSHARP i, SSHARP s, float4 c);
extern __attribute__((pure)) float OCKL_MANGLE_T(image_sample,2Dd)(TSHARP i, SSHARP s, float2 c);
extern __attribute__((pure)) float4 OCKL_MANGLE_T(image_sample,3D)(TSHARP i, SSHARP s, float4 c);
extern __attribute__((pure)) float4 OCKL_MANGLE_T(image_sample,CM)(TSHARP i, SSHARP s, float4 c);
extern __attribute__((pure)) float4 OCKL_MANGLE_T(image_sample,CMa)(TSHARP i, SSHARP s, float4 c);
extern __attribute__((pure)) float4 OCKL_MANGLE_T(image_sample_grad,1D)(TSHARP i, SSHARP s, float c, float dx, float dy);
extern __attribute__((pure)) float4 OCKL_MANGLE_T(image_sample_grad,1Da)(TSHARP i, SSHARP s, float2 c, float dx, float dy);
extern __attribute__((pure)) float4 OCKL_MANGLE_T(image_sample_grad,2D)(TSHARP i, SSHARP s, float2 c, float2 dx, float2 dy);
Expand All @@ -310,11 +328,15 @@ extern __attribute__((pure)) float4 OCKL_MANGLE_T(image_sample_lod,2Da)(TSHARP i
extern __attribute__((pure)) float OCKL_MANGLE_T(image_sample_lod,2Dad)(TSHARP i, SSHARP s, float4 c, float l);
extern __attribute__((pure)) float OCKL_MANGLE_T(image_sample_lod,2Dd)(TSHARP i, SSHARP s, float2 c, float l);
extern __attribute__((pure)) float4 OCKL_MANGLE_T(image_sample_lod,3D)(TSHARP i, SSHARP s, float4 c, float l);
extern __attribute__((pure)) float4 OCKL_MANGLE_T(image_sample_lod,CM)(TSHARP i, SSHARP s, float4 c, float l);
extern __attribute__((pure)) float4 OCKL_MANGLE_T(image_sample_lod,CMa)(TSHARP i, SSHARP s, float4 c, float l);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_sampleh,1D)(TSHARP i, SSHARP s, float c);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_sampleh,1Da)(TSHARP i, SSHARP s, float2 c);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_sampleh,2D)(TSHARP i, SSHARP s, float2 c);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_sampleh,2Da)(TSHARP i, SSHARP s, float4 c);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_sampleh,3D)(TSHARP i, SSHARP s, float4 c);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_sampleh,CM)(TSHARP i, SSHARP s, float4 c);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_sampleh,CMa)(TSHARP i, SSHARP s, float4 c);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_sampleh_grad,1D)(TSHARP i, SSHARP s, float c, float dx, float dy);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_sampleh_grad,1Da)(TSHARP i, SSHARP s, float2 c, float dx, float dy);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_sampleh_grad,2D)(TSHARP i, SSHARP s, float2 c, float2 dx, float2 dy);
Expand All @@ -325,6 +347,8 @@ extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_sampleh_lod,1Da)(TSHARP i
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_sampleh_lod,2D)(TSHARP i, SSHARP s, float2 c, float l);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_sampleh_lod,2Da)(TSHARP i, SSHARP s, float4 c, float l);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_sampleh_lod,3D)(TSHARP i, SSHARP s, float4 c, float l);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_sampleh_lod,CM)(TSHARP i, SSHARP s, float4 c, float l);
extern __attribute__((pure)) half4 OCKL_MANGLE_T(image_sampleh_lod,CMa)(TSHARP i, SSHARP s, float4 c, float l);

extern __attribute__((pure)) float4 OCKL_MANGLE_T(image_gather4r,2D)(TSHARP i, SSHARP s, float2 c);
extern __attribute__((pure)) float4 OCKL_MANGLE_T(image_gather4g,2D)(TSHARP i, SSHARP s, float2 c);
Expand All @@ -334,6 +358,7 @@ extern __attribute__((pure)) float4 OCKL_MANGLE_T(image_gather4a,2D)(TSHARP i, S
extern __attribute__((const)) int OCKL_MANGLE_T(image_array_size,1Da)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_array_size,2Da)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_array_size,2Dad)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_array_size,CMa)(TSHARP i);

extern __attribute__((const)) int OCKL_MANGLE_T(image_channel_data_type,1D)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_channel_data_type,1Da)(TSHARP i);
Expand All @@ -343,6 +368,8 @@ extern __attribute__((const)) int OCKL_MANGLE_T(image_channel_data_type,2Da)(TSH
extern __attribute__((const)) int OCKL_MANGLE_T(image_channel_data_type,2Dad)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_channel_data_type,2Dd)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_channel_data_type,3D)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_channel_data_type,CM)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_channel_data_type,CMa)(TSHARP i);

extern __attribute__((const)) int OCKL_MANGLE_T(image_channel_order,1D)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_channel_order,1Da)(TSHARP i);
Expand All @@ -352,6 +379,8 @@ extern __attribute__((const)) int OCKL_MANGLE_T(image_channel_order,2Da)(TSHARP
extern __attribute__((const)) int OCKL_MANGLE_T(image_channel_order,2Dad)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_channel_order,2Dd)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_channel_order,3D)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_channel_order,CM)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_channel_order,CMa)(TSHARP i);

extern __attribute__((const)) int OCKL_MANGLE_T(image_depth,3D)(TSHARP i);

Expand All @@ -360,6 +389,8 @@ extern __attribute__((const)) int OCKL_MANGLE_T(image_height,2Da)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_height,2Dad)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_height,2Dd)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_height,3D)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_height,CM)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_height,CMa)(TSHARP i);

extern __attribute__((const)) int OCKL_MANGLE_T(image_num_mip_levels,1D)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_num_mip_levels,1Da)(TSHARP i);
Expand All @@ -368,6 +399,8 @@ extern __attribute__((const)) int OCKL_MANGLE_T(image_num_mip_levels,2Da)(TSHARP
extern __attribute__((const)) int OCKL_MANGLE_T(image_num_mip_levels,2Dad)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_num_mip_levels,2Dd)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_num_mip_levels,3D)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_num_mip_levels,CM)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_num_mip_levels,CMa)(TSHARP i);

extern __attribute__((const)) int OCKL_MANGLE_T(image_width,1D)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_width,1Da)(TSHARP i);
Expand All @@ -377,6 +410,8 @@ extern __attribute__((const)) int OCKL_MANGLE_T(image_width,2Da)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_width,2Dad)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_width,2Dd)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_width,3D)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_width,CM)(TSHARP i);
extern __attribute__((const)) int OCKL_MANGLE_T(image_width,CMa)(TSHARP i);

extern __attribute__((const)) size_t __ockl_get_global_offset(uint);
extern __attribute__((const)) size_t __ockl_get_global_id(uint);
Expand Down
Loading

0 comments on commit e807ef7

Please sign in to comment.