add SPIR-V 1.4 testing for OpPtrEqual, OpPtrNotEqual, OpPtrDiff (#2054)

This PR adds targeted testing for the SPIR-V 1.4 instructions
OpPtrEqual, OpPtrNotEqual, and OpPtrDiff.

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

@@ -0,0 +1,34 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos LLVM/SPIR-V Translator; 14
+; Bound: 61
+; Schema: 0
+               OpCapability Addresses
+               OpCapability Kernel
+               OpMemoryModel Physical32 OpenCL
+               OpEntryPoint Kernel %kernel "ptrops_test"
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+     %uint_1 = OpConstant %uint 1
+     %uint_2 = OpConstant %uint 2
+       %void = OpTypeVoid
+   %ptr_uint = OpTypePointer CrossWorkgroup %uint
+ %kernel_sig = OpTypeFunction %void %ptr_uint %ptr_uint
+       %bool = OpTypeBool
+     %kernel = OpFunction %void None %kernel_sig
+        %dst = OpFunctionParameter %ptr_uint
+        %tst = OpFunctionParameter %ptr_uint
+      %entry = OpLabel
+        %cmp = OpPtrEqual %bool %dst %tst
+      %bool0 = OpSelect %uint %cmp %uint_1 %uint_0
+       %dst0 = OpInBoundsPtrAccessChain %ptr_uint %dst %uint_0
+               OpStore %dst0 %bool0 Aligned 4
+       %cmp1 = OpPtrNotEqual %bool %dst %tst
+      %bool1 = OpSelect %uint %cmp1 %uint_1 %uint_0
+       %dst1 = OpInBoundsPtrAccessChain %ptr_uint %dst %uint_1
+               OpStore %dst1 %bool1 Aligned 4
+      %delta = OpPtrDiff %uint %dst %tst
+       %dst2 = OpInBoundsPtrAccessChain %ptr_uint %dst %uint_2
+               OpStore %dst2 %delta Aligned 4
+               OpReturn
+               OpFunctionEnd

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

@@ -0,0 +1,39 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos LLVM/SPIR-V Translator; 14
+; Bound: 61
+; Schema: 0
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability Int64
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %kernel "ptrops_test"
+       %uint = OpTypeInt 32 0
+      %ulong = OpTypeInt 64 0
+     %uint_0 = OpConstant %uint 0
+     %uint_1 = OpConstant %uint 1
+    %ulong_0 = OpConstant %ulong 0
+    %ulong_1 = OpConstant %ulong 1
+    %ulong_2 = OpConstant %ulong 2
+       %void = OpTypeVoid
+   %ptr_uint = OpTypePointer CrossWorkgroup %uint
+ %kernel_sig = OpTypeFunction %void %ptr_uint %ptr_uint
+       %bool = OpTypeBool
+     %kernel = OpFunction %void None %kernel_sig
+        %dst = OpFunctionParameter %ptr_uint
+        %tst = OpFunctionParameter %ptr_uint
+      %entry = OpLabel
+        %cmp = OpPtrEqual %bool %dst %tst
+      %bool0 = OpSelect %uint %cmp %uint_1 %uint_0
+       %dst0 = OpInBoundsPtrAccessChain %ptr_uint %dst %ulong_0
+               OpStore %dst0 %bool0 Aligned 4
+       %cmp1 = OpPtrNotEqual %bool %dst %tst
+      %bool1 = OpSelect %uint %cmp1 %uint_1 %uint_0
+       %dst1 = OpInBoundsPtrAccessChain %ptr_uint %dst %ulong_1
+               OpStore %dst1 %bool1 Aligned 4
+      %delta = OpPtrDiff %ulong %dst %tst
+    %deltaui = OpUConvert %uint %delta
+       %dst2 = OpInBoundsPtrAccessChain %ptr_uint %dst %ulong_2
+               OpStore %dst2 %deltaui Aligned 4
+               OpReturn
+               OpFunctionEnd

test_conformance/spirv_new/test_spirv_14.cpp

@@ -209,3 +209,80 @@ TEST_SPIRV_FUNC(spirv14_loop_control_partialcount)
     return test_loop_control_helper(deviceID, context, queue,
                                     "loop_control_partialcount");
 }
+
+TEST_SPIRV_FUNC(spirv14_ptrops)
+{
+    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/ptrops");
+    SPIRV_CHECK_ERROR(error, "Failed to compile spv program");
+
+    clKernelWrapper kernel = clCreateKernel(prog, "ptrops_test", &error);
+    SPIRV_CHECK_ERROR(error, "Failed to create spv kernel");
+
+    std::vector<cl_int> results(3);
+
+    clMemWrapper dst =
+        clCreateBuffer(context, CL_MEM_READ_WRITE,
+                       results.size() * sizeof(cl_int), NULL, &error);
+    SPIRV_CHECK_ERROR(error, "Failed to create dst buffer");
+
+    clMemWrapper tst = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                      sizeof(cl_int), NULL, &error);
+    SPIRV_CHECK_ERROR(error, "Failed to create tst buffer");
+
+    // Test with different pointers:
+    error |= clSetKernelArg(kernel, 0, sizeof(dst), &dst);
+    error |= clSetKernelArg(kernel, 1, sizeof(tst), &tst);
+    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,
+                                results.size() * sizeof(cl_int), results.data(),
+                                0, NULL, NULL);
+    SPIRV_CHECK_ERROR(error, "Unable to read destination buffer");
+
+    if (results[0] != (dst == tst) || results[1] != (dst != tst)
+        || results[2] == 0 /* dst - tst */)
+    {
+        log_error(
+            "Results mismatch with different pointers!  Got: %i, %i, %i\n",
+            results[0], results[1], results[2]);
+        return TEST_FAIL;
+    }
+
+    // Test with equal pointers:
+    error |= clSetKernelArg(kernel, 0, sizeof(dst), &dst);
+    error |= clSetKernelArg(kernel, 1, sizeof(dst), &dst);
+    SPIRV_CHECK_ERROR(error, "Failed to set kernel args");
+
+    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,
+                                results.size() * sizeof(cl_int), results.data(),
+                                0, NULL, NULL);
+    SPIRV_CHECK_ERROR(error, "Unable to read destination buffer");
+
+    if (results[0] != (dst == dst) || results[1] != (dst != dst)
+        || results[2] != 0 /* dst - dst */)
+    {
+        log_error("Results mismatch with equal pointers!  Got: %i, %i, %i\n",
+                  results[0], results[1], results[2]);
+        return TEST_FAIL;
+    }
+
+    return TEST_PASS;
+}