diff --git a/test_common/harness/testHarness.cpp b/test_common/harness/testHarness.cpp index 3e5d7c95..34e179d9 100644 --- a/test_common/harness/testHarness.cpp +++ b/test_common/harness/testHarness.cpp @@ -1196,18 +1196,21 @@ Version get_device_spirv_il_version(cl_device_id device) ASSERT_SUCCESS(err, "clGetDeviceInfo"); } - if (strstr(str.data(), "SPIR-V_1.0") != NULL) - return Version(1, 0); - else if (strstr(str.data(), "SPIR-V_1.1") != NULL) - return Version(1, 1); - else if (strstr(str.data(), "SPIR-V_1.2") != NULL) - return Version(1, 2); - else if (strstr(str.data(), "SPIR-V_1.3") != NULL) - return Version(1, 3); + // Because this query returns a space-separated list of IL version strings + // we should check for SPIR-V versions in reverse order, to return the + // highest version supported. + if (strstr(str.data(), "SPIR-V_1.5") != NULL) + return Version(1, 5); else if (strstr(str.data(), "SPIR-V_1.4") != NULL) return Version(1, 4); - else if (strstr(str.data(), "SPIR-V_1.5") != NULL) - return Version(1, 5); + else if (strstr(str.data(), "SPIR-V_1.3") != NULL) + return Version(1, 3); + else if (strstr(str.data(), "SPIR-V_1.2") != NULL) + return Version(1, 2); + else if (strstr(str.data(), "SPIR-V_1.1") != NULL) + return Version(1, 1); + else if (strstr(str.data(), "SPIR-V_1.0") != NULL) + return Version(1, 0); throw std::runtime_error(std::string("Unknown SPIR-V version: ") + str.data()); diff --git a/test_conformance/spirv_new/assemble_spirv.py b/test_conformance/spirv_new/assemble_spirv.py index 99b16adf..d02e5421 100755 --- a/test_conformance/spirv_new/assemble_spirv.py +++ b/test_conformance/spirv_new/assemble_spirv.py @@ -1,7 +1,7 @@ #!/usr/bin/env python3 ##################################################################### -# Copyright (c) 2020 The Khronos Group Inc. All Rights Reserved. +# Copyright (c) 2020-2023 The Khronos Group Inc. All Rights Reserved. # # Licensed under the Apache License, Version 2.0 (the "License"); # you may not use this file except in compliance with the License. @@ -30,6 +30,16 @@ import subprocess import sys from textwrap import wrap +# sub-directories for specific SPIR-V environments +spirv_envs = [ + '', # all files in the root directory are considered SPIR-V 1.0 + 'spv1.1', + 'spv1.2', + 'spv1.3', + 'spv1.4', + 'spv1.5', + 'spv1.6', +] def fatal(message): """Print an error message and exit with a non-zero status, to @@ -39,7 +49,7 @@ def fatal(message): sys.exit(1) -def assemble_spirv(asm_dir, bin_dir, spirv_as, verbose): +def assemble_spirv(asm_dir, bin_dir, spirv_as, spirv_env, verbose): """Assemble SPIR-V source into binaries.""" if not os.path.exists(bin_dir): @@ -57,8 +67,8 @@ def assemble_spirv(asm_dir, bin_dir, spirv_as, verbose): bin_file = asm_file_root + asm_file_ext.replace('asm', '') bin_file_path = os.path.join(bin_dir, bin_file) - command = '"{}" --target-env spv1.0 "{}" -o "{}"'.format( - spirv_as, asm_file_path, bin_file_path) + command = '"{}" --target-env "{}" "{}" -o "{}"'.format( + spirv_as, spirv_env, asm_file_path, bin_file_path) if subprocess.call(command, shell=True) != 0: assembly_failures = True print('ERROR: Failure assembling {}: ' @@ -72,7 +82,7 @@ def assemble_spirv(asm_dir, bin_dir, spirv_as, verbose): 'messages from the assembler, if any.'))) -def validate_spirv(bin_dir, spirv_val, verbose): +def validate_spirv(bin_dir, spirv_val, spirv_env, verbose): """Validates SPIR-V binaries. Ignores known failures.""" validation_failures = False @@ -83,8 +93,8 @@ def validate_spirv(bin_dir, spirv_val, verbose): if verbose: print(' Validating {}'.format(bin_file)) - command = '"{}" "{}"'.format( - spirv_val, bin_file_path) + command = '"{}" --target-env "{}" "{}"'.format( + spirv_val, spirv_env, bin_file_path) if subprocess.call(command, shell=True) != 0: print('ERROR: Failure validating {}: ' 'see above output.'.format( @@ -95,8 +105,6 @@ def validate_spirv(bin_dir, spirv_val, verbose): if validation_failures: fatal('ERROR: Validation failure(s) found. ' 'See above for validation output.') - else: - print('All SPIR-V binaries validated successfully.') def parse_args(): @@ -144,18 +152,26 @@ def main(): args = parse_args() - print('Assembling SPIR-V source into binaries...') - assemble_spirv(args.source_dir, args.output_dir, args.assembler, - args.verbose) - print('Finished assembling SPIR-V binaries.') - print() + for subdir in spirv_envs: + src_dir = os.path.join(args.source_dir, subdir) + out_dir = os.path.join(args.output_dir, subdir) + spirv_env = 'spv1.0' if subdir == '' else subdir + print('Assembling SPIR-V source into binaries for target {}...'. + format(spirv_env)) + assemble_spirv(src_dir, out_dir, args.assembler, + spirv_env, args.verbose) + print('Finished assembling SPIR-V binaries.') + print() - if args.skip_validation: - print('Skipping validation of SPIR-V binaries as requested.') - else: - print('Validating SPIR-V binaries...') - validate_spirv(args.output_dir, args.validator, args.verbose) - print() + if args.skip_validation: + print('Skipping validation of SPIR-V binaries as requested.') + else: + print('Validating SPIR-V binaries for target {}...'. + format(spirv_env)) + validate_spirv(out_dir, args.validator, + spirv_env, args.verbose) + print('All SPIR-V binaries validated successfully.') + print() print('Done.') diff --git a/test_conformance/spirv_new/main.cpp b/test_conformance/spirv_new/main.cpp index 3c93baca..fc3c0bec 100644 --- a/test_conformance/spirv_new/main.cpp +++ b/test_conformance/spirv_new/main.cpp @@ -33,9 +33,12 @@ const std::string slash = "/"; #endif const std::string spvExt = ".spv"; +bool gVersionSkip = false; std::string gAddrWidth = ""; std::string spvBinariesPath = "spirv_bin"; -std::string spvBinariesPathArg = "--spirv-binaries-path"; + +const std::string spvBinariesPathArg = "--spirv-binaries-path"; +const std::string spvVersionSkipArg = "--skip-spirv-version-check"; std::vector readBinary(const char *file_name) { @@ -227,7 +230,10 @@ test_status InitCL(cl_device_id id) void printUsage() { log_info("Reading SPIR-V files from default '%s' path.\n", spvBinariesPath.c_str()); - log_info("In case you want to set other directory use '%s' argument.\n", spvBinariesPathArg.c_str()); + log_info("In case you want to set other directory use '%s' argument.\n", + spvBinariesPathArg.c_str()); + log_info("To skip the SPIR-V version check use the '%s' argument.\n", + spvVersionSkipArg.c_str()); } int main(int argc, const char *argv[]) @@ -246,6 +252,11 @@ int main(int argc, const char *argv[]) modifiedSpvBinariesPath = true; } } + if (argv[i] == spvVersionSkipArg) + { + gVersionSkip = true; + argsRemoveNum++; + } if (argsRemoveNum > 0) { for (int j = i; j < (argc - argsRemoveNum); ++j) diff --git a/test_conformance/spirv_new/spirv_asm/basic.spvasm32 b/test_conformance/spirv_new/spirv_asm/basic.spvasm32 new file mode 100644 index 00000000..a640c8c9 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/basic.spvasm32 @@ -0,0 +1,33 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 18 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %9 "test_basic" %gl_GlobalInvocationID + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %gl_GlobalInvocationID Constant + OpDecorate %gl_GlobalInvocationID LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %8 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint +%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input + %9 = OpFunction %void None %8 + %10 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %11 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %12 = OpLabel + %13 = OpLoad %v3uint %gl_GlobalInvocationID Aligned 16 + %14 = OpCompositeExtract %uint %13 0 + %15 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %11 %14 + %16 = OpLoad %uint %15 Aligned 4 + %17 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %10 %14 + OpStore %17 %16 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/basic.spvasm64 b/test_conformance/spirv_new/spirv_asm/basic.spvasm64 new file mode 100644 index 00000000..662bffbf --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/basic.spvasm64 @@ -0,0 +1,38 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 22 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %10 "test_basic" %gl_GlobalInvocationID + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %gl_GlobalInvocationID Constant + OpDecorate %gl_GlobalInvocationID LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + %ulong = OpTypeInt 64 0 + %uint = OpTypeInt 32 0 + %v3ulong = OpTypeVector %ulong 3 +%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %9 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint +%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3ulong Input + %10 = OpFunction %void None %9 + %11 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %12 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %13 = OpLabel + %14 = OpLoad %v3ulong %gl_GlobalInvocationID Aligned 32 + %15 = OpCompositeExtract %ulong %14 0 + %16 = OpUConvert %uint %15 + %17 = OpSConvert %ulong %16 + %18 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %12 %17 + %19 = OpLoad %uint %18 Aligned 4 + %20 = OpSConvert %ulong %16 + %21 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %11 %20 + OpStore %21 %19 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.1/basic.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.1/basic.spvasm32 new file mode 100644 index 00000000..2388c840 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.1/basic.spvasm32 @@ -0,0 +1,33 @@ +; SPIR-V +; Version: 1.1 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 18 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %9 "test_basic" %gl_GlobalInvocationID + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %gl_GlobalInvocationID Constant + OpDecorate %gl_GlobalInvocationID LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %8 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint +%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input + %9 = OpFunction %void None %8 + %10 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %11 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %12 = OpLabel + %13 = OpLoad %v3uint %gl_GlobalInvocationID Aligned 16 + %14 = OpCompositeExtract %uint %13 0 + %15 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %11 %14 + %16 = OpLoad %uint %15 Aligned 4 + %17 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %10 %14 + OpStore %17 %16 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.1/basic.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.1/basic.spvasm64 new file mode 100644 index 00000000..80bc770a --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.1/basic.spvasm64 @@ -0,0 +1,38 @@ +; SPIR-V +; Version: 1.1 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 22 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %10 "test_basic" %gl_GlobalInvocationID + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %gl_GlobalInvocationID Constant + OpDecorate %gl_GlobalInvocationID LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + %ulong = OpTypeInt 64 0 + %uint = OpTypeInt 32 0 + %v3ulong = OpTypeVector %ulong 3 +%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %9 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint +%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3ulong Input + %10 = OpFunction %void None %9 + %11 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %12 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %13 = OpLabel + %14 = OpLoad %v3ulong %gl_GlobalInvocationID Aligned 32 + %15 = OpCompositeExtract %ulong %14 0 + %16 = OpUConvert %uint %15 + %17 = OpSConvert %ulong %16 + %18 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %12 %17 + %19 = OpLoad %uint %18 Aligned 4 + %20 = OpSConvert %ulong %16 + %21 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %11 %20 + OpStore %21 %19 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.2/basic.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.2/basic.spvasm32 new file mode 100644 index 00000000..f3233224 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.2/basic.spvasm32 @@ -0,0 +1,33 @@ +; SPIR-V +; Version: 1.2 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 18 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %9 "test_basic" %gl_GlobalInvocationID + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %gl_GlobalInvocationID Constant + OpDecorate %gl_GlobalInvocationID LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %8 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint +%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input + %9 = OpFunction %void None %8 + %10 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %11 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %12 = OpLabel + %13 = OpLoad %v3uint %gl_GlobalInvocationID Aligned 16 + %14 = OpCompositeExtract %uint %13 0 + %15 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %11 %14 + %16 = OpLoad %uint %15 Aligned 4 + %17 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %10 %14 + OpStore %17 %16 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.2/basic.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.2/basic.spvasm64 new file mode 100644 index 00000000..fcf8b44e --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.2/basic.spvasm64 @@ -0,0 +1,38 @@ +; SPIR-V +; Version: 1.2 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 22 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %10 "test_basic" %gl_GlobalInvocationID + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %gl_GlobalInvocationID Constant + OpDecorate %gl_GlobalInvocationID LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + %ulong = OpTypeInt 64 0 + %uint = OpTypeInt 32 0 + %v3ulong = OpTypeVector %ulong 3 +%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %9 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint +%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3ulong Input + %10 = OpFunction %void None %9 + %11 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %12 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %13 = OpLabel + %14 = OpLoad %v3ulong %gl_GlobalInvocationID Aligned 32 + %15 = OpCompositeExtract %ulong %14 0 + %16 = OpUConvert %uint %15 + %17 = OpSConvert %ulong %16 + %18 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %12 %17 + %19 = OpLoad %uint %18 Aligned 4 + %20 = OpSConvert %ulong %16 + %21 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %11 %20 + OpStore %21 %19 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.3/basic.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.3/basic.spvasm32 new file mode 100644 index 00000000..b47fbbf5 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.3/basic.spvasm32 @@ -0,0 +1,33 @@ +; SPIR-V +; Version: 1.3 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 18 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %9 "test_basic" %gl_GlobalInvocationID + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %gl_GlobalInvocationID Constant + OpDecorate %gl_GlobalInvocationID LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %8 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint +%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input + %9 = OpFunction %void None %8 + %10 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %11 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %12 = OpLabel + %13 = OpLoad %v3uint %gl_GlobalInvocationID Aligned 16 + %14 = OpCompositeExtract %uint %13 0 + %15 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %11 %14 + %16 = OpLoad %uint %15 Aligned 4 + %17 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %10 %14 + OpStore %17 %16 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.3/basic.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.3/basic.spvasm64 new file mode 100644 index 00000000..ba5d232c --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.3/basic.spvasm64 @@ -0,0 +1,38 @@ +; SPIR-V +; Version: 1.3 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 22 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %10 "test_basic" %gl_GlobalInvocationID + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %gl_GlobalInvocationID Constant + OpDecorate %gl_GlobalInvocationID LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + %ulong = OpTypeInt 64 0 + %uint = OpTypeInt 32 0 + %v3ulong = OpTypeVector %ulong 3 +%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %9 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint +%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3ulong Input + %10 = OpFunction %void None %9 + %11 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %12 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %13 = OpLabel + %14 = OpLoad %v3ulong %gl_GlobalInvocationID Aligned 32 + %15 = OpCompositeExtract %ulong %14 0 + %16 = OpUConvert %uint %15 + %17 = OpSConvert %ulong %16 + %18 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %12 %17 + %19 = OpLoad %uint %18 Aligned 4 + %20 = OpSConvert %ulong %16 + %21 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %11 %20 + OpStore %21 %19 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/basic.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.4/basic.spvasm32 new file mode 100644 index 00000000..407ef51d --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/basic.spvasm32 @@ -0,0 +1,33 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 18 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %9 "test_basic" %gl_GlobalInvocationID + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %gl_GlobalInvocationID Constant + OpDecorate %gl_GlobalInvocationID LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %8 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint +%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input + %9 = OpFunction %void None %8 + %10 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %11 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %12 = OpLabel + %13 = OpLoad %v3uint %gl_GlobalInvocationID Aligned 16 + %14 = OpCompositeExtract %uint %13 0 + %15 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %11 %14 + %16 = OpLoad %uint %15 Aligned 4 + %17 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %10 %14 + OpStore %17 %16 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/basic.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.4/basic.spvasm64 new file mode 100644 index 00000000..c2debf9c --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/basic.spvasm64 @@ -0,0 +1,38 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 22 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %10 "test_basic" %gl_GlobalInvocationID + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %gl_GlobalInvocationID Constant + OpDecorate %gl_GlobalInvocationID LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + %ulong = OpTypeInt 64 0 + %uint = OpTypeInt 32 0 + %v3ulong = OpTypeVector %ulong 3 +%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %9 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint +%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3ulong Input + %10 = OpFunction %void None %9 + %11 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %12 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %13 = OpLabel + %14 = OpLoad %v3ulong %gl_GlobalInvocationID Aligned 32 + %15 = OpCompositeExtract %ulong %14 0 + %16 = OpUConvert %uint %15 + %17 = OpSConvert %ulong %16 + %18 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %12 %17 + %19 = OpLoad %uint %18 Aligned 4 + %20 = OpSConvert %ulong %16 + %21 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %11 %20 + OpStore %21 %19 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.5/basic.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.5/basic.spvasm32 new file mode 100644 index 00000000..6b51ad5f --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.5/basic.spvasm32 @@ -0,0 +1,33 @@ +; SPIR-V +; Version: 1.5 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 18 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %9 "test_basic" %gl_GlobalInvocationID + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %gl_GlobalInvocationID Constant + OpDecorate %gl_GlobalInvocationID LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %8 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint +%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input + %9 = OpFunction %void None %8 + %10 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %11 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %12 = OpLabel + %13 = OpLoad %v3uint %gl_GlobalInvocationID Aligned 16 + %14 = OpCompositeExtract %uint %13 0 + %15 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %11 %14 + %16 = OpLoad %uint %15 Aligned 4 + %17 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %10 %14 + OpStore %17 %16 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.5/basic.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.5/basic.spvasm64 new file mode 100644 index 00000000..fefc130c --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.5/basic.spvasm64 @@ -0,0 +1,38 @@ +; SPIR-V +; Version: 1.5 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 22 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %10 "test_basic" %gl_GlobalInvocationID + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %gl_GlobalInvocationID Constant + OpDecorate %gl_GlobalInvocationID LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + %ulong = OpTypeInt 64 0 + %uint = OpTypeInt 32 0 + %v3ulong = OpTypeVector %ulong 3 +%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %9 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint +%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3ulong Input + %10 = OpFunction %void None %9 + %11 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %12 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %13 = OpLabel + %14 = OpLoad %v3ulong %gl_GlobalInvocationID Aligned 32 + %15 = OpCompositeExtract %ulong %14 0 + %16 = OpUConvert %uint %15 + %17 = OpSConvert %ulong %16 + %18 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %12 %17 + %19 = OpLoad %uint %18 Aligned 4 + %20 = OpSConvert %ulong %16 + %21 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %11 %20 + OpStore %21 %19 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.6/basic.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.6/basic.spvasm32 new file mode 100644 index 00000000..ff5745ae --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.6/basic.spvasm32 @@ -0,0 +1,33 @@ +; SPIR-V +; Version: 1.6 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 18 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %9 "test_basic" %gl_GlobalInvocationID + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %gl_GlobalInvocationID Constant + OpDecorate %gl_GlobalInvocationID LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %8 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint +%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input + %9 = OpFunction %void None %8 + %10 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %11 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %12 = OpLabel + %13 = OpLoad %v3uint %gl_GlobalInvocationID Aligned 16 + %14 = OpCompositeExtract %uint %13 0 + %15 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %11 %14 + %16 = OpLoad %uint %15 Aligned 4 + %17 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %10 %14 + OpStore %17 %16 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/spirv_asm/spv1.6/basic.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.6/basic.spvasm64 new file mode 100644 index 00000000..bed9d3aa --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.6/basic.spvasm64 @@ -0,0 +1,38 @@ +; SPIR-V +; Version: 1.6 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 22 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical64 OpenCL + OpEntryPoint Kernel %10 "test_basic" %gl_GlobalInvocationID + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %gl_GlobalInvocationID Constant + OpDecorate %gl_GlobalInvocationID LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + %ulong = OpTypeInt 64 0 + %uint = OpTypeInt 32 0 + %v3ulong = OpTypeVector %ulong 3 +%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %9 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint +%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3ulong Input + %10 = OpFunction %void None %9 + %11 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %12 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %13 = OpLabel + %14 = OpLoad %v3ulong %gl_GlobalInvocationID Aligned 32 + %15 = OpCompositeExtract %ulong %14 0 + %16 = OpUConvert %uint %15 + %17 = OpSConvert %ulong %16 + %18 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %12 %17 + %19 = OpLoad %uint %18 Aligned 4 + %20 = OpSConvert %ulong %16 + %21 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %11 %20 + OpStore %21 %19 Aligned 4 + OpReturn + OpFunctionEnd diff --git a/test_conformance/spirv_new/test_basic_versions.cpp b/test_conformance/spirv_new/test_basic_versions.cpp new file mode 100644 index 00000000..afe17390 --- /dev/null +++ b/test_conformance/spirv_new/test_basic_versions.cpp @@ -0,0 +1,123 @@ +// +// Copyright (c) 2023 The Khronos Group Inc. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. +// + +#include "testBase.h" +#include "types.hpp" + +#include +#include +#include + +extern bool gVersionSkip; + +TEST_SPIRV_FUNC(basic_versions) +{ + cl_int error = CL_SUCCESS; + + MTdataHolder d(gRandomSeed); + + std::vector h_src(num_elements); + generate_random_data(kInt, h_src.size(), d, h_src.data()); + + clMemWrapper src = + clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, + h_src.size() * sizeof(cl_int), h_src.data(), &error); + test_error(error, "Unable to create source buffer"); + + clMemWrapper dst = + clCreateBuffer(context, 0, h_src.size() * sizeof(cl_int), NULL, &error); + test_error(error, "Unable to create destination buffer"); + + std::map mapILtoSubdir({ + { "SPIR-V_1.0", "" }, // SPIR-V 1.0 files are in the base directory + { "SPIR-V_1.1", "spv1.1" }, + { "SPIR-V_1.2", "spv1.2" }, + { "SPIR-V_1.3", "spv1.3" }, + { "SPIR-V_1.4", "spv1.4" }, + { "SPIR-V_1.5", "spv1.5" }, + { "SPIR-V_1.6", "spv1.6" }, + }); + + size_t sz = 0; + error = clGetDeviceInfo(deviceID, CL_DEVICE_IL_VERSION, 0, NULL, &sz); + test_error(error, "Unable to query device IL versions size"); + + std::string ilVersions; + ilVersions.resize(sz); + error = clGetDeviceInfo(deviceID, CL_DEVICE_IL_VERSION, sz, &ilVersions[0], + NULL); + test_error(error, "Unable to query device IL versions string"); + + for (auto& testCase : mapILtoSubdir) + { + if (gVersionSkip) + { + log_info(" Skipping version check for %s.\n", + testCase.first.c_str()); + } + else if (ilVersions.find(testCase.first) == std::string::npos) + { + log_info(" Version %s is not supported; skipping test.\n", + testCase.first.c_str()); + continue; + } + else + { + log_info(" testing %s...\n", testCase.first.c_str()); + } + + const cl_int zero = 0; + error = + clEnqueueFillBuffer(queue, dst, &zero, sizeof(zero), 0, + h_src.size() * sizeof(cl_int), 0, NULL, NULL); + test_error(error, "Unable to initialize destination buffer"); + + std::string filename = testCase.second + "/basic"; + + clProgramWrapper prog; + error = get_program_with_il(prog, deviceID, context, filename.c_str()); + test_error(error, "Unable to build SPIR-V program"); + + clKernelWrapper kernel = clCreateKernel(prog, "test_basic", &error); + test_error(error, "Unable to create SPIR-V kernel"); + + error |= clSetKernelArg(kernel, 0, sizeof(dst), &dst); + error |= clSetKernelArg(kernel, 1, sizeof(src), &src); + test_error(error, "Unable to set kernel arguments"); + + size_t global = num_elements; + error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0, + NULL, NULL); + test_error(error, "Unable to enqueue kernel"); + + std::vector h_dst(num_elements); + error = clEnqueueReadBuffer(queue, dst, CL_TRUE, 0, + h_dst.size() * sizeof(cl_int), h_dst.data(), + 0, NULL, NULL); + test_error(error, "Unable to read destination buffer"); + + for (int i = 0; i < num_elements; i++) + { + if (h_dst[i] != h_src[i]) + { + log_error("Values do not match at location %d\n", i); + return TEST_FAIL; + } + } + } + + return TEST_PASS; +}