1 //
2 // Copyright (c) 2017 The Khronos Group Inc.
3 //
4 // Licensed under the Apache License, Version 2.0 (the "License");
5 // you may not use this file except in compliance with the License.
6 // You may obtain a copy of the License at
7 //
8 // http://www.apache.org/licenses/LICENSE-2.0
9 //
10 // Unless required by applicable law or agreed to in writing, software
11 // distributed under the License is distributed on an "AS IS" BASIS,
12 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13 // See the License for the specific language governing permissions and
14 // limitations under the License.
15 //
16 #include "common.h"
17
18 const char *SVMPointerPassing_test_kernel[] = {
19 "__kernel void verify_char(__global uchar* pChar, volatile __global uint* num_correct, uchar expected)\n"
20 "{\n"
21 " if(0 == get_global_id(0))\n"
22 " {\n"
23 " *num_correct = 0;\n"
24 " if(*pChar == expected)\n"
25 " {\n"
26 " *num_correct=1;\n"
27 " }\n"
28 " }\n"
29 "}\n"
30 };
31
32
33 // Test that arbitrarily aligned char pointers into shared buffers can be passed directly to a kernel.
34 // This iterates through a buffer passing a pointer to each location to the kernel.
35 // The buffer is initialized to known values at each location.
36 // The kernel checks that it finds the expected value at each location.
37 // TODO: possibly make this work across all base types (including typeN?), also check ptr arithmetic ++,--.
test_svm_pointer_passing(cl_device_id deviceID,cl_context context2,cl_command_queue queue,int num_elements)38 int test_svm_pointer_passing(cl_device_id deviceID, cl_context context2, cl_command_queue queue, int num_elements)
39 {
40 clContextWrapper context = NULL;
41 clProgramWrapper program = NULL;
42 cl_uint num_devices = 0;
43 cl_int error = CL_SUCCESS;
44 clCommandQueueWrapper queues[MAXQ];
45
46 error = create_cl_objects(deviceID, &SVMPointerPassing_test_kernel[0], &context, &program, &queues[0], &num_devices, CL_DEVICE_SVM_COARSE_GRAIN_BUFFER);
47 if(error) return -1;
48
49 clKernelWrapper kernel_verify_char = clCreateKernel(program, "verify_char", &error);
50 test_error(error,"clCreateKernel failed");
51
52 size_t bufSize = 256;
53 cl_uchar *pbuf_svm_alloc = (cl_uchar*) clSVMAlloc(context, CL_MEM_READ_WRITE, sizeof(cl_uchar)*bufSize, 0);
54
55 cl_int *pNumCorrect = NULL;
56 pNumCorrect = (cl_int*) clSVMAlloc(context, CL_MEM_READ_WRITE, sizeof(cl_int), 0);
57
58 {
59 clMemWrapper buf = clCreateBuffer(context, CL_MEM_USE_HOST_PTR, sizeof(cl_uchar)*bufSize, pbuf_svm_alloc, &error);
60 test_error(error, "clCreateBuffer failed.");
61
62 clMemWrapper num_correct = clCreateBuffer(context, CL_MEM_USE_HOST_PTR, sizeof(cl_int), pNumCorrect, &error);
63 test_error(error, "clCreateBuffer failed.");
64
65 error = clSetKernelArg(kernel_verify_char, 1, sizeof(void*), (void *) &num_correct);
66 test_error(error, "clSetKernelArg failed");
67
68 // put values into buf so that we can expect to see these values in the kernel when we pass a pointer to them.
69 cl_command_queue cmdq = queues[0];
70 cl_uchar* pbuf_map_buffer = (cl_uchar*) clEnqueueMapBuffer(cmdq, buf, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0, sizeof(cl_uchar)*bufSize, 0, NULL,NULL, &error);
71 test_error2(error, pbuf_map_buffer, "clEnqueueMapBuffer failed");
72 for(int i = 0; i<(int)bufSize; i++)
73 {
74 pbuf_map_buffer[i]= (cl_uchar)i;
75 }
76 error = clEnqueueUnmapMemObject(cmdq, buf, pbuf_map_buffer, 0,NULL,NULL);
77 test_error(error, "clEnqueueUnmapMemObject failed.");
78
79 for (cl_uint ii = 0; ii<num_devices; ++ii) // iterate over all devices in the platform.
80 {
81 cmdq = queues[ii];
82 for(int i = 0; i<(int)bufSize; i++)
83 {
84 cl_uchar* pChar = &pbuf_svm_alloc[i];
85 error = clSetKernelArgSVMPointer(kernel_verify_char, 0, pChar); // pass a pointer to a location within the buffer
86 test_error(error, "clSetKernelArg failed");
87 error = clSetKernelArg(kernel_verify_char, 2, sizeof(cl_uchar), (void *) &i ); // pass the expected value at the above location.
88 test_error(error, "clSetKernelArg failed");
89 error = clEnqueueNDRangeKernel(cmdq, kernel_verify_char, 1, NULL, &bufSize, NULL, 0, NULL, NULL);
90 test_error(error,"clEnqueueNDRangeKernel failed");
91
92 pNumCorrect = (cl_int*) clEnqueueMapBuffer(cmdq, num_correct, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0, sizeof(cl_int), 0, NULL,NULL, &error);
93 test_error2(error, pNumCorrect, "clEnqueueMapBuffer failed");
94 cl_int correct_count = *pNumCorrect;
95 error = clEnqueueUnmapMemObject(cmdq, num_correct, pNumCorrect, 0,NULL,NULL);
96 test_error(error, "clEnqueueUnmapMemObject failed.");
97
98 if(correct_count != 1)
99 {
100 log_error("Passing pointer directly to kernel for byte #%d failed on device %d\n", i, ii);
101 return -1;
102 }
103 }
104 }
105
106 error = clFinish(cmdq);
107 test_error(error, "clFinish failed");
108 }
109
110
111 clSVMFree(context, pbuf_svm_alloc);
112 clSVMFree(context, pNumCorrect);
113
114 return 0;
115 }
116