val->constant = vtn_null_constant(b, val->type);
break;
- case SpvOpConstantSampler:
- vtn_fail("OpConstantSampler requires Kernel Capability");
- break;
-
default:
vtn_fail_with_opcode("Unhandled opcode", opcode);
}
vtn_foreach_decoration(b, val, handle_workgroup_size_decoration_cb, NULL);
}
-SpvMemorySemanticsMask
-vtn_storage_class_to_memory_semantics(SpvStorageClass sc)
-{
- switch (sc) {
- case SpvStorageClassStorageBuffer:
- case SpvStorageClassPhysicalStorageBuffer:
- return SpvMemorySemanticsUniformMemoryMask;
- case SpvStorageClassWorkgroup:
- return SpvMemorySemanticsWorkgroupMemoryMask;
- default:
- return SpvMemorySemanticsMaskNone;
- }
-}
-
static void
vtn_split_barrier_semantics(struct vtn_builder *b,
SpvMemorySemanticsMask semantics,
SpvMemorySemanticsOutputMemoryMask);
const SpvMemorySemanticsMask other_semantics =
- semantics & ~(order_semantics | av_vis_semantics | storage_semantics);
+ semantics & ~(order_semantics | av_vis_semantics | storage_semantics |
+ SpvMemorySemanticsVolatileMask);
if (other_semantics)
vtn_warn("Ignoring unhandled memory semantics: %u\n", other_semantics);
image.lod = NULL;
break;
+ case SpvOpImageQueryFormat:
+ case SpvOpImageQueryOrder:
+ res_val = vtn_untyped_value(b, w[3]);
+ image.image = vtn_get_image(b, w[3]);
+ image.coord = NULL;
+ image.sample = NULL;
+ image.lod = NULL;
+ break;
+
case SpvOpImageRead: {
res_val = vtn_untyped_value(b, w[3]);
image.image = vtn_get_image(b, w[3]);
vtn_fail_with_opcode("Invalid image opcode", opcode);
}
+ if (semantics & SpvMemorySemanticsVolatileMask)
+ access |= ACCESS_VOLATILE;
+
nir_intrinsic_op op;
switch (opcode) {
#define OP(S, N) case SpvOp##S: op = nir_intrinsic_image_deref_##N; break;
OP(AtomicOr, atomic_or)
OP(AtomicXor, atomic_xor)
OP(AtomicFAddEXT, atomic_fadd)
+ OP(ImageQueryFormat, format)
+ OP(ImageQueryOrder, order)
#undef OP
default:
vtn_fail_with_opcode("Invalid image opcode", opcode);
intrin->src[0] = nir_src_for_ssa(&image.image->dest.ssa);
- /* size doesn't take non-lod coordinate parameters */
- if (opcode != SpvOpImageQuerySize && opcode != SpvOpImageQuerySizeLod) {
+ switch (opcode) {
+ case SpvOpImageQuerySize:
+ case SpvOpImageQuerySizeLod:
+ case SpvOpImageQueryFormat:
+ case SpvOpImageQueryOrder:
+ break;
+ default:
/* The image coordinate is always 4 components but we may not have that
* many. Swizzle to compensate.
*/
intrin->src[1] = nir_src_for_ssa(expand_to_vec4(&b->nb, image.coord));
intrin->src[2] = nir_src_for_ssa(image.sample);
+ break;
}
/* The Vulkan spec says:
nir_intrinsic_set_access(intrin, access);
switch (opcode) {
+ case SpvOpImageQueryFormat:
+ case SpvOpImageQueryOrder:
+ /* No additional sources */
+ break;
case SpvOpImageQuerySize:
intrin->src[1] = nir_src_for_ssa(nir_imm_int(&b->nb, 0));
break;
SpvScope scope = SpvScopeInvocation;
SpvMemorySemanticsMask semantics = 0;
+ enum gl_access_qualifier access = 0;
switch (opcode) {
case SpvOpAtomicLoad:
vtn_fail_with_opcode("Invalid SPIR-V atomic", opcode);
}
+ if (semantics & SpvMemorySemanticsVolatileMask)
+ access |= ACCESS_VOLATILE;
+
/* uniform as "atomic counter uniform" */
if (ptr->mode == vtn_variable_mode_atomic_counter) {
nir_deref_instr *deref = vtn_pointer_to_deref(b, ptr);
nir_intrinsic_op op = get_ssbo_nir_atomic_op(b, opcode);
atomic = nir_intrinsic_instr_create(b->nb.shader, op);
- nir_intrinsic_set_access(atomic, ACCESS_COHERENT);
+ nir_intrinsic_set_access(atomic, access | ACCESS_COHERENT);
int src = 0;
switch (opcode) {
atomic->src[0] = nir_src_for_ssa(&deref->dest.ssa);
if (ptr->mode != vtn_variable_mode_workgroup)
- nir_intrinsic_set_access(atomic, ACCESS_COHERENT);
+ access |= ACCESS_COHERENT;
+
+ nir_intrinsic_set_access(atomic, access);
switch (opcode) {
case SpvOpAtomicLoad:
/* Atomic ordering operations will implicitly apply to the atomic operation
* storage class, so include that too.
*/
- semantics |= vtn_storage_class_to_memory_semantics(ptr->ptr_type->storage_class);
+ semantics |= vtn_mode_to_memory_semantics(ptr->mode);
SpvMemorySemanticsMask before_semantics;
SpvMemorySemanticsMask after_semantics;
spv_check_supported(kernel_image, cap);
break;
+ case SpvCapabilityLiteralSampler:
+ spv_check_supported(literal_sampler, cap);
+ break;
+
case SpvCapabilityImageReadWrite:
case SpvCapabilityImageMipmap:
case SpvCapabilityPipes:
case SpvCapabilityDeviceEnqueue:
- case SpvCapabilityLiteralSampler:
case SpvCapabilityGenericPointer:
vtn_warn("Unsupported OpenCL-style SPIR-V capability: %s",
spirv_capability_to_string(cap));
case SpvOpConstantFalse:
case SpvOpConstant:
case SpvOpConstantComposite:
- case SpvOpConstantSampler:
case SpvOpConstantNull:
case SpvOpSpecConstantTrue:
case SpvOpSpecConstantFalse:
case SpvOpUndef:
case SpvOpVariable:
+ case SpvOpConstantSampler:
vtn_handle_variables(b, opcode, w, count);
break;
case SpvOpImageRead:
case SpvOpImageWrite:
case SpvOpImageTexelPointer:
+ case SpvOpImageQueryFormat:
+ case SpvOpImageQueryOrder:
vtn_handle_image(b, opcode, w, count);
break;