4 // Licensed under the Apache License, Version 2.0 (the "License");
8 //    http://www.apache.org/licenses/LICENSE-2.0
38   return -1;  in AtomicLoadExplicit()
50   return -1;  in AtomicFetchAddExplicit()
62   return -1;  in AtomicExchangeExplicit()
68   "typedef struct Node {\n"
71   "    __global struct Node* pNext;\n"
72   "} Node;\n"
74   // The allocation_index parameter must be initialized on the host to N work-items
76 …"__kernel void create_linked_lists(__global Node* pNodes, volatile __attribute__((nosvm)) __global…
79   "    __global Node *pNode = &pNodes[i];\n"
81   "    pNode->global_id = i;\n"
82   "    pNode->position_in_list = 0;\n"
84   "    __global Node *pNew;\n"
87   "        pNew = &pNodes[ atomic_inc(allocation_index) ];// allocate a new node\n"
88   "        pNew->global_id = i;\n"
89   "        pNew->position_in_list = j;\n"
90   "        pNode->pNext = pNew;  // link new node onto end of list\n"
95 …"__kernel void verify_linked_lists(__global Node* pNodes, volatile __global uint* num_correct, int…
98   "    __global Node *pNode = &pNodes[i];\n"
102   "        if( pNode->global_id == i && pNode->position_in_list == j)\n"
109   "        pNode = pNode->pNext;\n"
116 void create_linked_lists(Node* pNodes, size_t num_lists, int list_length)  in create_linked_lists()
122     Node *pNode = &pNodes[i];  in create_linked_lists()
123     pNode->global_id = i;  in create_linked_lists()
124     pNode->position_in_list = 0;  in create_linked_lists()
125     Node *pNew;  in create_linked_lists()
128       pNew = &pNodes[ allocation_index++ ];// allocate a new node  in create_linked_lists()
129       pNew->global_id = i;  in create_linked_lists()
130       pNew->position_in_list = j;  in create_linked_lists()
131       pNode->pNext = pNew;  // link new node onto end of list  in create_linked_lists()
137 cl_int verify_linked_lists(Node* pNodes, size_t num_lists, int list_length)  in verify_linked_lists()
145     Node *pNode = &pNodes[i];  in verify_linked_lists()
148       if( pNode->global_id == i && pNode->position_in_list == j)  in verify_linked_lists()
155       pNode = pNode->pNext;  in verify_linked_lists()
160     error = -1;  in verify_linked_lists()
201     Version version = get_device_cl_version(devices[i]);  in create_cl_objects()  local
203     if(device_from_harness != devices[i] && version < Version(2,0))  in create_cl_objects()
213       return -1;  in create_cl_objects()
220             log_error("Required extension not found - device id %d - %s\n", i,  in create_cl_objects()
277     ADD_TEST_VERSION( svm_migrate, Version(2, 1)),
283   auto version = get_device_cl_version(device);  in InitCL()  local
284   auto expected_min_version = Version(2, 0);  in InitCL()
285   if (version < expected_min_version)  in InitCL()
289                             version.to_string().c_str());  in InitCL()
302   if ((svm_caps == 0) && (version >= Version(3, 0)))  in InitCL()