diff options
author | Bill Hollings <bill.hollings@brenwill.com> | 2021-08-16 20:56:05 +0300 |
---|---|---|
committer | Bill Hollings <bill.hollings@brenwill.com> | 2021-08-16 20:56:05 +0300 |
commit | e76fcf93099334ed47277b4427e89fc2b8956241 (patch) | |
tree | d9d4ad2880a341b968f47ecec35c8493da6bfbb1 | |
parent | a75fe07546c2d5d584208d76ae84147b413a3fb0 (diff) |
MSL: Add test for fixes to MSL constant expression type down-casting.
-rw-r--r-- | reference/opt/shaders-msl/comp/type_casting_i64.msl22.comp | 27 | ||||
-rw-r--r-- | reference/shaders-msl/comp/type_casting_i64.msl22.comp | 27 | ||||
-rw-r--r-- | shaders-msl/comp/type_casting_i64.msl22.comp | 23 |
3 files changed, 77 insertions, 0 deletions
diff --git a/reference/opt/shaders-msl/comp/type_casting_i64.msl22.comp b/reference/opt/shaders-msl/comp/type_casting_i64.msl22.comp new file mode 100644 index 00000000..6820b077 --- /dev/null +++ b/reference/opt/shaders-msl/comp/type_casting_i64.msl22.comp @@ -0,0 +1,27 @@ +#include <metal_stdlib> +#include <simd/simd.h> + +using namespace metal; + +struct dst_buff_t +{ + int m0[1]; +}; + +struct src_buff_t +{ + int m0[1]; +}; + +constant int base_val_tmp [[function_constant(0)]]; +constant int base_val = is_function_constant_defined(base_val_tmp) ? base_val_tmp : 0; +constant long shift_val_tmp [[function_constant(1)]]; +constant long shift_val = is_function_constant_defined(shift_val_tmp) ? shift_val_tmp : 0l; +constant int offset = (base_val >> int(shift_val)); +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + +kernel void main0(device dst_buff_t& dst_buff [[buffer(0)]], device src_buff_t& src_buff [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + dst_buff.m0[gl_GlobalInvocationID.x] = src_buff.m0[gl_GlobalInvocationID.x] + offset; +} + diff --git a/reference/shaders-msl/comp/type_casting_i64.msl22.comp b/reference/shaders-msl/comp/type_casting_i64.msl22.comp new file mode 100644 index 00000000..6820b077 --- /dev/null +++ b/reference/shaders-msl/comp/type_casting_i64.msl22.comp @@ -0,0 +1,27 @@ +#include <metal_stdlib> +#include <simd/simd.h> + +using namespace metal; + +struct dst_buff_t +{ + int m0[1]; +}; + +struct src_buff_t +{ + int m0[1]; +}; + +constant int base_val_tmp [[function_constant(0)]]; +constant int base_val = is_function_constant_defined(base_val_tmp) ? base_val_tmp : 0; +constant long shift_val_tmp [[function_constant(1)]]; +constant long shift_val = is_function_constant_defined(shift_val_tmp) ? shift_val_tmp : 0l; +constant int offset = (base_val >> int(shift_val)); +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + +kernel void main0(device dst_buff_t& dst_buff [[buffer(0)]], device src_buff_t& src_buff [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + dst_buff.m0[gl_GlobalInvocationID.x] = src_buff.m0[gl_GlobalInvocationID.x] + offset; +} + diff --git a/shaders-msl/comp/type_casting_i64.msl22.comp b/shaders-msl/comp/type_casting_i64.msl22.comp new file mode 100644 index 00000000..45e682e5 --- /dev/null +++ b/shaders-msl/comp/type_casting_i64.msl22.comp @@ -0,0 +1,23 @@ +#version 450 +#extension GL_ARB_gpu_shader_int64 : require +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +layout(constant_id = 0) const int base_val = 0; +layout(constant_id = 1) const int64_t shift_val = 0; +const int offset = base_val >> shift_val; + +layout(set = 0, binding = 0, std430) buffer src_buff_t +{ + int m0[]; +} src_buff; + +layout(set = 0, binding = 1, std430) buffer dst_buff_t +{ + int m0[]; +} dst_buff; + +void main() +{ + dst_buff.m0[gl_GlobalInvocationID.x] = src_buff.m0[gl_GlobalInvocationID.x] + offset; +} + |