After enabling physical addressing in CLSPV, I found that integer type with bit size 40 is getting generated using following logic
File: /clspv/src/lib/SPIRVProducerPass.cpp
Type *SPIRVProducerPassImpl::CanonicalType(Type *type) { if (type->isIntegerTy()) { auto bit_width = static_cast<uint32_t>(type->getPrimitiveSizeInBits()); if (bit_width > 1) { // round up bit_width to a multiple of 8 bit_width = ((bit_width + 7) / 8) * 8; } return IntegerType::get(type->getContext(), bit_width); }
and error from Vulkan validation layer are
ERR: vk_renderer.cpp:856 (DebugUtilsMessenger): [ VUID-VkShaderModuleCreateInfo-pCode-08737 ] Validation Error: [ VUID-VkShaderModuleCreateInfo-pCode-08737 ] | MessageID = 0xa5625282 | vkCreateShaderModule(): pCreateInfo->pCode (spirv-val produced an error):
Invalid number of bits (40) used for OpTypeInt.
%u40 = OpTypeInt 40 0
The Vulkan spec states: If pCode is a pointer to SPIR-V code, pCode must adhere to the validation rules described by the Validation Rules within a Module section of the SPIR-V Environment appendix (https://docs.vulkan.org/spec/latest/chapters/shaders.html#VUID-VkShaderModuleCreateInfo-pCode-08737)
ERR: vk_renderer.cpp:856 (DebugUtilsMessenger): [ VUID-VkShaderModuleCreateInfo-pCode-08742 ] Validation Error: [ VUID-VkShaderModuleCreateInfo-pCode-08742 ] | MessageID = 0xfc68be96 | vkCreateShaderModule(): SPIR-V Extension SPV_KHR_non_semantic_info was declared, but one of the following requirements is required (1.3.0 (0x00403000) or VK_KHR_shader_non_semantic_info).
The Vulkan spec states: If pCode is a pointer to SPIR-V code, and pCode declares any of the SPIR-V extensions listed in the SPIR-V Environment appendix, one of the corresponding requirements must be satisfied (https://docs.vulkan.org/spec/latest/chapters/shaders.html#VUID-VkShaderModuleCreateInfo-pCode-08742)
ERR: vk_renderer.cpp:856 (DebugUtilsMessenger): [ VUID-VkPipelineShaderStageCreateInfo-pSpecializationInfo-06849 ] Validation Error: [ VUID-VkPipelineShaderStageCreateInfo-pSpecializationInfo-06849 ] | MessageID = 0x437c19d3 | vkCreateComputePipelines(): pCreateInfos[0].stage After specialization was applied, VkShaderModule 0xc000000000c[] produces a spirv-val error (stage VK_SHADER_STAGE_COMPUTE_BIT):
Invalid number of bits (40) used for OpTypeInt.
%u40 = OpTypeInt 40 0
The Vulkan spec states: If a shader module identifier is not specified, the shader code used by the pipeline must be valid as described by the Khronos SPIR-V Specification after applying the specializations provided in pSpecializationInfo, if any, and then converting all specialization constants into fixed constants (https://docs.vulkan.org/spec/latest/chapters/pipelines.html#VUID-VkPipelineShaderStageCreateInfo-pSpecializationInfo-06849
This logic was added in patch on 18-07-22.
4ada0aa roundup integer type to be a multiple of 8 (#884)
roundup integer type to be a multiple of 8 (#884)
Most Vulkan drivers do not support integer type not being multiple
of 8.
Make sure we always produce integer type with a size multiple of 8.
This will fix a compilation issue with libclc lgamma and lgamma_r
function.
I'm curious to know the logic behind this patch. It looks to me at that time some Vulkan drivers were able to handle integer type of size which is multiple of 8 but it's not the case anymore.
To recreate this issue, I used following kernel
__kernel void minimal(__global int *output)
{
if (get_global_id(1) == 0) output[0] = 7;
else if (get_global_id(1) == 1) output[1] = 8;
else if (get_global_id(1) == 2) output[2] = 9;
}
and following clspv options
-arch=spir64 -physical-storage-buffers -cl-arm-non-uniform-work-group-size --spv-version=1.4 --long-vector --global-offset --enable-printf --cl-kernel-arg-info --int8 --rewrite-packed-structs --no-8bit-storage=pushconstant --no-16bit-storage=pushconstant --std430-ubo-layout --fp16 --fp64=0 -module-constants-in-storage-buffer --rounding-mode-rte=32,16 --enable-feature-macros=__opencl_c_atomic_order_acq_rel,__opencl_c_atomic_order_seq_cst,__opencl_c_atomic_scope_device,__opencl_c_images,__opencl_c_3d_image_writes,__opencl_c_read_write_images,__opencl_c_int64,__opencl_c_integer_dot_product_input_4x8bit,__opencl_c_integer_dot_product_input_4x8bit_packed,__opencl_c_subgroups --use-native-builtins=fma,half_exp2,exp2, -vectorize-loops=false
After enabling physical addressing in CLSPV, I found that integer type with bit size 40 is getting generated using following logic
File: /clspv/src/lib/SPIRVProducerPass.cpp
Type *SPIRVProducerPassImpl::CanonicalType(Type *type) { if (type->isIntegerTy()) { auto bit_width = static_cast<uint32_t>(type->getPrimitiveSizeInBits()); if (bit_width > 1) { // round up bit_width to a multiple of 8 bit_width = ((bit_width + 7) / 8) * 8; } return IntegerType::get(type->getContext(), bit_width); }and error from Vulkan validation layer are
ERR: vk_renderer.cpp:856 (DebugUtilsMessenger): [ VUID-VkShaderModuleCreateInfo-pCode-08737 ] Validation Error: [ VUID-VkShaderModuleCreateInfo-pCode-08737 ] | MessageID = 0xa5625282 | vkCreateShaderModule(): pCreateInfo->pCode (spirv-val produced an error):
Invalid number of bits (40) used for OpTypeInt.
%u40 = OpTypeInt 40 0
The Vulkan spec states: If pCode is a pointer to SPIR-V code, pCode must adhere to the validation rules described by the Validation Rules within a Module section of the SPIR-V Environment appendix (https://docs.vulkan.org/spec/latest/chapters/shaders.html#VUID-VkShaderModuleCreateInfo-pCode-08737)
ERR: vk_renderer.cpp:856 (DebugUtilsMessenger): [ VUID-VkShaderModuleCreateInfo-pCode-08742 ] Validation Error: [ VUID-VkShaderModuleCreateInfo-pCode-08742 ] | MessageID = 0xfc68be96 | vkCreateShaderModule(): SPIR-V Extension SPV_KHR_non_semantic_info was declared, but one of the following requirements is required (1.3.0 (0x00403000) or VK_KHR_shader_non_semantic_info).
The Vulkan spec states: If pCode is a pointer to SPIR-V code, and pCode declares any of the SPIR-V extensions listed in the SPIR-V Environment appendix, one of the corresponding requirements must be satisfied (https://docs.vulkan.org/spec/latest/chapters/shaders.html#VUID-VkShaderModuleCreateInfo-pCode-08742)
ERR: vk_renderer.cpp:856 (DebugUtilsMessenger): [ VUID-VkPipelineShaderStageCreateInfo-pSpecializationInfo-06849 ] Validation Error: [ VUID-VkPipelineShaderStageCreateInfo-pSpecializationInfo-06849 ] | MessageID = 0x437c19d3 | vkCreateComputePipelines(): pCreateInfos[0].stage After specialization was applied, VkShaderModule 0xc000000000c[] produces a spirv-val error (stage VK_SHADER_STAGE_COMPUTE_BIT):
Invalid number of bits (40) used for OpTypeInt.
%u40 = OpTypeInt 40 0
The Vulkan spec states: If a shader module identifier is not specified, the shader code used by the pipeline must be valid as described by the Khronos SPIR-V Specification after applying the specializations provided in pSpecializationInfo, if any, and then converting all specialization constants into fixed constants (https://docs.vulkan.org/spec/latest/chapters/pipelines.html#VUID-VkPipelineShaderStageCreateInfo-pSpecializationInfo-06849
This logic was added in patch on 18-07-22.
4ada0aa roundup integer type to be a multiple of 8 (#884)
roundup integer type to be a multiple of 8 (#884)
Most Vulkan drivers do not support integer type not being multiple
of 8.
Make sure we always produce integer type with a size multiple of 8.
This will fix a compilation issue with libclc lgamma and lgamma_r
function.
I'm curious to know the logic behind this patch. It looks to me at that time some Vulkan drivers were able to handle integer type of size which is multiple of 8 but it's not the case anymore.
To recreate this issue, I used following kernel
__kernel void minimal(__global int *output)
{
if (get_global_id(1) == 0) output[0] = 7;
else if (get_global_id(1) == 1) output[1] = 8;
else if (get_global_id(1) == 2) output[2] = 9;
}
and following clspv options
-arch=spir64 -physical-storage-buffers -cl-arm-non-uniform-work-group-size --spv-version=1.4 --long-vector --global-offset --enable-printf --cl-kernel-arg-info --int8 --rewrite-packed-structs --no-8bit-storage=pushconstant --no-16bit-storage=pushconstant --std430-ubo-layout --fp16 --fp64=0 -module-constants-in-storage-buffer --rounding-mode-rte=32,16 --enable-feature-macros=__opencl_c_atomic_order_acq_rel,__opencl_c_atomic_order_seq_cst,__opencl_c_atomic_scope_device,__opencl_c_images,__opencl_c_3d_image_writes,__opencl_c_read_write_images,__opencl_c_int64,__opencl_c_integer_dot_product_input_4x8bit,__opencl_c_integer_dot_product_input_4x8bit_packed,__opencl_c_subgroups --use-native-builtins=fma,half_exp2,exp2, -vectorize-loops=false