1 /******************************************************************
2 Copyright (c) 2016 The Khronos Group Inc. All Rights Reserved.
3
4 This code is protected by copyright laws and contains material proprietary to the Khronos Group, Inc.
5 This is UNPUBLISHED PROPRIETARY SOURCE CODE that may not be disclosed h_in whole or h_in part to
6 third parties, and may not be reproduced, republished, distributed, transmitted, displayed,
7 broadcast or otherwise exploited h_in any manner without the express prior written permission
8 of Khronos Group. The receipt or possession of this code does not convey any rights to reproduce,
9 disclose, or distribute its contents, or to manufacture, use, or sell anything that it may describe,
10 h_in whole or h_in part other than under the terms of the Khronos Adopters Agreement
11 or Khronos Conformance Test Source License Agreement as executed between Khronos and the recipient.
12 ******************************************************************/
13
14 #include "testBase.h"
15 #include "types.hpp"
16
17 #include <sstream>
18 #include <string>
19
test_linkage_compile(cl_device_id deviceID,cl_context context,cl_command_queue queue,const char * fname,clProgramWrapper & prog)20 static int test_linkage_compile(cl_device_id deviceID,
21 cl_context context,
22 cl_command_queue queue,
23 const char *fname,
24 clProgramWrapper &prog)
25 {
26 cl_int err = CL_SUCCESS;
27 std::vector<unsigned char> buffer_vec = readSPIRV(fname);
28
29 int file_bytes = buffer_vec.size();
30 if (file_bytes == 0) {
31 log_error("File not found\n");
32 return -1;
33 }
34 unsigned char *buffer = &buffer_vec[0];
35
36 if (gCoreILProgram)
37 {
38 prog = clCreateProgramWithIL(context, buffer, file_bytes, &err);
39 SPIRV_CHECK_ERROR(
40 err, "Failed to create program with clCreateProgramWithIL");
41 }
42 else
43 {
44 cl_platform_id platform;
45 err = clGetDeviceInfo(deviceID, CL_DEVICE_PLATFORM,
46 sizeof(cl_platform_id), &platform, NULL);
47 SPIRV_CHECK_ERROR(err,
48 "Failed to get platform info with clGetDeviceInfo");
49 clCreateProgramWithILKHR_fn clCreateProgramWithILKHR = NULL;
50
51 clCreateProgramWithILKHR = (clCreateProgramWithILKHR_fn)
52 clGetExtensionFunctionAddressForPlatform(
53 platform, "clCreateProgramWithILKHR");
54 if (clCreateProgramWithILKHR == NULL)
55 {
56 log_error(
57 "ERROR: clGetExtensionFunctionAddressForPlatform failed\n");
58 return -1;
59 }
60 prog = clCreateProgramWithILKHR(context, buffer, file_bytes, &err);
61 SPIRV_CHECK_ERROR(
62 err, "Failed to create program with clCreateProgramWithILKHR");
63 }
64
65 err = clCompileProgram(prog, 1, &deviceID,
66 NULL, // options
67 0, // num headers
68 NULL, // input headers
69 NULL, // header include names
70 NULL, // callback
71 NULL // User data
72 );
73 SPIRV_CHECK_ERROR(err, "Failed to compile spv program");
74 return 0;
75 }
76
TEST_SPIRV_FUNC(linkage_export_function_compile)77 TEST_SPIRV_FUNC(linkage_export_function_compile)
78 {
79 clProgramWrapper prog;
80 return test_linkage_compile(deviceID, context, queue, "linkage_export", prog);
81 }
82
TEST_SPIRV_FUNC(linkage_import_function_compile)83 TEST_SPIRV_FUNC(linkage_import_function_compile)
84 {
85 clProgramWrapper prog;
86 return test_linkage_compile(deviceID, context, queue, "linkage_import", prog);
87 }
88
TEST_SPIRV_FUNC(linkage_import_function_link)89 TEST_SPIRV_FUNC(linkage_import_function_link)
90 {
91 int err = 0;
92
93 clProgramWrapper prog_export;
94 err = test_linkage_compile(deviceID, context, queue, "linkage_export", prog_export);
95 SPIRV_CHECK_ERROR(err, "Failed to compile export program");
96
97 clProgramWrapper prog_import;
98 err = test_linkage_compile(deviceID, context, queue, "linkage_import", prog_import);
99 SPIRV_CHECK_ERROR(err, "Failed to compile import program");
100
101 cl_program progs[] = {prog_export, prog_import};
102
103 clProgramWrapper prog = clLinkProgram(context, 1, &deviceID, NULL, 2, progs, NULL, NULL, &err);
104 SPIRV_CHECK_ERROR(err, "Failed to link programs");
105
106 clKernelWrapper kernel = clCreateKernel(prog, "test_linkage", &err);
107 SPIRV_CHECK_ERROR(err, "Failed to create spv kernel");
108
109 const int num = 1 << 20;
110 std::vector<cl_float> h_in(num);
111 RandomSeed seed(gRandomSeed);
112 for (int i = 0; i < num; i++) {
113 h_in[i] = genrand<cl_float>(seed);
114 }
115
116 size_t bytes = sizeof(cl_float) * num;
117 clMemWrapper in = clCreateBuffer(context, CL_MEM_READ_WRITE, bytes, NULL, &err);
118 SPIRV_CHECK_ERROR(err, "Failed to create in buffer");
119
120 err = clEnqueueWriteBuffer(queue, in, CL_TRUE, 0, bytes, &h_in[0], 0, NULL, NULL);
121 SPIRV_CHECK_ERROR(err, "Failed to copy to in buffer");
122
123 err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &in);
124 SPIRV_CHECK_ERROR(err, "Failed to set arg 1");
125
126
127 size_t global = num;
128 err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0, NULL, NULL);
129 SPIRV_CHECK_ERROR(err, "Failed to enqueue cl kernel");
130
131 std::vector<cl_float> h_out(num);
132 err = clEnqueueReadBuffer(queue, in, CL_TRUE, 0, bytes, &h_out[0], 0, NULL, NULL);
133 SPIRV_CHECK_ERROR(err, "Failed to read to output");
134
135 for (int i = 0; i < num; i++) {
136 if (h_out[i] != -h_in[i]) {
137 log_error("Values do not match at location %d\n", i);
138 return -1;
139 }
140 }
141
142 return 0;
143 }
144