From a73bd10cc75df9ff34a209f2361806ffb8ab008c Mon Sep 17 00:00:00 2001 From: Duncan Brawley Date: Tue, 23 Jan 2024 15:01:29 +0000 Subject: [PATCH] [SYCL][Bindless][Exp] Change coords for 3D images to accept 3 component vecs instead of 4 component vecs This commit changes read/write image functions to only accept coords with three arguments instead of the current four that is enforced. Updates tests and bindless spec to reflect this change. --- libclc/ptx-nvidiacl/libspirv/images/image.cl | 194 +++++++++--------- .../sycl_ext_oneapi_bindless_images.asciidoc | 12 +- .../sycl/ext/oneapi/bindless_images.hpp | 28 +-- .../bindless_images/mipmap/mipmap_read_3D.cpp | 7 +- sycl/test-e2e/bindless_images/read_3D.cpp | 4 +- .../bindless_images/read_norm_types.cpp | 5 +- .../bindless_images/read_write_3D.cpp | 6 +- .../read_write_3D_subregion.cpp | 6 +- .../bindless_images/read_write_unsampled.cpp | 12 +- sycl/test-e2e/bindless_images/sampling_3D.cpp | 2 +- .../vulkan_interop/sampled_images.cpp | 2 +- .../vulkan_interop/unsampled_images.cpp | 18 +- 12 files changed, 145 insertions(+), 151 deletions(-) diff --git a/libclc/ptx-nvidiacl/libspirv/images/image.cl b/libclc/ptx-nvidiacl/libspirv/images/image.cl index f55f0c435cf35..1d642cecdced0 100644 --- a/libclc/ptx-nvidiacl/libspirv/images/image.cl +++ b/libclc/ptx-nvidiacl/libspirv/images/image.cl @@ -1660,179 +1660,179 @@ void __nvvm_sust_3d_v4f16_clamp_s(unsigned long imageHandle, int x, int y, // Int _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(int, 1, i, i32, i, int x, x * sizeof(int)) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(int, 2, i, i32, Dv2_i, int2 coord, coord.x * sizeof(int), coord.y) -_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(int, 3, i, i32, Dv4_i, int4 coord, coord.x * sizeof(int), coord.y, coord.z) +_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(int, 3, i, i32, Dv3_i, int3 coord, coord.x * sizeof(int), coord.y, coord.z) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(int2, 1, Dv2_i, v2i32, i, int x, x * sizeof(int2)) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(int2, 2, Dv2_i, v2i32, S0_, int2 coord, coord.x * sizeof(int2), coord.y) -_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(int2, 3, Dv2_i, v2i32, Dv4_i, int4 coord, coord.x * sizeof(int2), coord.y, coord.z) +_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(int2, 3, Dv2_i, v2i32, Dv3_i, int3 coord, coord.x * sizeof(int2), coord.y, coord.z) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(int4, 1, Dv4_i, v4i32, i, int x, x * sizeof(int4)) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(int4, 2, Dv4_i, v4i32, Dv2_i, int2 coord, coord.x * sizeof(int4), coord.y) -_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(int4, 3, Dv4_i, v4i32, S0_, int4 coord, coord.x * sizeof(int4), coord.y, coord.z) +_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(int4, 3, Dv4_i, v4i32, Dv3_i, int3 coord, coord.x * sizeof(int4), coord.y, coord.z) // Unsigned Int _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(unsigned int, 1, j, j32, i, int x, x * sizeof(unsigned int)) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(unsigned int, 2, j, j32, Dv2_i, int2 coord, coord.x * sizeof(unsigned int), coord.y) -_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(unsigned int, 3, j, j32, Dv4_i, int4 coord, coord.x * sizeof(unsigned int), coord.y, coord.z) +_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(unsigned int, 3, j, j32, Dv3_i, int3 coord, coord.x * sizeof(unsigned int), coord.y, coord.z) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(uint2, 1, Dv2_j, v2j32, i, int x, x * sizeof(uint2)) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(uint2, 2, Dv2_j, v2j32, Dv2_i, int2 coord, coord.x * sizeof(uint2), coord.y) -_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(uint2, 3, Dv2_j, v2j32, Dv4_i, int4 coord, coord.x * sizeof(uint2), coord.y, coord.z) +_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(uint2, 3, Dv2_j, v2j32, Dv3_i, int3 coord, coord.x * sizeof(uint2), coord.y, coord.z) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(uint4, 1, Dv4_j, v4j32, i, int x, x * sizeof(uint4)) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(uint4, 2, Dv4_j, v4j32, Dv2_i, int2 coord, coord.x * sizeof(uint4), coord.y) -_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(uint4, 3, Dv4_j, v4j32, Dv4_i, int4 coord, coord.x * sizeof(uint4), coord.y, coord.z) +_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(uint4, 3, Dv4_j, v4j32, Dv3_i, int3 coord, coord.x * sizeof(uint4), coord.y, coord.z) // Short _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(short, 1, s, i16, i, int x, x * sizeof(short)) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(short, 2, s, i16, Dv2_i, int2 coord, coord.x * sizeof(short), coord.y) -_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(short, 3, s, i16, Dv4_i, int4 coord, coord.x * sizeof(short), coord.y, coord.z) +_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(short, 3, s, i16, Dv3_i, int3 coord, coord.x * sizeof(short), coord.y, coord.z) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(short2, 1, Dv2_s, v2i16, i, int x, x * sizeof(short2)) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(short2, 2, Dv2_s, v2i16, Dv2_i, int2 coord, coord.x * sizeof(short2), coord.y) -_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(short2, 3, Dv2_s, v2i16, Dv4_i, int4 coord, coord.x * sizeof(short2), coord.y, coord.z) +_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(short2, 3, Dv2_s, v2i16, Dv3_i, int3 coord, coord.x * sizeof(short2), coord.y, coord.z) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(short4, 1, Dv4_s, v4i16, i, int x, x * sizeof(short4)) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(short4, 2, Dv4_s, v4i16, Dv2_i, int2 coord, coord.x * sizeof(short4), coord.y) -_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(short4, 3, Dv4_s, v4i16, Dv4_i, int4 coord, coord.x * sizeof(short4), coord.y, coord.z) +_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(short4, 3, Dv4_s, v4i16, Dv3_i, int3 coord, coord.x * sizeof(short4), coord.y, coord.z) // Unsigned Short _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(ushort, 1, t, t16, i, int x, x * sizeof(ushort)) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(ushort, 2, t, t16, Dv2_i, int2 coord, coord.x * sizeof(ushort), coord.y) -_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(ushort, 3, t, t16, Dv4_i, int4 coord, coord.x * sizeof(ushort), coord.y, coord.z) +_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(ushort, 3, t, t16, Dv3_i, int3 coord, coord.x * sizeof(ushort), coord.y, coord.z) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(ushort2, 1, Dv2_t, v2t16, i, int x, x * sizeof(ushort2)) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(ushort2, 2, Dv2_t, v2t16, Dv2_i, int2 coord, coord.x * sizeof(ushort2), coord.y) -_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(ushort2, 3, Dv2_t, v2t16, Dv4_i, int4 coord, coord.x * sizeof(ushort2), coord.y, coord.z) +_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(ushort2, 3, Dv2_t, v2t16, Dv3_i, int3 coord, coord.x * sizeof(ushort2), coord.y, coord.z) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(ushort4, 1, Dv4_t, v4t16, i, int x, x * sizeof(ushort4)) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(ushort4, 2, Dv4_t, v4t16, Dv2_i, int2 coord, coord.x * sizeof(ushort4), coord.y) -_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(ushort4, 3, Dv4_t, v4t16, Dv4_i, int4 coord, coord.x * sizeof(ushort4), coord.y, coord.z) +_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(ushort4, 3, Dv4_t, v4t16, Dv3_i, int3 coord, coord.x * sizeof(ushort4), coord.y, coord.z) // Char _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(char, 1, a, i8, i, int x, x * sizeof(char)) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(char, 2, a, i8, Dv2_i, int2 coord, coord.x * sizeof(char), coord.y) -_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(char, 3, a, i8, Dv4_i, int4 coord, coord.x * sizeof(char), coord.y, coord.z) +_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(char, 3, a, i8, Dv3_i, int3 coord, coord.x * sizeof(char), coord.y, coord.z) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(char2, 1, Dv2_a, v2i8, i, int x, x * sizeof(char2)) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(char2, 2, Dv2_a, v2i8, Dv2_i, int2 coord, coord.x * sizeof(char2), coord.y) -_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(char2, 3, Dv2_a, v2i8, Dv4_i, int4 coord, coord.x * sizeof(char2), coord.y, coord.z) +_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(char2, 3, Dv2_a, v2i8, Dv3_i, int3 coord, coord.x * sizeof(char2), coord.y, coord.z) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(char4, 1, Dv4_a, v4i8, i, int x, x * sizeof(char4)) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(char4, 2, Dv4_a, v4i8, Dv2_i, int2 coord, coord.x * sizeof(char4), coord.y) -_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(char4, 3, Dv4_a, v4i8, Dv4_i, int4 coord, coord.x * sizeof(char4), coord.y, coord.z) +_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(char4, 3, Dv4_a, v4i8, Dv3_i, int3 coord, coord.x * sizeof(char4), coord.y, coord.z) // Unsigned Char _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(uchar, 1, h, h8, i, int x, x * sizeof(uchar)) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(uchar, 2, h, h8, Dv2_i, int2 coord, coord.x * sizeof(uchar), coord.y) -_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(uchar, 3, h, h8, Dv4_i, int4 coord, coord.x * sizeof(uchar), coord.y, coord.z) +_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(uchar, 3, h, h8, Dv3_i, int3 coord, coord.x * sizeof(uchar), coord.y, coord.z) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(uchar2, 1, Dv2_h, v2h8, i, int x, x * sizeof(uchar2)) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(uchar2, 2, Dv2_h, v2h8, Dv2_i, int2 coord, coord.x * sizeof(uchar2), coord.y) -_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(uchar2, 3, Dv2_h, v2h8, Dv4_i, int4 coord, coord.x * sizeof(uchar2), coord.y, coord.z) +_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(uchar2, 3, Dv2_h, v2h8, Dv3_i, int3 coord, coord.x * sizeof(uchar2), coord.y, coord.z) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(uchar4, 1, Dv4_h, v4h8, i, int x, x * sizeof(uchar4)) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(uchar4, 2, Dv4_h, v4h8, Dv2_i, int2 coord, coord.x * sizeof(uchar4), coord.y) -_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(uchar4, 3, Dv4_h, v4h8, Dv4_i, int4 coord, coord.x * sizeof(uchar4), coord.y, coord.z) +_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(uchar4, 3, Dv4_h, v4h8, Dv3_i, int3 coord, coord.x * sizeof(uchar4), coord.y, coord.z) // Float _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(float, 1, f, f32, i, int x, x * sizeof(float)) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(float, 2, f, f32, Dv2_i, int2 coord, coord.x * sizeof(float), coord.y) -_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(float, 3, f, f32, Dv4_i, int4 coord, coord.x * sizeof(float), coord.y, coord.z) +_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(float, 3, f, f32, Dv3_i, int3 coord, coord.x * sizeof(float), coord.y, coord.z) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(float2, 1, Dv2_f, v2f32, i, int x, x * sizeof(float2)) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(float2, 2, Dv2_f, v2f32, Dv2_i, int2 coord, coord.x * sizeof(float2), coord.y) -_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(float2, 3, Dv2_f, v2f32, Dv4_i, int4 coord, coord.x * sizeof(float2), coord.y, coord.z) +_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(float2, 3, Dv2_f, v2f32, Dv3_i, int3 coord, coord.x * sizeof(float2), coord.y, coord.z) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(float4, 1, Dv4_f, v4f32, i, int x, x * sizeof(float4)) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(float4, 2, Dv4_f, v4f32, Dv2_i, int2 coord, coord.x * sizeof(float4), coord.y) -_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(float4, 3, Dv4_f, v4f32, Dv4_i, int4 coord, coord.x * sizeof(float4), coord.y, coord.z) +_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(float4, 3, Dv4_f, v4f32, Dv3_i, int3 coord, coord.x * sizeof(float4), coord.y, coord.z) // Half _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(half, 1, DF16_, f16, i, int x, x * sizeof(half)) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(half, 2, DF16_, f16, Dv2_i, int2 coord, coord.x * sizeof(half), coord.y) -_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(half, 3, DF16_, f16, Dv4_i, int4 coord, coord.x * sizeof(half), coord.y, coord.z) +_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(half, 3, DF16_, f16, Dv3_i, int3 coord, coord.x * sizeof(half), coord.y, coord.z) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(half2, 1, Dv2_DF16_, v2f16, i, int x, x * sizeof(half2)) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(half2, 2, Dv2_DF16_, v2f16, Dv2_i, int2 coord, coord.x * sizeof(half2), coord.y) -_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(half2, 3, Dv2_DF16_, v2f16, Dv4_i, int4 coord, coord.x * sizeof(half2), coord.y, coord.z) +_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(half2, 3, Dv2_DF16_, v2f16, Dv3_i, int3 coord, coord.x * sizeof(half2), coord.y, coord.z) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(half4, 1, Dv4_DF16_, v4f16, i, int x, x * sizeof(half4)) _CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(half4, 2, Dv4_DF16_, v4f16, Dv2_i, int2 coord, coord.x * sizeof(half4), coord.y) -_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(half4, 3, Dv4_DF16_, v4f16, Dv4_i, int4 coord, coord.x * sizeof(half4), coord.y, coord.z) +_CLC_DEFINE_IMAGE_BINDLESS_READ_BUILTIN(half4, 3, Dv4_DF16_, v4f16, Dv3_i, int3 coord, coord.x * sizeof(half4), coord.y, coord.z) // WRITES // Int _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(int, 1, i, i, i32, int x, x * sizeof(int), c) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(int, 2, Dv2_i, i, i32, int2 coord, coord.x * sizeof(int), coord.y, c) -_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(int, 3, Dv4_i, i, i32, int4 coord, coord.x * sizeof(int), coord.y, coord.z, c) +_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(int, 3, Dv3_i, i, i32, int3 coord, coord.x * sizeof(int), coord.y, coord.z, c) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(int2, 1, i, Dv2_i, v2i32, int x, x * sizeof(int2), c.x, c.y) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(int2, 2, Dv2_i, S0_, v2i32, int2 coord, coord.x * sizeof(int2), coord.y, c.x, c.y) -_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(int2, 3, Dv4_i, Dv2_i, v2i32, int4 coord, coord.x * sizeof(int2), coord.y, coord.z, c.x, c.y) +_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(int2, 3, Dv3_i, Dv2_i, v2i32, int3 coord, coord.x * sizeof(int2), coord.y, coord.z, c.x, c.y) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(int4, 1, i, Dv4_i, v4i32, int x, x * sizeof(int4), c.x, c.y, c.z, c.w) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(int4, 2, Dv2_i, Dv4_i, v4i32, int2 coord, coord.x * sizeof(int4), coord.y, c.x, c.y, c.z, c.w) -_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(int4, 3, Dv4_i, S0_, v4i32, int4 coord, coord.x * sizeof(int4), coord.y, coord.z, c.x, c.y, c.z, c.w) +_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(int4, 3, Dv3_i, Dv4_i, v4i32, int3 coord, coord.x * sizeof(int4), coord.y, coord.z, c.x, c.y, c.z, c.w) // Unsigned Int _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(unsigned int, 1, i, j, j32, int x, x * sizeof(unsigned int), c) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(unsigned int, 2, Dv2_i, j, j32, int2 coord, coord.x * sizeof(unsigned int), coord.y, c) -_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(unsigned int, 3, Dv4_i, j, j32, int4 coord, coord.x * sizeof(unsigned int), coord.y, coord.z, c) +_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(unsigned int, 3, Dv3_i, j, j32, int3 coord, coord.x * sizeof(unsigned int), coord.y, coord.z, c) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(uint2, 1, i, Dv2_j, v2j32, int x, x * sizeof(uint2), c.x, c.y) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(uint2, 2, Dv2_i, Dv2_j, v2j32, int2 coord, coord.x * sizeof(uint2), coord.y, c.x, c.y) -_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(uint2, 3, Dv4_i, Dv2_j, v2j32, int4 coord, coord.x * sizeof(uint2), coord.y, coord.z, c.x, c.y) +_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(uint2, 3, Dv3_i, Dv2_j, v2j32, int3 coord, coord.x * sizeof(uint2), coord.y, coord.z, c.x, c.y) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(uint4, 1, i, Dv4_j, v4j32, int x, x * sizeof(uint4), c.x, c.y, c.z, c.w) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(uint4, 2, Dv2_i, Dv4_j, v4j32, int2 coord, coord.x * sizeof(uint4), coord.y, c.x, c.y, c.z, c.w) -_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(uint4, 3, Dv4_i, Dv4_j, v4j32, int4 coord, coord.x * sizeof(uint4), coord.y, coord.z, c.x, c.y, c.z, c.w) +_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(uint4, 3, Dv3_i, Dv4_j, v4j32, int3 coord, coord.x * sizeof(uint4), coord.y, coord.z, c.x, c.y, c.z, c.w) // Short _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(short, 1, i, s, i16, int x, x * sizeof(short), c) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(short, 2, Dv2_i, s, i16, int2 coord, coord.x * sizeof(short), coord.y, c) -_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(short, 3, Dv4_i, s, i16, int4 coord, coord.x * sizeof(short), coord.y, coord.z, c) +_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(short, 3, Dv3_i, s, i16, int3 coord, coord.x * sizeof(short), coord.y, coord.z, c) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(short2, 1, i, Dv2_s, v2i16, int x, x * sizeof(short2), c.x, c.y) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(short2, 2, Dv2_i, Dv2_s, v2i16, int2 coord, coord.x * sizeof(short2), coord.y, c.x, c.y) -_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(short2, 3, Dv4_i, Dv2_s, v2i16, int4 coord, coord.x * sizeof(short2), coord.y, coord.z, c.x, c.y) +_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(short2, 3, Dv3_i, Dv2_s, v2i16, int3 coord, coord.x * sizeof(short2), coord.y, coord.z, c.x, c.y) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(short4, 1, i, Dv4_s, v4i16, int x, x * sizeof(short4), c.x, c.y, c.z, c.w) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(short4, 2, Dv2_i, Dv4_s, v4i16, int2 coord, coord.x * sizeof(short4), coord.y, c.x, c.y, c.z, c.w) -_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(short4, 3, Dv4_i, Dv4_s, v4i16, int4 coord, coord.x * sizeof(short4), coord.y, coord.z, c.x, c.y, c.z, c.w) +_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(short4, 3, Dv3_i, Dv4_s, v4i16, int3 coord, coord.x * sizeof(short4), coord.y, coord.z, c.x, c.y, c.z, c.w) // Unsigned Short _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(ushort, 1, i, t, t16, int x, x * sizeof(ushort), c) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(ushort, 2, Dv2_i, t, t16, int2 coord, coord.x * sizeof(ushort), coord.y, c) -_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(ushort, 3, Dv4_i, t, t16, int4 coord, coord.x * sizeof(ushort), coord.y, coord.z, c) +_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(ushort, 3, Dv3_i, t, t16, int3 coord, coord.x * sizeof(ushort), coord.y, coord.z, c) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(ushort2, 1, i, Dv2_t, v2t16, int x, x * sizeof(ushort2), c.x, c.y) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(ushort2, 2, Dv2_i, Dv2_t, v2t16, int2 coord, coord.x * sizeof(ushort2), coord.y, c.x, c.y) -_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(ushort2, 3, Dv4_i, Dv2_t, v2t16, int4 coord, coord.x * sizeof(ushort2), coord.y, coord.z, c.x, c.y) +_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(ushort2, 3, Dv3_i, Dv2_t, v2t16, int3 coord, coord.x * sizeof(ushort2), coord.y, coord.z, c.x, c.y) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(ushort4, 1, i, Dv4_t, v4t16, int x, x * sizeof(ushort4), c.x, c.y, c.z, c.w) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(ushort4, 2, Dv2_i, Dv4_t, v4t16, int2 coord, coord.x * sizeof(ushort4), coord.y, c.x, c.y, c.z, c.w) -_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(ushort4, 3, Dv4_i, Dv4_t, v4t16, int4 coord, coord.x * sizeof(ushort4), coord.y, coord.z, c.x, c.y, c.z, c.w) +_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(ushort4, 3, Dv3_i, Dv4_t, v4t16, int3 coord, coord.x * sizeof(ushort4), coord.y, coord.z, c.x, c.y, c.z, c.w) // Char _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(char, 1, i, a, i8, int x, x * sizeof(char), c) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(char, 2, Dv2_i, a, i8, int2 coord, coord.x * sizeof(char), coord.y, c) -_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(char, 3, Dv4_i, a, i8, int4 coord, coord.x * sizeof(char), coord.y, coord.z, c) +_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(char, 3, Dv3_i, a, i8, int3 coord, coord.x * sizeof(char), coord.y, coord.z, c) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(char2, 1, i, Dv2_a, v2i8, int x, x * sizeof(char2), c.x, c.y) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(char2, 2, Dv2_i, Dv2_a, v2i8, int2 coord, coord.x * sizeof(char2), coord.y, c.x, c.y) -_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(char2, 3, Dv4_i, Dv2_a, v2i8, int4 coord, coord.x * sizeof(char2), coord.y, coord.z, c.x, c.y) +_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(char2, 3, Dv3_i, Dv2_a, v2i8, int3 coord, coord.x * sizeof(char2), coord.y, coord.z, c.x, c.y) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(char4, 1, i, Dv4_a, v4i8, int x, x * sizeof(char4), c.x, c.y, c.z, c.w) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(char4, 2, Dv2_i, Dv4_a, v4i8, int2 coord, coord.x * sizeof(char4), coord.y, c.x, c.y, c.z, c.w) -_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(char4, 3, Dv4_i, Dv4_a, v4i8, int4 coord, coord.x * sizeof(char4), coord.y, coord.z, c.x, c.y, c.z, c.w) +_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(char4, 3, Dv3_i, Dv4_a, v4i8, int3 coord, coord.x * sizeof(char4), coord.y, coord.z, c.x, c.y, c.z, c.w) // Unsigned Char _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(uchar, 1, i, h, h8, int x, x * sizeof(uchar), c) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(uchar, 2, Dv2_i, h, h8, int2 coord, coord.x * sizeof(uchar), coord.y, c) -_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(uchar, 3, Dv4_i, h, h8, int4 coord, coord.x * sizeof(uchar), coord.y, coord.z, c) +_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(uchar, 3, Dv3_i, h, h8, int3 coord, coord.x * sizeof(uchar), coord.y, coord.z, c) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(uchar2, 1, i, Dv2_h, v2h8, int x, x * sizeof(uchar2), c.x, c.y) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(uchar2, 2, Dv2_i, Dv2_h, v2h8, int2 coord, coord.x * sizeof(uchar2), coord.y, c.x, c.y) -_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(uchar2, 3, Dv4_i, Dv2_h, v2h8, int4 coord, coord.x * sizeof(uchar2), coord.y, coord.z, c.x, c.y) +_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(uchar2, 3, Dv3_i, Dv2_h, v2h8, int3 coord, coord.x * sizeof(uchar2), coord.y, coord.z, c.x, c.y) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(uchar4, 1, i, Dv4_h, v4h8, int x, x * sizeof(uchar4), c.x, c.y, c.z, c.w) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(uchar4, 2, Dv2_i, Dv4_h, v4h8, int2 coord, coord.x * sizeof(uchar4), coord.y, c.x, c.y, c.z, c.w) -_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(uchar4, 3, Dv4_i, Dv4_h, v4h8, int4 coord, coord.x * sizeof(uchar4), coord.y, coord.z, c.x, c.y, c.z, c.w) +_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(uchar4, 3, Dv3_i, Dv4_h, v4h8, int3 coord, coord.x * sizeof(uchar4), coord.y, coord.z, c.x, c.y, c.z, c.w) // Float _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(float, 1, i, f, f32, int x, x * sizeof(float), c) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(float, 2, Dv2_i, f, f32, int2 coord, coord.x * sizeof(float), coord.y, c) -_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(float, 3, Dv4_i, f, f32, int4 coord, coord.x * sizeof(float), coord.y, coord.z, c) +_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(float, 3, Dv3_i, f, f32, int3 coord, coord.x * sizeof(float), coord.y, coord.z, c) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(float2, 1, i, Dv2_f, v2f32, int x, x * sizeof(float2), c.x, c.y) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(float2, 2, Dv2_i, Dv2_f, v2f32, int2 coord, coord.x * sizeof(float2), coord.y, c.x, c.y) -_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(float2, 3, Dv4_i, Dv2_f, v2f32, int4 coord, coord.x * sizeof(float2), coord.y, coord.z, c.x, c.y) +_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(float2, 3, Dv3_i, Dv2_f, v2f32, int3 coord, coord.x * sizeof(float2), coord.y, coord.z, c.x, c.y) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(float4, 1, i, Dv4_f, v4f32, int x, x * sizeof(float4), c.x, c.y, c.z, c.w) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(float4, 2, Dv2_i, Dv4_f, v4f32, int2 coord, coord.x * sizeof(float4), coord.y, c.x, c.y, c.z, c.w) -_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(float4, 3, Dv4_i, Dv4_f, v4f32, int4 coord, coord.x * sizeof(float4), coord.y, coord.z, c.x, c.y, c.z, c.w) +_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(float4, 3, Dv3_i, Dv4_f, v4f32, int3 coord, coord.x * sizeof(float4), coord.y, coord.z, c.x, c.y, c.z, c.w) // Half _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(half, 1, i, DF16_, f16, int x, x * sizeof(half), c) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(half, 2, Dv2_i, DF16_, f16, int2 coord, coord.x * sizeof(half), coord.y, c) -_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(half, 3, Dv4_i, DF16_, f16, int4 coord, coord.x * sizeof(half), coord.y, coord.z, c) +_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(half, 3, Dv3_i, DF16_, f16, int3 coord, coord.x * sizeof(half), coord.y, coord.z, c) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(half2, 1, i, Dv2_DF16_, v2f16, int x, x * sizeof(half2), c.x, c.y) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(half2, 2, Dv2_i, Dv2_DF16_, v2f16, int2 coord, coord.x * sizeof(half2), coord.y, c.x, c.y) -_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(half2, 3, Dv4_i, Dv2_DF16_, v2f16, int4 coord, coord.x * sizeof(half2), coord.y, coord.z, c.x, c.y) +_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(half2, 3, Dv3_i, Dv2_DF16_, v2f16, int3 coord, coord.x * sizeof(half2), coord.y, coord.z, c.x, c.y) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(half4, 1, i, Dv4_DF16_, v4f16, int x, x * sizeof(half4), c.x, c.y, c.z, c.w) _CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(half4, 2, Dv2_i, Dv4_DF16_, v4f16, int2 coord, coord.x * sizeof(half4), coord.y, c.x, c.y, c.z, c.w) -_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(half4, 3, Dv4_i, Dv4_DF16_, v4f16, int4 coord, coord.x * sizeof(half4), coord.y, coord.z, c.x, c.y, c.z, c.w) +_CLC_DEFINE_IMAGE_BINDLESS_WRITE_BUILTIN(half4, 3, Dv3_i, Dv4_DF16_, v4f16, int3 coord, coord.x * sizeof(half4), coord.y, coord.z, c.x, c.y, c.z, c.w) // <--- TEXTURES ---> @@ -2185,90 +2185,90 @@ half __nvvm_tex_3d_f16_f32(unsigned long imageHandle, float x, float y, // Int _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(int, 1, i, i32, f, float x, x) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(int, 2, i, i32, Dv2_f, float2 coord, coord.x, coord.y) -_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(int, 3, i, i32, Dv4_f, float4 coord, coord.x, coord.y, coord.z) +_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(int, 3, i, i32, Dv3_f, float3 coord, coord.x, coord.y, coord.z) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(int2, 1, Dv2_i, v2i32, f, float x, x) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(int2, 2, Dv2_i, v2i32, Dv2_f, float2 coord, coord.x, coord.y) -_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(int2, 3, Dv2_i, v2i32, Dv4_f, float4 coord, coord.x, coord.y, coord.z) +_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(int2, 3, Dv2_i, v2i32, Dv3_f, float3 coord, coord.x, coord.y, coord.z) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(int4, 1, Dv4_i, v4i32, f, float x, x) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(int4, 2, Dv4_i, v4i32, Dv2_f, float2 coord, coord.x, coord.y) -_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(int4, 3, Dv4_i, v4i32, Dv4_f, float4 coord, coord.x, coord.y, coord.z) +_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(int4, 3, Dv4_i, v4i32, Dv3_f, float3 coord, coord.x, coord.y, coord.z) // Unsigned int _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(uint, 1, j, j32, f, float x, x) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(uint, 2, j, j32, Dv2_f, float2 coord, coord.x, coord.y) -_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(uint, 3, j, j32, Dv4_f, float4 coord, coord.x, coord.y, coord.z) +_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(uint, 3, j, j32, Dv3_f, float3 coord, coord.x, coord.y, coord.z) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(uint2, 1, Dv2_j, v2j32, f, float x, x) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(uint2, 2, Dv2_j, v2j32, Dv2_f, float2 coord, coord.x, coord.y) -_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(uint2, 3, Dv2_j, v2j32, Dv4_f, float4 coord, coord.x, coord.y, coord.z) +_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(uint2, 3, Dv2_j, v2j32, Dv3_f, float3 coord, coord.x, coord.y, coord.z) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(uint4, 1, Dv4_j, v4j32, f, float x, x) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(uint4, 2, Dv4_j, v4j32, Dv2_f, float2 coord, coord.x, coord.y) -_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(uint4, 3, Dv4_j, v4j32, Dv4_f, float4 coord, coord.x, coord.y, coord.z) +_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(uint4, 3, Dv4_j, v4j32, Dv3_f, float3 coord, coord.x, coord.y, coord.z) // Short _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(short, 1, s, i16, f, float x, x) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(short, 2, s, i16, Dv2_f, float2 coord, coord.x, coord.y) -_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(short, 3, s, i16, Dv4_f, float4 coord, coord.x, coord.y, coord.z) +_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(short, 3, s, i16, Dv3_f, float3 coord, coord.x, coord.y, coord.z) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(short2, 1, Dv2_s, v2i16, f, float x, x) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(short2, 2, Dv2_s, v2i16, Dv2_f, float2 coord, coord.x, coord.y) -_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(short2, 3, Dv2_s, v2i16, Dv4_f, float4 coord, coord.x, coord.y, coord.z) +_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(short2, 3, Dv2_s, v2i16, Dv3_f, float3 coord, coord.x, coord.y, coord.z) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(short4, 1, Dv4_s, v4i16, f, float x, x) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(short4, 2, Dv4_s, v4i16, Dv2_f, float2 coord, coord.x, coord.y) -_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(short4, 3, Dv4_s, v4i16, Dv4_f, float4 coord, coord.x, coord.y, coord.z) +_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(short4, 3, Dv4_s, v4i16, Dv3_f, float3 coord, coord.x, coord.y, coord.z) // Unsigned short _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(ushort, 1, t, t16, f, float x, x) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(ushort, 2, t, t16, Dv2_f, float2 coord, coord.x, coord.y) -_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(ushort, 3, t, t16, Dv4_f, float4 coord, coord.x, coord.y, coord.z) +_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(ushort, 3, t, t16, Dv3_f, float3 coord, coord.x, coord.y, coord.z) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(ushort2, 1, Dv2_t, v2t16, f, float x, x) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(ushort2, 2, Dv2_t, v2t16, Dv2_f, float2 coord, coord.x, coord.y) -_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(ushort2, 3, Dv2_t, v2t16, Dv4_f, float4 coord, coord.x, coord.y, coord.z) +_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(ushort2, 3, Dv2_t, v2t16, Dv3_f, float3 coord, coord.x, coord.y, coord.z) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(ushort4, 1, Dv4_t, v4t16, f, float x, x) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(ushort4, 2, Dv4_t, v4t16, Dv2_f, float2 coord, coord.x, coord.y) -_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(ushort4, 3, Dv4_t, v4t16, Dv4_f, float4 coord, coord.x, coord.y, coord.z) +_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(ushort4, 3, Dv4_t, v4t16, Dv3_f, float3 coord, coord.x, coord.y, coord.z) // Char _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(char, 1, a, i8, f, float x, x) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(char, 2, a, i8, Dv2_f, float2 coord, coord.x, coord.y) -_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(char, 3, a, i8, Dv4_f, float4 coord, coord.x, coord.y, coord.z) +_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(char, 3, a, i8, Dv3_f, float3 coord, coord.x, coord.y, coord.z) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(char2, 1, Dv2_a, v2i8, f, float x, x) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(char2, 2, Dv2_a, v2i8, Dv2_f, float2 coord, coord.x, coord.y) -_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(char2, 3, Dv2_a, v2i8, Dv4_f, float4 coord, coord.x, coord.y, coord.z) +_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(char2, 3, Dv2_a, v2i8, Dv3_f, float3 coord, coord.x, coord.y, coord.z) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(char4, 1, Dv4_a, v4i8, f, float x, x) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(char4, 2, Dv4_a, v4i8, Dv2_f, float2 coord, coord.x, coord.y) -_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(char4, 3, Dv4_a, v4i8, Dv4_f, float4 coord, coord.x, coord.y, coord.z) +_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(char4, 3, Dv4_a, v4i8, Dv3_f, float3 coord, coord.x, coord.y, coord.z) // Unsigned Char _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(uchar, 1, h, h8, f, float x, x) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(uchar, 2, h, h8, Dv2_f, float2 coord, coord.x, coord.y) -_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(uchar, 3, h, h8, Dv4_f, float4 coord, coord.x, coord.y, coord.z) +_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(uchar, 3, h, h8, Dv3_f, float3 coord, coord.x, coord.y, coord.z) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(uchar2, 1, Dv2_h, v2h8, f, float x, x) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(uchar2, 2, Dv2_h, v2h8, Dv2_f, float2 coord, coord.x, coord.y) -_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(uchar2, 3, Dv2_h, v2h8, Dv4_f, float4 coord, coord.x, coord.y, coord.z) +_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(uchar2, 3, Dv2_h, v2h8, Dv3_f, float3 coord, coord.x, coord.y, coord.z) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(uchar4, 1, Dv4_h, v4h8, f, float x, x) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(uchar4, 2, Dv4_h, v4h8, Dv2_f, float2 coord, coord.x, coord.y) -_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(uchar4, 3, Dv4_h, v4h8, Dv4_f, float4 coord, coord.x, coord.y, coord.z) +_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(uchar4, 3, Dv4_h, v4h8, Dv3_f, float3 coord, coord.x, coord.y, coord.z) // Float _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(float, 1, f, f32, f, float x, x) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(float, 2, f, f32, Dv2_f, float2 coord, coord.x, coord.y) -_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(float, 3, f, f32, Dv4_f, float4 coord, coord.x, coord.y, coord.z) +_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(float, 3, f, f32, Dv3_f, float3 coord, coord.x, coord.y, coord.z) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(float2, 1, Dv2_f, v2f32, f, float x, x) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(float2, 2, Dv2_f, v2f32, S0_, float2 coord, coord.x, coord.y) -_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(float2, 3, Dv2_f, v2f32, Dv4_f, float4 coord, coord.x, coord.y, coord.z) +_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(float2, 3, Dv2_f, v2f32, Dv3_f, float3 coord, coord.x, coord.y, coord.z) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(float4, 1, Dv4_f, v4f32, f, float x, x) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(float4, 2, Dv4_f, v4f32, Dv2_f, float2 coord, coord.x, coord.y) -_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(float4, 3, Dv4_f, v4f32, S0_, float4 coord, coord.x, coord.y, coord.z) +_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(float4, 3, Dv4_f, v4f32, Dv3_f, float3 coord, coord.x, coord.y, coord.z) // Half _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(half, 1, DF16_, f16, f, float x, x) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(half, 2, DF16_, f16, Dv2_f, float2 coord, coord.x, coord.y) -_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(half, 3, DF16_, f16, Dv4_f, float4 coord, coord.x, coord.y, coord.z) +_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(half, 3, DF16_, f16, Dv3_f, float3 coord, coord.x, coord.y, coord.z) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(half2, 1, Dv2_DF16_, v2f16, f, float x, x) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(half2, 2, Dv2_DF16_, v2f16, Dv2_f, float2 coord, coord.x, coord.y) -_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(half2, 3, Dv2_DF16_, v2f16, Dv4_f, float4 coord, coord.x, coord.y, coord.z) +_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(half2, 3, Dv2_DF16_, v2f16, Dv3_f, float3 coord, coord.x, coord.y, coord.z) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(half4, 1, Dv4_DF16_, v4f16, f, float x, x) _CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(half4, 2, Dv4_DF16_, v4f16, Dv2_f, float2 coord, coord.x, coord.y) -_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(half4, 3, Dv4_DF16_, v4f16, Dv4_f, float4 coord, coord.x, coord.y, coord.z) +_CLC_DEFINE_SAMPLEDIMAGE_BINDLESS_READ_BUILTIN(half4, 3, Dv4_DF16_, v4f16, Dv3_f, float3 coord, coord.x, coord.y, coord.z) // <--- MIPMAP ---> @@ -2372,6 +2372,8 @@ _CLC_DEFINE_MIPMAP_BINDLESS_VEC4THUNK_READS_BUILTIN(half, float, 2, v4f16, v4f32 _CLC_DEFINE_MIPMAP_BINDLESS_VEC4THUNK_READS_BUILTIN(half, float, 3, v4f16, v4f32, COORD_INPUT_3D, COORD_PARAMS_3D, GRAD_INPUT_3D, dXx, dXy, dXz, dYx, dYy, dYz) +#undef _CLC_DEFINE_MIPMAP_BINDLESS_VEC4THUNK_READS_BUILTIN + // Macro to generate mipmap vec2 fetches #define _CLC_DEFINE_MIPMAP_BINDLESS_VEC2THUNK_READS_BUILTIN( \ elem_t, fetch_elem_t, dimension, vec_size, fetch_vec_size, coord_input, \ @@ -2498,90 +2500,90 @@ _CLC_DEFINE_MIPMAP_BINDLESS_THUNK_READS_BUILTIN(half, 3, f16, v4f32, COORD_INPUT // Int _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(int, 1, i, i32, f, float coord, COORD_PARAMS_1D, S2_S2_, , dX, dY) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(int, 2, i, i32, Dv2_f, float2 coord, COORD_PARAMS_2D, S3_S3_, 2, dX.x, dX.y, dY.x, dY.y) -_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(int, 3, i, i32, Dv4_f, float4 coord, COORD_PARAMS_3D, S3_S3_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) +_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(int, 3, i, i32, Dv3_f, float3 coord, COORD_PARAMS_3D, S3_S3_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(int2, 1, Dv2_i, v2i32, f, float coord, COORD_PARAMS_1D, S3_S3_, , dX, dY) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(int2, 2, Dv2_i, v2i32, Dv2_f, float2 coord, COORD_PARAMS_2D, S4_S4_, 2, dX.x, dX.y, dY.x, dY.y) -_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(int2, 3, Dv2_i, v2i32, Dv4_f, float4 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) +_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(int2, 3, Dv2_i, v2i32, Dv3_f, float3 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(int4, 1, Dv4_i, v4i32, f, float coord, coord, S3_S3_, , dX, dY) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(int4, 2, Dv4_i, v4i32, Dv2_f, float2 coord, COORD_PARAMS_2D, S4_S4_, 2, dX.x, dX.y, dY.x, dY.y) -_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(int4, 3, Dv4_i, v4i32, Dv4_f, float4 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) +_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(int4, 3, Dv4_i, v4i32, Dv3_f, float3 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) // UInt _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(uint, 1, j, j32, f, float coord, COORD_PARAMS_1D, S2_S2_, , dX, dY) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(uint, 2, j, j32, Dv2_f, float2 coord, COORD_PARAMS_2D, S3_S3_, 2, dX.x, dX.y, dY.x, dY.y) -_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(uint, 3, j, j32, Dv4_f, float4 coord, COORD_PARAMS_3D, S3_S3_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) +_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(uint, 3, j, j32, Dv3_f, float3 coord, COORD_PARAMS_3D, S3_S3_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(uint2, 1, Dv2_j, v2j32, f, float coord, COORD_PARAMS_1D, S3_S3_, , dX, dY) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(uint2, 2, Dv2_j, v2j32, Dv2_f, float2 coord, COORD_PARAMS_2D, S4_S4_, 2, dX.x, dX.y, dY.x, dY.y) -_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(uint2, 3, Dv2_j, v2j32, Dv4_f, float4 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) +_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(uint2, 3, Dv2_j, v2j32, Dv3_f, float3 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(uint4, 1, Dv4_j, v4j32, f, float coord, COORD_PARAMS_1D, S3_S3_, , dX, dY) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(uint4, 2, Dv4_j, v4j32, Dv2_f, float2 coord, COORD_PARAMS_2D, S4_S4_, 2, dX.x, dX.y, dY.x, dY.y) -_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(uint4, 3, Dv4_j, v4j32, Dv4_f, float4 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) +_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(uint4, 3, Dv4_j, v4j32, Dv3_f, float3 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) // Float _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(float, 1, f, f32, f, float coord, COORD_PARAMS_1D, S2_S2_, , dX, dY) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(float, 2, f, f32, Dv2_f, float2 coord, COORD_PARAMS_2D, S3_S3_, 2, dX.x, dX.y, dY.x, dY.y) -_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(float, 3, f, f32, Dv4_f, float4 coord, COORD_PARAMS_3D, S3_S3_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) +_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(float, 3, f, f32, Dv3_f, float3 coord, COORD_PARAMS_3D, S3_S3_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(float2, 1, Dv2_f, v2f32, f, float coord, COORD_PARAMS_1D, S3_S3_, , dX, dY) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(float2, 2, Dv2_f, v2f32, S0_, float2 coord, COORD_PARAMS_2D, S3_S3_, 2, dX.x, dX.y, dY.x, dY.y) -_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(float2, 3, Dv2_f, v2f32, Dv4_f, float4 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) +_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(float2, 3, Dv2_f, v2f32, Dv3_f, float3 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(float4, 1, Dv4_f, v4f32, f, float coord, COORD_PARAMS_1D, S3_S3_, , dX, dY) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(float4, 2, Dv4_f, v4f32, Dv2_f, float2 coord, COORD_PARAMS_2D, S4_S4_, 2, dX.x, dX.y, dY.x, dY.y) -_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(float4, 3, Dv4_f, v4f32, S0_, float4 coord, COORD_PARAMS_3D, S3_S3_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) +_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(float4, 3, Dv4_f, v4f32, Dv3_f, float3 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) // Short _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(short, 1, s, i16, f, float coord, COORD_PARAMS_1D, S2_S2_, , dX, dY) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(short, 2, s, i16, Dv2_f, float2 coord, COORD_PARAMS_2D, S3_S3_, 2, dX.x, dX.y, dY.x, dY.y) -_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(short, 3, s, i16, Dv4_f, float4 coord, COORD_PARAMS_3D, S3_S3_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) +_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(short, 3, s, i16, Dv3_f, float3 coord, COORD_PARAMS_3D, S3_S3_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(short2, 1, Dv2_s, v2i16, f, float coord, COORD_PARAMS_1D, S3_S3_, , dX, dY) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(short2, 2, Dv2_s, v2i16, Dv2_f, float2 coord, COORD_PARAMS_2D, S4_S4_, 2, dX.x, dX.y, dY.x, dY.y) -_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(short2, 3, Dv2_s, v2i16, Dv4_f, float4 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) +_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(short2, 3, Dv2_s, v2i16, Dv3_f, float3 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(short4, 1, Dv4_s, v4i16, f, float coord, COORD_PARAMS_1D, S3_S3_, , dX, dY) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(short4, 2, Dv4_s, v4i16, Dv2_f, float2 coord, COORD_PARAMS_2D, S4_S4_, 2, dX.x, dX.y, dY.x, dY.y) -_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(short4, 3, Dv4_s, v4i16, Dv4_f, float4 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) +_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(short4, 3, Dv4_s, v4i16, Dv3_f, float3 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) // Unsigned Short _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(ushort, 1, t, j16, f, float coord, COORD_PARAMS_1D, S2_S2_, , dX, dY) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(ushort, 2, t, j16, Dv2_f, float2 coord, COORD_PARAMS_2D, S3_S3_, 2, dX.x, dX.y, dY.x, dY.y) -_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(ushort, 3, t, j16, Dv4_f, float4 coord, COORD_PARAMS_3D, S3_S3_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) +_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(ushort, 3, t, j16, Dv3_f, float3 coord, COORD_PARAMS_3D, S3_S3_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(ushort2, 1, Dv2_t, v2j16, f, float coord, COORD_PARAMS_1D, S3_S3_, , dX, dY) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(ushort2, 2, Dv2_t, v2j16, Dv2_f, float2 coord, COORD_PARAMS_2D, S4_S4_, 2, dX.x, dX.y, dY.x, dY.y) -_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(ushort2, 3, Dv2_t, v2j16, Dv4_f, float4 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) +_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(ushort2, 3, Dv2_t, v2j16, Dv3_f, float3 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(ushort4, 1, Dv4_t, v4j16, f, float coord, COORD_PARAMS_1D, S3_S3_, , dX, dY) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(ushort4, 2, Dv4_t, v4j16, Dv2_f, float2 coord, COORD_PARAMS_2D, S4_S4_, 2, dX.x, dX.y, dY.x, dY.y) -_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(ushort4, 3, Dv4_t, v4j16, Dv4_f, float4 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) +_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(ushort4, 3, Dv4_t, v4j16, Dv3_f, float3 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) // Char _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(char, 1, a, i8, f, float coord, COORD_PARAMS_1D, S2_S2_, , dX, dY) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(char, 2, a, i8, Dv2_f, float2 coord, COORD_PARAMS_2D, S3_S3_, 2, dX.x, dX.y, dY.x, dY.y) -_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(char, 3, a, i8, Dv4_f, float4 coord, COORD_PARAMS_3D, S3_S3_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) +_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(char, 3, a, i8, Dv3_f, float3 coord, COORD_PARAMS_3D, S3_S3_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(char2, 1, Dv2_a, v2i8, f, float coord, COORD_PARAMS_1D, S3_S3_, , dX, dY) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(char2, 2, Dv2_a, v2i8, Dv2_f, float2 coord, COORD_PARAMS_2D, S4_S4_, 2, dX.x, dX.y, dY.x, dY.y) -_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(char2, 3, Dv2_a, v2i8, Dv4_f, float4 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) +_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(char2, 3, Dv2_a, v2i8, Dv3_f, float3 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(char4, 1, Dv4_a, v4i8, f, float coord, COORD_PARAMS_1D, S3_S3_, , dX, dY) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(char4, 2, Dv4_a, v4i8, Dv2_f, float2 coord, COORD_PARAMS_2D, S4_S4_, 2, dX.x, dX.y, dY.x, dY.y) -_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(char4, 3, Dv4_a, v4i8, Dv4_f, float4 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) +_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(char4, 3, Dv4_a, v4i8, Dv3_f, float3 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) // Unsigned Char _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(uchar, 1, h, j8, f, float coord, COORD_PARAMS_1D, S2_S2_, , dX, dY) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(uchar, 2, h, j8, Dv2_f, float2 coord, COORD_PARAMS_2D, S3_S3_, 2, dX.x, dX.y, dY.x, dY.y) -_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(uchar, 3, h, j8, Dv4_f, float4 coord, COORD_PARAMS_3D, S3_S3_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) +_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(uchar, 3, h, j8, Dv3_f, float3 coord, COORD_PARAMS_3D, S3_S3_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(uchar2, 1, Dv2_h, v2j8, f, float coord, COORD_PARAMS_1D, S3_S3_, , dX, dY) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(uchar2, 2, Dv2_h, v2j8, Dv2_f, float2 coord, COORD_PARAMS_2D, S4_S4_, 2, dX.x, dX.y, dY.x, dY.y) -_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(uchar2, 3, Dv2_h, v2j8, Dv4_f, float4 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) +_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(uchar2, 3, Dv2_h, v2j8, Dv3_f, float3 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(uchar4, 1, Dv4_h, v4j8, f, float coord, COORD_PARAMS_1D, S3_S3_, , dX, dY) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(uchar4, 2, Dv4_h, v4j8, Dv2_f, float2 coord, COORD_PARAMS_2D, S4_S4_, 2, dX.x, dX.y, dY.x, dY.y) -_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(uchar4, 3, Dv4_h, v4j8, Dv4_f, float4 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) +_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(uchar4, 3, Dv4_h, v4j8, Dv3_f, float3 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) // Half _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(half, 1, DF16_, f16, f, float coord, COORD_PARAMS_1D, S2_S2_, , dX, dY) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(half, 2, DF16_, f16, Dv2_f, float2 coord, COORD_PARAMS_2D, S3_S3_, 2, dX.x, dX.y, dY.x, dY.y) -_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(half, 3, DF16_, f16, Dv4_f, float4 coord, COORD_PARAMS_3D, S3_S3_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) +_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(half, 3, DF16_, f16, Dv3_f, float3 coord, COORD_PARAMS_3D, S3_S3_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(half2, 1, Dv2_DF16_, v2f16, f, float coord, COORD_PARAMS_1D, S3_S3_, , dX, dY) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(half2, 2, Dv2_DF16_, v2f16, Dv2_f, float2 coord, COORD_PARAMS_2D, S4_S4_, 2, dX.x, dX.y, dY.x, dY.y) -_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(half2, 3, Dv2_DF16_, v2f16, Dv4_f, float4 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) +_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(half2, 3, Dv2_DF16_, v2f16, Dv3_f, float3 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(half4, 1, Dv4_DF16_, v4f16, f, float coord, COORD_PARAMS_1D, S3_S3_, , dX, dY) _CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(half4, 2, Dv4_DF16_, v4f16, Dv2_f, float2 coord, COORD_PARAMS_2D, S4_S4_, 2, dX.x, dX.y, dY.x, dY.y) -_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(half4, 3, Dv4_DF16_, v4f16, Dv4_f, float4 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) +_CLC_DEFINE_MIPMAP_BINDLESS_READS_BUILTIN(half4, 3, Dv4_DF16_, v4f16, Dv3_f, float3 coord, COORD_PARAMS_3D, S4_S4_, 4, dX.x, dX.y, dX.z, dY.x, dY.y, dY.z) #undef COORD_PARAMS_1D #undef COORD_PARAMS_2D diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc index e5e0e5951f2d1..28ceead32b48b 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc @@ -1030,15 +1030,11 @@ of the <>. Sampled images cannot be written to using `write_image`. For reading and writing of unsampled images, coordinates are specified by `int`, -`sycl::vec`, and `sycl::vec` for 1D, 2D, and 3D images, +`sycl::vec`, and `sycl::vec` for 1D, 2D, and 3D images, respectively. Sampled image reads take `float`, `sycl::vec`, and -`sycl::vec` coordinate types for 1D, 2D, and 3D images, respectively. - -Note that in the case of 3D reads or writes, coordinates for 3D images take a -vector of size 4, not 3, as the fourth element in the coordinate vector is -ignored. +`sycl::vec` coordinate types for 1D, 2D, and 3D images, respectively. Note also that all images must be used in either read-only or write-only fashion within a single kernel invocation; read/write images are not supported. @@ -1061,7 +1057,7 @@ standard types. * All POD types (`char`, `short`, `int`, `float`, etc.) excluding `double` * `sycl::half` -* Variants of `sycl::vec` where `T` is one of the above, and `N` is `1`, `2`, or `4` +* Variants of `sycl::vec` where `T` is one of the above, and `N` is `1`, `2`, or `3` Any other types are classified as user-defined types. @@ -1080,7 +1076,7 @@ struct my_short2 { ``` When providing the above types as `DataT` parameters to an image read function, -the corresponding `HintT` parameters to use would be `sycl::vec` and +the corresponding `HintT` parameters to use would be `sycl::vec` and `sycl::vec`, respectively. == Mipmapped images diff --git a/sycl/include/sycl/ext/oneapi/bindless_images.hpp b/sycl/include/sycl/ext/oneapi/bindless_images.hpp index 847f53ea2547f..c68e0992bb670 100644 --- a/sycl/include/sycl/ext/oneapi/bindless_images.hpp +++ b/sycl/include/sycl/ext/oneapi/bindless_images.hpp @@ -783,8 +783,8 @@ DataT read_image(const unsampled_image_handle &imageHandle [[maybe_unused]], const CoordT &coords [[maybe_unused]]) { detail::assert_unsampled_coords(); constexpr size_t coordSize = detail::coord_size(); - static_assert(coordSize == 1 || coordSize == 2 || coordSize == 4, - "Expected input coordinate to be have 1, 2, or 4 components " + static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3, + "Expected input coordinate to be have 1, 2, or 3 components " "for 1D, 2D and 3D images, respectively."); #ifdef __SYCL_DEVICE_ONLY__ @@ -829,8 +829,8 @@ DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]], const CoordT &coords [[maybe_unused]]) { detail::assert_sampled_coords(); constexpr size_t coordSize = detail::coord_size(); - static_assert(coordSize == 1 || coordSize == 2 || coordSize == 4, - "Expected input coordinate to be have 1, 2, or 4 components " + static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3, + "Expected input coordinate to be have 1, 2, or 3 components " "for 1D, 2D and 3D images, respectively."); #ifdef __SYCL_DEVICE_ONLY__ @@ -871,8 +871,8 @@ DataT read_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]], const float level [[maybe_unused]]) { detail::assert_sampled_coords(); constexpr size_t coordSize = detail::coord_size(); - static_assert(coordSize == 1 || coordSize == 2 || coordSize == 4, - "Expected input coordinate to be have 1, 2, or 4 components " + static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3, + "Expected input coordinate to be have 1, 2, or 3 components " "for 1D, 2D and 3D images, respectively."); #ifdef __SYCL_DEVICE_ONLY__ @@ -915,8 +915,8 @@ DataT read_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]], const CoordT &dY [[maybe_unused]]) { detail::assert_sampled_coords(); constexpr size_t coordSize = detail::coord_size(); - static_assert(coordSize == 1 || coordSize == 2 || coordSize == 4, - "Expected input coordinates and gradients to have 1, 2, or 4 " + static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3, + "Expected input coordinates and gradients to have 1, 2, or 3 " "components for 1D, 2D, and 3D images, respectively."); #ifdef __SYCL_DEVICE_ONLY__ @@ -961,8 +961,8 @@ DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]], const float level [[maybe_unused]]) { detail::assert_sampled_coords(); constexpr size_t coordSize = detail::coord_size(); - static_assert(coordSize == 1 || coordSize == 2 || coordSize == 4, - "Expected input coordinate to be have 1, 2, or 4 components " + static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3, + "Expected input coordinate to be have 1, 2, or 3 components " "for 1D, 2D and 3D images, respectively."); #ifdef __SYCL_DEVICE_ONLY__ @@ -1008,8 +1008,8 @@ DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]], const CoordT &dY [[maybe_unused]]) { detail::assert_sampled_coords(); constexpr size_t coordSize = detail::coord_size(); - static_assert(coordSize == 1 || coordSize == 2 || coordSize == 4, - "Expected input coordinates and gradients to have 1, 2, or 4 " + static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3, + "Expected input coordinates and gradients to have 1, 2, or 3 " "components for 1D, 2D, and 3D images, respectively."); #ifdef __SYCL_DEVICE_ONLY__ @@ -1045,8 +1045,8 @@ void write_image(const unsampled_image_handle &imageHandle [[maybe_unused]], const DataT &color [[maybe_unused]]) { detail::assert_unsampled_coords(); constexpr size_t coordSize = detail::coord_size(); - static_assert(coordSize == 1 || coordSize == 2 || coordSize == 4, - "Expected input coordinate to be have 1, 2, or 4 components " + static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3, + "Expected input coordinate to be have 1, 2, or 3 components " "for 1D, 2D and 3D images, respectively."); #ifdef __SYCL_DEVICE_ONLY__ diff --git a/sycl/test-e2e/bindless_images/mipmap/mipmap_read_3D.cpp b/sycl/test-e2e/bindless_images/mipmap/mipmap_read_3D.cpp index 10a9798f64667..fd08eaa729987 100644 --- a/sycl/test-e2e/bindless_images/mipmap/mipmap_read_3D.cpp +++ b/sycl/test-e2e/bindless_images/mipmap/mipmap_read_3D.cpp @@ -8,7 +8,7 @@ #include // Uncomment to print additional test information -#define VERBOSE_PRINT +// #define VERBOSE_PRINT template class kernel; @@ -102,9 +102,8 @@ template bool runTest() { // Extension: read mipmap with anisotropic filtering with zero // viewing gradients VecType px1 = sycl::ext::oneapi::experimental::read_mipmap( - mipHandle, sycl::float4(fdim0, fdim1, fdim2, (float)0), - sycl::float4(0.0f, 0.0f, 0.0f, 0.0f), - sycl::float4(0.0f, 0.0f, 0.0f, 0.0f)); + mipHandle, sycl::float3(fdim0, fdim1, fdim2), + sycl::float3(0.0f, 0.0f, 0.0f), sycl::float3(0.0f, 0.0f, 0.0f)); outAcc[sycl::id<3>{dim2, dim1, dim0}] = px1[0]; }); diff --git a/sycl/test-e2e/bindless_images/read_3D.cpp b/sycl/test-e2e/bindless_images/read_3D.cpp index b71f0b5828db1..2b45aadba3a1c 100644 --- a/sycl/test-e2e/bindless_images/read_3D.cpp +++ b/sycl/test-e2e/bindless_images/read_3D.cpp @@ -76,10 +76,10 @@ int main() { // Extension: read image data from handle sycl::float4 px1 = sycl::ext::oneapi::experimental::read_image( - imgHandle1, sycl::int4(dim0, dim1, dim2, 0)); + imgHandle1, sycl::int3(dim0, dim1, dim2)); sycl::float4 px2 = sycl::ext::oneapi::experimental::read_image( - imgHandle2, sycl::int4(dim0, dim1, dim2, 0)); + imgHandle2, sycl::int3(dim0, dim1, dim2)); sum = px1[0] + px2[0]; outAcc[sycl::id<3>{dim2, dim1, dim0}] = sum; diff --git a/sycl/test-e2e/bindless_images/read_norm_types.cpp b/sycl/test-e2e/bindless_images/read_norm_types.cpp index 2cbebb2fef886..9c0ce065114c0 100644 --- a/sycl/test-e2e/bindless_images/read_norm_types.cpp +++ b/sycl/test-e2e/bindless_images/read_norm_types.cpp @@ -71,9 +71,8 @@ bool run_test(sycl::range globalSize, sycl::range localSize) { syclexp::write_image(imgOut, sycl::int2(dim0, dim1), pixel); } else if constexpr (NDims == 3) { OutputType pixel = syclexp::read_image( - imgIn, sycl::float4(dim0, dim1, dim2, 0)); - syclexp::write_image(imgOut, sycl::int4(dim0, dim1, dim2, 0), - pixel); + imgIn, sycl::float3(dim0, dim1, dim2)); + syclexp::write_image(imgOut, sycl::int3(dim0, dim1, dim2), pixel); } }); }); diff --git a/sycl/test-e2e/bindless_images/read_write_3D.cpp b/sycl/test-e2e/bindless_images/read_write_3D.cpp index 38216599fe735..3e362b09e9135 100644 --- a/sycl/test-e2e/bindless_images/read_write_3D.cpp +++ b/sycl/test-e2e/bindless_images/read_write_3D.cpp @@ -76,15 +76,15 @@ int main() { // Extension: read image data from handle sycl::float4 px1 = sycl::ext::oneapi::experimental::read_image( - imgIn1, sycl::int4(dim0, dim1, dim2, 0)); + imgIn1, sycl::int3(dim0, dim1, dim2)); sycl::float4 px2 = sycl::ext::oneapi::experimental::read_image( - imgIn2, sycl::int4(dim0, dim1, dim2, 0)); + imgIn2, sycl::int3(dim0, dim1, dim2)); sum = px1[0] + px2[0]; // Extension: write to image with handle sycl::ext::oneapi::experimental::write_image( - imgOut, sycl::int4(dim0, dim1, dim2, 0), sycl::float4(sum)); + imgOut, sycl::int3(dim0, dim1, dim2), sycl::float4(sum)); }); }); diff --git a/sycl/test-e2e/bindless_images/read_write_3D_subregion.cpp b/sycl/test-e2e/bindless_images/read_write_3D_subregion.cpp index 6044162a8887d..706a3ef64065e 100644 --- a/sycl/test-e2e/bindless_images/read_write_3D_subregion.cpp +++ b/sycl/test-e2e/bindless_images/read_write_3D_subregion.cpp @@ -111,14 +111,14 @@ int main() { float sum = 0; // Extension: read image data from handle float px1 = sycl::ext::oneapi::experimental::read_image( - imgHandle1, sycl::int4(dim0, dim1, dim2, 0)); + imgHandle1, sycl::int3(dim0, dim1, dim2)); float px2 = sycl::ext::oneapi::experimental::read_image( - imgHandle2, sycl::int4(dim0, dim1, dim2, 0)); + imgHandle2, sycl::int3(dim0, dim1, dim2)); sum = px1 + px2; // Extension: write to image with handle sycl::ext::oneapi::experimental::write_image( - imgHandle3, sycl::int4(dim0, dim1, dim2, 0), sum); + imgHandle3, sycl::int3(dim0, dim1, dim2), sum); }); }); diff --git a/sycl/test-e2e/bindless_images/read_write_unsampled.cpp b/sycl/test-e2e/bindless_images/read_write_unsampled.cpp index 855da584d7846..38adcac3294d0 100644 --- a/sycl/test-e2e/bindless_images/read_write_unsampled.cpp +++ b/sycl/test-e2e/bindless_images/read_write_unsampled.cpp @@ -99,24 +99,24 @@ struct util { if constexpr (NChannels >= 1) { VecType px1 = sycl::ext::oneapi::experimental::read_image( - input_0, sycl::int4(dim0, dim1, dim2, 0)); + input_0, sycl::int3(dim0, dim1, dim2)); VecType px2 = sycl::ext::oneapi::experimental::read_image( - input_1, sycl::int4(dim0, dim1, dim2, 0)); + input_1, sycl::int3(dim0, dim1, dim2)); auto sum = VecType(util::add_kernel(px1, px2)); sycl::ext::oneapi::experimental::write_image( - output, sycl::int4(dim0, dim1, dim2, 0), VecType(sum)); + output, sycl::int3(dim0, dim1, dim2), VecType(sum)); } else { DType px1 = sycl::ext::oneapi::experimental::read_image( - input_0, sycl::int4(dim0, dim1, dim2, 0)); + input_0, sycl::int3(dim0, dim1, dim2)); DType px2 = sycl::ext::oneapi::experimental::read_image( - input_1, sycl::int4(dim0, dim1, dim2, 0)); + input_1, sycl::int3(dim0, dim1, dim2)); auto sum = DType(util::add_kernel(px1, px2)); sycl::ext::oneapi::experimental::write_image( - output, sycl::int4(dim0, dim1, dim2, 0), DType(sum)); + output, sycl::int3(dim0, dim1, dim2), DType(sum)); } }); }); diff --git a/sycl/test-e2e/bindless_images/sampling_3D.cpp b/sycl/test-e2e/bindless_images/sampling_3D.cpp index 18193617affe1..192fa1f556927 100644 --- a/sycl/test-e2e/bindless_images/sampling_3D.cpp +++ b/sycl/test-e2e/bindless_images/sampling_3D.cpp @@ -80,7 +80,7 @@ int main() { // Extension: read image data from handle sycl::float4 px1 = sycl::ext::oneapi::experimental::read_image( - imgHandle, sycl::float4(fdim0, fdim1, fdim2, (float)0)); + imgHandle, sycl::float3(fdim0, fdim1, fdim2)); outAcc[sycl::id<3>{dim2, dim1, dim0}] = px1[0]; }); diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/sampled_images.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/sampled_images.cpp index eb775b4a02155..e3d28b202d213 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/sampled_images.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/sampled_images.cpp @@ -151,7 +151,7 @@ bool run_sycl(sycl::range globalSize, sycl::range localSize, VecType pixel; pixel = syclexp::read_image< std::conditional_t>( - handles.imgInput, sycl::float4(fdim0, fdim1, fdim2, 0)); + handles.imgInput, sycl::float3(fdim0, fdim1, fdim2)); pixel *= static_cast(10.1f); outAcc[sycl::id{dim2, dim1, dim0}] = pixel; diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/unsampled_images.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/unsampled_images.cpp index d8d79d2c167f1..c1d16567fad57 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/unsampled_images.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/unsampled_images.cpp @@ -209,25 +209,23 @@ void run_ndim_test(sycl::range global_size, if constexpr (NChannels > 1) { VecType px1 = syclexp::read_image( - handles.input_1, sycl::int4(dim0, dim1, dim2, 0)); + handles.input_1, sycl::int3(dim0, dim1, dim2)); VecType px2 = syclexp::read_image( - handles.input_2, sycl::int4(dim0, dim1, dim2, 0)); + handles.input_2, sycl::int3(dim0, dim1, dim2)); auto sum = VecType(util::add_kernel(px1, px2)); - syclexp::write_image(handles.output, - sycl::int4(dim0, dim1, dim2, 0), - VecType(sum)); + syclexp::write_image( + handles.output, sycl::int3(dim0, dim1, dim2), VecType(sum)); } else { DType px1 = syclexp::read_image( - handles.input_1, sycl::int4(dim0, dim1, dim2, 0)); + handles.input_1, sycl::int3(dim0, dim1, dim2)); DType px2 = syclexp::read_image( - handles.input_2, sycl::int4(dim0, dim1, dim2, 0)); + handles.input_2, sycl::int3(dim0, dim1, dim2)); auto sum = DType(util::add_kernel(px1, px2)); - syclexp::write_image(handles.output, - sycl::int4(dim0, dim1, dim2, 0), - DType(sum)); + syclexp::write_image( + handles.output, sycl::int3(dim0, dim1, dim2), DType(sum)); } } });