diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/copymemory_memory_operands.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.4/copymemory_memory_operands.spvasm32 new file mode 100644 index 00000000..f4d82f48 --- /dev/null +++ b/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 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/copymemory_memory_operands.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.4/copymemory_memory_operands.spvasm64 new file mode 100644 index 00000000..c615213f --- /dev/null +++ b/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 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/nonwriteable_decoration_function_storage_class.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.4/nonwriteable_decoration_function_storage_class.spvasm32 new file mode 100644 index 00000000..c94befbe --- /dev/null +++ b/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 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/nonwriteable_decoration_function_storage_class.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.4/nonwriteable_decoration_function_storage_class.spvasm64 new file mode 100644 index 00000000..7171438b --- /dev/null +++ b/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 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/select_struct.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.4/select_struct.spvasm32 new file mode 100644 index 00000000..516a4107 --- /dev/null +++ b/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 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/select_struct.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.4/select_struct.spvasm64 new file mode 100644 index 00000000..cbcbbba4 --- /dev/null +++ b/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 diff --git a/test_conformance/spirv_new/test_spirv_14.cpp b/test_conformance/spirv_new/test_spirv_14.cpp index 60d74e57..3df78b45 100644 --- a/test_conformance/spirv_new/test_spirv_14.cpp +++ b/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 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 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; +}