add SPIR-V 1.4 testing for OpCopyLogical

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

@@ -0,0 +1,24 @@
+; SPIR-V
+; Version: 1.4
+               OpCapability Addresses
+               OpCapability Kernel
+               OpMemoryModel Physical32 OpenCL
+               OpEntryPoint Kernel %kernel "copylogical_test"
+       %uint = OpTypeInt 32 0
+      %float = OpTypeFloat 32
+       %void = OpTypeVoid
+   %struct_a = OpTypeStruct %uint %float
+%ptr_struct_a = OpTypePointer CrossWorkgroup %struct_a
+   %struct_b = OpTypeStruct %uint %float
+%ptr_struct_b = OpTypePointer CrossWorkgroup %struct_b
+ %kernel_sig = OpTypeFunction %void %ptr_struct_b
+  %uint_1024 = OpConstant %uint 1024
+   %float_pi = OpConstant %float 3.1415
+%struct_a_src = OpConstantComposite %struct_a %uint_1024 %float_pi
+     %kernel = OpFunction %void None %kernel_sig
+        %dst = OpFunctionParameter %ptr_struct_b
+      %entry = OpLabel
+%struct_b_dst = OpCopyLogical %struct_b %struct_a_src
+               OpStore %dst %struct_b_dst
+               OpReturn
+               OpFunctionEnd

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

@@ -0,0 +1,24 @@
+; SPIR-V
+; Version: 1.4
+               OpCapability Addresses
+               OpCapability Kernel
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %kernel "copylogical_test"
+       %uint = OpTypeInt 32 0
+      %float = OpTypeFloat 32
+       %void = OpTypeVoid
+   %struct_a = OpTypeStruct %uint %float
+%ptr_struct_a = OpTypePointer CrossWorkgroup %struct_a
+   %struct_b = OpTypeStruct %uint %float
+%ptr_struct_b = OpTypePointer CrossWorkgroup %struct_b
+ %kernel_sig = OpTypeFunction %void %ptr_struct_b
+  %uint_1024 = OpConstant %uint 1024
+   %float_pi = OpConstant %float 3.1415
+%struct_a_src = OpConstantComposite %struct_a %uint_1024 %float_pi
+     %kernel = OpFunction %void None %kernel_sig
+        %dst = OpFunctionParameter %ptr_struct_b
+      %entry = OpLabel
+%struct_b_dst = OpCopyLogical %struct_b %struct_a_src
+               OpStore %dst %struct_b_dst
+               OpReturn
+               OpFunctionEnd

test_conformance/spirv_new/test_spirv_14.cpp

@@ -511,3 +511,53 @@ TEST_SPIRV_FUNC(spirv14_select_composite)
 
     return TEST_PASS;
 }
+
+TEST_SPIRV_FUNC(spirv14_copylogical)
+{
+    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/copylogical_struct");
+    SPIRV_CHECK_ERROR(error, "Failed to compile spv program");
+
+    clKernelWrapper kernel = clCreateKernel(prog, "copylogical_test", &error);
+    SPIRV_CHECK_ERROR(error, "Failed to create spv kernel");
+
+    struct TestStruct
+    {
+        cl_int i;
+        cl_float f;
+    };
+    TestStruct results{ 0, 0.0f };
+    clMemWrapper dst = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                      sizeof(results), NULL, &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, NULL, &global, NULL, 0,
+                                   NULL, NULL);
+    SPIRV_CHECK_ERROR(error, "Failed to enqueue kernel");
+
+    error = clEnqueueReadBuffer(queue, dst, CL_TRUE, 0, sizeof(results),
+                                &results, 0, NULL, NULL);
+    SPIRV_CHECK_ERROR(error, "Unable to read destination buffer");
+
+    if (results.i != 1024 || results.f != 3.1415f)
+    {
+        log_error(
+            "Results mismatch!  Got: { %d, %f }\n",
+            results.i, results.f);
+        return TEST_FAIL;
+    }
+
+    return TEST_PASS;
+}