1 /******************************************************************
2 Copyright (c) 2020 The Khronos Group Inc. All Rights Reserved.
3
4 This code is protected by copyright laws and contains material proprietary to
5 the Khronos Group, Inc. This is UNPUBLISHED PROPRIETARY SOURCE CODE that may not
6 be disclosed in whole or in part to third parties, and may not be reproduced,
7 republished, distributed, transmitted, displayed, broadcast or otherwise
8 exploited in any manner without the express prior written permission of Khronos
9 Group. The receipt or possession of this code does not convey any rights to
10 reproduce, disclose, or distribute its contents, or to manufacture, use, or sell
11 anything that it may describe, in whole or in part other than under the terms of
12 the Khronos Adopters Agreement or Khronos Conformance Test Source License
13 Agreement as executed between Khronos and the recipient.
14 ******************************************************************/
15
16 #include "testBase.h"
17
18 const char *sample_kernel_code_single_line[] = {
19 "__kernel void sample_test(__global float *src, __global int *dst)\n"
20 "{\n"
21 " int tid = get_global_id(0);\n"
22 "\n"
23 " dst[tid] = (int)src[tid];\n"
24 "\n"
25 "}\n"
26 };
27
TEST_SPIRV_FUNC(get_program_il)28 TEST_SPIRV_FUNC(get_program_il)
29 {
30 clProgramWrapper source_program;
31 size_t il_size = -1;
32 int error;
33
34 /* If a program has been created with clCreateProgramWithIL, CL_PROGRAM_IL
35 * should return the program IL it was created with and it's size */
36 if (gCoreILProgram || is_extension_available(deviceID, "cl_khr_il_program"))
37 {
38 clProgramWrapper il_program;
39 std::string spvStr = "op_function_none";
40 const char *spvName = spvStr.c_str();
41
42 std::vector<unsigned char> spirv_binary = readSPIRV(spvName);
43
44 int file_bytes = spirv_binary.size();
45 if (file_bytes == 0)
46 {
47 test_fail("ERROR: SPIRV file %s not found!\n", spvName);
48 }
49
50 /* Create program with IL */
51 unsigned char *spirv_buffer = &spirv_binary[0];
52
53 error = get_program_with_il(il_program, deviceID, context, spvName);
54
55 SPIRV_CHECK_ERROR(error, "Unable to create program with IL.");
56 if (il_program == NULL)
57 {
58 test_fail("ERROR: Unable to create test program!\n");
59 }
60
61 /* Check program IL is the same as the source IL */
62 unsigned char *buffer = new unsigned char[file_bytes];
63 error = clGetProgramInfo(il_program, CL_PROGRAM_IL, file_bytes, buffer,
64 &il_size);
65 SPIRV_CHECK_ERROR(error, "Unable to get program info.");
66
67 if (il_size != file_bytes)
68 {
69 test_fail("ERROR: Returned IL size is not the same as source IL "
70 "size (%lu "
71 "!= %lu)!\n",
72 il_size, file_bytes);
73 }
74
75 if (memcmp(buffer, spirv_buffer, file_bytes) != 0)
76 {
77 test_fail("ERROR: Returned IL is not the same as source IL!\n");
78 }
79
80 delete[] buffer;
81 }
82
83 /* CL_PROGRAM_IL shouldn't return IL value unless program is created with
84 * clCreateProgramWithIL */
85 error = create_single_kernel_helper_create_program(
86 context, &source_program, 1, sample_kernel_code_single_line);
87 if (source_program == NULL)
88 {
89 test_fail("ERROR: Unable to create test program!\n");
90 }
91
92 if (gCompilationMode != kSpir_v)
93 {
94 error =
95 clGetProgramInfo(source_program, CL_PROGRAM_IL, 0, NULL, &il_size);
96 SPIRV_CHECK_ERROR(error, "Unable to get program il length");
97 if (il_size != 0)
98 {
99 test_fail(
100 "ERROR: Returned length of non-IL program IL is non-zero!\n");
101 }
102 }
103
104 return 0;
105 }