The documentation says "an unsupported feature which is to store/load
pointers to/from memory[...]We plan to fix it in next major release
1.1.0" (docs/Beignet.mdwn:216-219)
Given that 1.1.0 has now been released, and does appear to fix this
(creating an array of pointers works in 1.1.2 but not in 1.0.3;
https://cgit.freedesktop.org/beignet/commit/?id=48e22ae4536fbf944f39ab13e8f1133d3f5edcc3),
should this note now be removed?
(This is the test case I used: would you like it making into a utest?)
/*output in 1.1.2:
built 0 built 0 kernel created 00 returned (7, 8, 7)
1 returned (2, 7, 6)
2 returned (0, 1, 0)
output in 1.0.3:
store i32 7, i32 addrspace(1)* %26, align 4, !tbaa !11
built 0 Illegal pointer which is not from a valid memory space.
Aborting...
*/
#include <stdio.h>
#include <CL/cl.h>
int main(int argc,char** argv)
{
cl_platform_id platforms[4];
cl_uint num_platforms;
cl_device_id device;
cl_int status, ret;
cl_context ctx;
cl_command_queue queue;
cl_program program,program2;
cl_kernel kernel;
cl_mem buffer[3];
cl_event kernel_finished;
FILE *f;
size_t n = 3,program_length;
int i;
cl_int test_data[3][3] = {{3, 8, 5},{2, 4, 6},{0, 1, 0}};
const char *kernel_source="__kernel void pointer_store_load"
"(__global int *p0, __global int *p1, __global int *p2)"
"{__global int * __private a[2];int i;"
"i=get_global_id(0);a[0]=p0;a[1]=p1;a[p2[i]][i]=7;}";
char *build_log;
ret = 2;
build_log=calloc(1001,1);
clGetPlatformIDs(4,platforms,&num_platforms);
for(i=0;i<num_platforms;i++){
status=clGetDeviceIDs(platforms[i],CL_DEVICE_TYPE_GPU,
1,&device,NULL);
if (status == CL_SUCCESS) {
break;
}
}
if (status != CL_SUCCESS) {
printf("Device not found");
exit(1);
}
ctx = clCreateContext(NULL, 1, &device, NULL, NULL,
&status);
if (status == CL_SUCCESS) {
queue = clCreateCommandQueue(ctx, device, 0,
&status);
if (status == CL_SUCCESS) {
program = clCreateProgramWithSource(ctx, 1,
&kernel_source, NULL, &status);
if (status == CL_SUCCESS) {
status = clCompileProgram(program, 1, &device,
"", 0, NULL, NULL, NULL, NULL);
clGetProgramBuildInfo(program,device,
CL_PROGRAM_BUILD_LOG,1000,build_log,NULL);
printf("built %i %s",status,build_log);
program2 = clLinkProgram(ctx, 1, &device,
"", 1,&program, NULL, NULL, NULL);
clGetProgramBuildInfo(program,device,
CL_PROGRAM_BUILD_LOG,1000,build_log,NULL);
printf("built %i %s",status,build_log);
if (status == CL_SUCCESS) {
kernel = clCreateKernel(program2, "pointer_store_load",
&status);
printf("kernel created %i",status);
if (status == CL_SUCCESS) {
for(i=0;i<3;i++){
buffer[i] = clCreateBuffer(ctx, CL_MEM_COPY_HOST_PTR,
4*n, test_data[i], &status);
}
if (status == CL_SUCCESS) {
for(i=0;i<3;i++){
status = clSetKernelArg(kernel, i,
sizeof(cl_mem), &buffer[i]);
}
if (status == CL_SUCCESS) {
status = clEnqueueNDRangeKernel(queue, kernel, 1,
NULL, &n, &n, 0, NULL, &kernel_finished);
if (status == CL_SUCCESS) {
for(i=0;i<3;i++){
status = clEnqueueReadBuffer(queue, buffer[i],
CL_TRUE, 0, n*4, test_data[i], 1,
&kernel_finished, NULL);
}
if (status == CL_SUCCESS) {
for(i=0;i<3;i++){
printf("%i returned (%i, %i, %i)\n", i,
test_data[i][0], test_data[i][1],
test_data[i][2]);
}
}
}
}
}
for(i=0;i<3;i++){
clReleaseMemObject(buffer[i]);
}
}
clReleaseKernel(kernel);
}
}
clReleaseProgram(program);
}
clReleaseCommandQueue(queue);
}
clReleaseContext(ctx);
}
_______________________________________________
Beignet mailing list
Beignet@lists.freedesktop.org
https://lists.freedesktop.org/mailman/listinfo/beignet
