diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/ptrops.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.4/ptrops.spvasm32 new file mode 100644 index 000000000..492cae203 --- /dev/null +++ b/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 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/ptrops.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.4/ptrops.spvasm64 new file mode 100644 index 000000000..f9c5ae5a5 --- /dev/null +++ b/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 diff --git a/test_conformance/spirv_new/test_spirv_14.cpp b/test_conformance/spirv_new/test_spirv_14.cpp index 83a72d44d..6c6b890f4 100644 --- a/test_conformance/spirv_new/test_spirv_14.cpp +++ b/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 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; +}