add SPIR-V 1.4 testing for various miscellaneous additions (#2122)

This PR adds targeted testing for a few remaining miscellaneous SPIR-V
1.4 features:

* Variables in the Function storage class can have a NonWriteable
decoration.
* OpCopyMemorySized can have two optional memory operands.
* OpSelect can have composite operands that are not vectors.

test_conformance/spirv_new/spirv_asm/spv1.4/copymemory_memory_operands.spvasm32

@@ -0,0 +1,34 @@
+; SPIR-V
+; Version: 1.4
+               OpCapability Addresses
+               OpCapability Kernel
+               OpMemoryModel Physical32 OpenCL
+               OpEntryPoint Kernel %kernel "copymemory_test"
+       %uint = OpTypeInt 32 0
+       %void = OpTypeVoid
+   %gptr_int = OpTypePointer CrossWorkgroup %uint
+   %pptr_int = OpTypePointer Function %uint
+ %kernel_sig = OpTypeFunction %void %gptr_int
+    %uint_42 = OpConstant %uint 42
+     %uint_1 = OpConstant %uint 1
+     %uint_2 = OpConstant %uint 2
+     %uint_3 = OpConstant %uint 3
+     %uint_4 = OpConstant %uint 4
+     %uint_5 = OpConstant %uint 5
+     %kernel = OpFunction %void None %kernel_sig
+        %dst = OpFunctionParameter %gptr_int
+      %entry = OpLabel
+     %pvalue = OpVariable %pptr_int Function %uint_42
+               OpCopyMemory %dst %pvalue                        ; no memory operands
+       %dst1 = OpInBoundsPtrAccessChain %gptr_int %dst %uint_1
+               OpCopyMemory %dst1 %pvalue Volatile              ; one memory operand
+       %dst2 = OpInBoundsPtrAccessChain %gptr_int %dst %uint_2
+               OpCopyMemory %dst2 %pvalue Volatile Volatile     ; two memory operands
+       %dst3 = OpInBoundsPtrAccessChain %gptr_int %dst %uint_3
+               OpCopyMemorySized %dst3 %pvalue %uint_4                   ; no memory operands
+       %dst4 = OpInBoundsPtrAccessChain %gptr_int %dst %uint_4
+               OpCopyMemorySized %dst4 %pvalue %uint_4 Volatile          ; one memory operand
+       %dst5 = OpInBoundsPtrAccessChain %gptr_int %dst %uint_5
+               OpCopyMemorySized %dst5 %pvalue %uint_4 Volatile Volatile ; two memory operands
+               OpReturn
+               OpFunctionEnd

test_conformance/spirv_new/spirv_asm/spv1.4/copymemory_memory_operands.spvasm64

@@ -0,0 +1,36 @@
+; SPIR-V
+; Version: 1.4
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability Int64
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %kernel "copymemory_test"
+       %uint = OpTypeInt 32 0
+      %ulong = OpTypeInt 64 0
+       %void = OpTypeVoid
+   %gptr_int = OpTypePointer CrossWorkgroup %uint
+   %pptr_int = OpTypePointer Function %uint
+ %kernel_sig = OpTypeFunction %void %gptr_int
+    %uint_42 = OpConstant %uint 42
+    %ulong_1 = OpConstant %ulong 1
+    %ulong_2 = OpConstant %ulong 2
+    %ulong_3 = OpConstant %ulong 3
+    %ulong_4 = OpConstant %ulong 4
+    %ulong_5 = OpConstant %ulong 5
+     %kernel = OpFunction %void None %kernel_sig
+        %dst = OpFunctionParameter %gptr_int
+      %entry = OpLabel
+     %pvalue = OpVariable %pptr_int Function %uint_42
+               OpCopyMemory %dst %pvalue                        ; no memory operands
+       %dst1 = OpInBoundsPtrAccessChain %gptr_int %dst %ulong_1
+               OpCopyMemory %dst1 %pvalue Volatile              ; one memory operand
+       %dst2 = OpInBoundsPtrAccessChain %gptr_int %dst %ulong_2
+               OpCopyMemory %dst2 %pvalue Volatile Volatile     ; two memory operands
+       %dst3 = OpInBoundsPtrAccessChain %gptr_int %dst %ulong_3
+               OpCopyMemorySized %dst3 %pvalue %ulong_4                   ; no memory operands
+       %dst4 = OpInBoundsPtrAccessChain %gptr_int %dst %ulong_4
+               OpCopyMemorySized %dst4 %pvalue %ulong_4 Volatile          ; one memory operand
+       %dst5 = OpInBoundsPtrAccessChain %gptr_int %dst %ulong_5
+               OpCopyMemorySized %dst5 %pvalue %ulong_4 Volatile Volatile ; two memory operands
+               OpReturn
+               OpFunctionEnd

test_conformance/spirv_new/spirv_asm/spv1.4/nonwriteable_decoration_function_storage_class.spvasm32

@@ -0,0 +1,21 @@
+; SPIR-V
+; Version: 1.4
+               OpCapability Addresses
+               OpCapability Kernel
+               OpMemoryModel Physical32 OpenCL
+               OpEntryPoint Kernel %kernel "nonwriteable_test"
+               OpDecorate %pvalue NonWritable
+       %uint = OpTypeInt 32 0
+       %void = OpTypeVoid
+   %gptr_int = OpTypePointer CrossWorkgroup %uint
+   %pptr_int = OpTypePointer Function %uint
+ %kernel_sig = OpTypeFunction %void %gptr_int
+    %uint_42 = OpConstant %uint 42
+     %kernel = OpFunction %void None %kernel_sig
+        %dst = OpFunctionParameter %gptr_int
+      %entry = OpLabel
+     %pvalue = OpVariable %pptr_int Function %uint_42
+      %value = OpLoad %uint %pvalue
+               OpStore %dst %value
+               OpReturn
+               OpFunctionEnd

test_conformance/spirv_new/spirv_asm/spv1.4/nonwriteable_decoration_function_storage_class.spvasm64

@@ -0,0 +1,21 @@
+; SPIR-V
+; Version: 1.4
+               OpCapability Addresses
+               OpCapability Kernel
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %kernel "nonwriteable_test"
+               OpDecorate %pvalue NonWritable
+       %uint = OpTypeInt 32 0
+       %void = OpTypeVoid
+   %gptr_int = OpTypePointer CrossWorkgroup %uint
+   %pptr_int = OpTypePointer Function %uint
+ %kernel_sig = OpTypeFunction %void %gptr_int
+    %uint_42 = OpConstant %uint 42
+     %kernel = OpFunction %void None %kernel_sig
+        %dst = OpFunctionParameter %gptr_int
+      %entry = OpLabel
+     %pvalue = OpVariable %pptr_int Function %uint_42
+      %value = OpLoad %uint %pvalue
+               OpStore %dst %value
+               OpReturn
+               OpFunctionEnd

test_conformance/spirv_new/spirv_asm/spv1.4/select_struct.spvasm32

@@ -0,0 +1,38 @@
+; SPIR-V
+; Version: 1.4
+               OpCapability Addresses
+               OpCapability Kernel
+               OpMemoryModel Physical32 OpenCL
+               OpEntryPoint Kernel %kernel "select_struct_test" %pgid
+               OpDecorate %pgid BuiltIn GlobalInvocationId
+               OpDecorate %pgid Constant
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+      %float = OpTypeFloat 32
+       %void = OpTypeVoid
+       %bool = OpTypeBool
+     %struct = OpTypeStruct %uint %float
+%gptr_struct = OpTypePointer CrossWorkgroup %struct
+%iptr_v3uint = OpTypePointer Input %v3uint
+ %kernel_sig = OpTypeFunction %void %gptr_struct
+     %uint_0 = OpConstant %uint 0
+     %uint_1 = OpConstant %uint 1
+  %uint_1024 = OpConstant %uint 1024
+   %float_pi = OpConstant %float 3.1415
+  %uint_2048 = OpConstant %uint 2048
+    %float_e = OpConstant %float 2.7128
+   %struct_a = OpConstantComposite %struct %uint_1024 %float_pi
+   %struct_b = OpConstantComposite %struct %uint_2048 %float_e
+       %pgid = OpVariable %iptr_v3uint Input
+     %kernel = OpFunction %void None %kernel_sig
+   %dst_base = OpFunctionParameter %gptr_struct
+      %entry = OpLabel
+        %gid = OpLoad %v3uint %pgid Aligned 32
+       %gid0 = OpCompositeExtract %uint %gid 0
+       %test = OpBitwiseAnd %uint %gid0 %uint_1
+       %cond = OpIEqual %bool %test %uint_1
+     %result = OpSelect %struct %cond %struct_a %struct_b
+        %dst = OpInBoundsPtrAccessChain %gptr_struct %dst_base %gid0
+               OpStore %dst %result
+               OpReturn
+               OpFunctionEnd

test_conformance/spirv_new/spirv_asm/spv1.4/select_struct.spvasm64

@@ -0,0 +1,40 @@
+; SPIR-V
+; Version: 1.4
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability Int64
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %kernel "select_struct_test" %pgid
+               OpDecorate %pgid BuiltIn GlobalInvocationId
+               OpDecorate %pgid Constant
+       %uint = OpTypeInt 32 0
+      %ulong = OpTypeInt 64 0
+    %v3ulong = OpTypeVector %ulong 3
+      %float = OpTypeFloat 32
+       %void = OpTypeVoid
+       %bool = OpTypeBool
+     %struct = OpTypeStruct %uint %float
+%gptr_struct = OpTypePointer CrossWorkgroup %struct
+%iptr_v3ulong = OpTypePointer Input %v3ulong
+ %kernel_sig = OpTypeFunction %void %gptr_struct
+    %ulong_0 = OpConstant %ulong 0
+    %ulong_1 = OpConstant %ulong 1
+  %uint_1024 = OpConstant %uint 1024
+   %float_pi = OpConstant %float 3.1415
+  %uint_2048 = OpConstant %uint 2048
+    %float_e = OpConstant %float 2.7128
+   %struct_a = OpConstantComposite %struct %uint_1024 %float_pi
+   %struct_b = OpConstantComposite %struct %uint_2048 %float_e
+       %pgid = OpVariable %iptr_v3ulong Input
+     %kernel = OpFunction %void None %kernel_sig
+   %dst_base = OpFunctionParameter %gptr_struct
+      %entry = OpLabel
+        %gid = OpLoad %v3ulong %pgid Aligned 32
+       %gid0 = OpCompositeExtract %ulong %gid 0
+       %test = OpBitwiseAnd %ulong %gid0 %ulong_1
+       %cond = OpIEqual %bool %test %ulong_1
+     %result = OpSelect %struct %cond %struct_a %struct_b
+        %dst = OpInBoundsPtrAccessChain %gptr_struct %dst_base %gid0
+               OpStore %dst %result
+               OpReturn
+               OpFunctionEnd

test_conformance/spirv_new/test_spirv_14.cpp

@@ -352,3 +352,162 @@ TEST_SPIRV_FUNC(spirv14_usersemantic_memberdecoratestring)
 
     return test_usersemantic_decoration(deviceID, context, queue, true);
 }
+
+TEST_SPIRV_FUNC(spirv14_nonwriteable_decoration)
+{
+    if (!is_spirv_version_supported(deviceID, "SPIR-V_1.4"))
+    {
+        log_info("SPIR-V 1.4 not supported; skipping tests.\n");
+        return TEST_SKIPPED_ITSELF;
+    }
+
+    cl_int error = CL_SUCCESS;
+
+    clProgramWrapper prog;
+    error = get_program_with_il(
+        prog, deviceID, context,
+        "spv1.4/nonwriteable_decoration_function_storage_class");
+    SPIRV_CHECK_ERROR(error, "Failed to compile spv program");
+
+    clKernelWrapper kernel = clCreateKernel(prog, "nonwriteable_test", &error);
+    SPIRV_CHECK_ERROR(error, "Failed to create spv kernel");
+
+    int result = 0;
+    clMemWrapper dst = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                      sizeof(result), nullptr, &error);
+    SPIRV_CHECK_ERROR(error, "Failed to create dst buffer");
+
+    error |= clSetKernelArg(kernel, 0, sizeof(dst), &dst);
+    SPIRV_CHECK_ERROR(error, "Failed to set kernel args");
+
+    size_t global = 1;
+    error = clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, &global, nullptr,
+                                   0, nullptr, nullptr);
+    SPIRV_CHECK_ERROR(error, "Failed to enqueue kernel");
+
+    error = clEnqueueReadBuffer(queue, dst, CL_TRUE, 0, sizeof(result), &result,
+                                0, nullptr, nullptr);
+    SPIRV_CHECK_ERROR(error, "Unable to read destination buffer");
+
+    int expected = 42;
+    if (result != expected)
+    {
+        log_error("Result mismatch!  Got %d, Wanted %d\n", result, expected);
+        return TEST_FAIL;
+    }
+
+    return TEST_PASS;
+}
+
+TEST_SPIRV_FUNC(spirv14_copymemory_memory_operands)
+{
+    if (!is_spirv_version_supported(deviceID, "SPIR-V_1.4"))
+    {
+        log_info("SPIR-V 1.4 not supported; skipping tests.\n");
+        return TEST_SKIPPED_ITSELF;
+    }
+
+    cl_int error = CL_SUCCESS;
+
+    clProgramWrapper prog;
+    error = get_program_with_il(prog, deviceID, context,
+                                "spv1.4/copymemory_memory_operands");
+    SPIRV_CHECK_ERROR(error, "Failed to compile spv program");
+
+    clKernelWrapper kernel = clCreateKernel(prog, "copymemory_test", &error);
+    SPIRV_CHECK_ERROR(error, "Failed to create spv kernel");
+
+    std::vector<int> results(6);
+    clMemWrapper dst =
+        clCreateBuffer(context, CL_MEM_READ_WRITE,
+                       results.size() * sizeof(results[0]), nullptr, &error);
+    SPIRV_CHECK_ERROR(error, "Failed to create dst buffer");
+
+    error |= clSetKernelArg(kernel, 0, sizeof(dst), &dst);
+    SPIRV_CHECK_ERROR(error, "Failed to set kernel args");
+
+    size_t global = 1;
+    error = clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, &global, nullptr,
+                                   0, nullptr, nullptr);
+    SPIRV_CHECK_ERROR(error, "Failed to enqueue kernel");
+
+    error = clEnqueueReadBuffer(queue, dst, CL_TRUE, 0,
+                                results.size() * sizeof(results[0]),
+                                results.data(), 0, nullptr, nullptr);
+    SPIRV_CHECK_ERROR(error, "Unable to read destination buffer");
+
+    const int expected = 42;
+    for (auto result : results)
+    {
+        if (result != expected)
+        {
+            log_error("Result mismatch!  Got %d, Wanted %d\n", result,
+                      expected);
+            return TEST_FAIL;
+        }
+    }
+
+    return TEST_PASS;
+}
+
+TEST_SPIRV_FUNC(spirv14_select_composite)
+{
+    constexpr size_t global_size = 16;
+
+    if (!is_spirv_version_supported(deviceID, "SPIR-V_1.4"))
+    {
+        log_info("SPIR-V 1.4 not supported; skipping tests.\n");
+        return TEST_SKIPPED_ITSELF;
+    }
+
+    cl_int error = CL_SUCCESS;
+
+    clProgramWrapper prog;
+    error =
+        get_program_with_il(prog, deviceID, context, "spv1.4/select_struct");
+    SPIRV_CHECK_ERROR(error, "Failed to compile spv program");
+
+    clKernelWrapper kernel = clCreateKernel(prog, "select_struct_test", &error);
+    SPIRV_CHECK_ERROR(error, "Failed to create spv kernel");
+
+    struct TestStruct
+    {
+        cl_int i;
+        cl_float f;
+    };
+
+    std::vector<TestStruct> results(global_size);
+    clMemWrapper dst =
+        clCreateBuffer(context, CL_MEM_READ_WRITE,
+                       results.size() * sizeof(results[0]), nullptr, &error);
+    SPIRV_CHECK_ERROR(error, "Failed to create dst buffer");
+
+    error |= clSetKernelArg(kernel, 0, sizeof(dst), &dst);
+    SPIRV_CHECK_ERROR(error, "Failed to set kernel args");
+
+    error = clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, &global_size,
+                                   nullptr, 0, nullptr, nullptr);
+    SPIRV_CHECK_ERROR(error, "Failed to enqueue kernel");
+
+    error = clEnqueueReadBuffer(queue, dst, CL_TRUE, 0,
+                                results.size() * sizeof(results[0]),
+                                results.data(), 0, nullptr, nullptr);
+    SPIRV_CHECK_ERROR(error, "Unable to read destination buffer");
+
+    const TestStruct struct_a{ 1024, 3.1415f };
+    const TestStruct struct_b{ 2048, 2.7128f };
+
+    for (size_t i = 0; i < global_size; i++)
+    {
+        const TestStruct& expected = (i & 1) ? struct_a : struct_b;
+        if (results[i].i != expected.i || results[i].f != expected.f)
+        {
+            log_error("Result mismatch at index %zu!  Got {%d, %f}, Wanted "
+                      "{%d, %f}\n",
+                      i, results[i].i, results[i].f, expected.i, expected.f);
+            return TEST_FAIL;
+        }
+    }
+
+    return TEST_PASS;
+}