Function BOOST_AUTO_TEST_CASE

compute/test/test_local_buffer.cpp:24–78

namespace compute = boost::compute;

BOOST_AUTO_TEST_CASE(local_buffer_arg)
{
    if(device.get_info<CL_DEVICE_LOCAL_MEM_TYPE>() != CL_LOCAL){
        std::cerr << "skipping local buffer arg test: "
                  << "device does not support local memory" << std::endl;
        return;
    }

    const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
        __kernel void foo(__local float *local_buffer,
                          __global float *global_buffer)
        {
            const uint gid = get_global_id(0);
            const uint lid = get_local_id(0);

            local_buffer[lid] = gid;

            global_buffer[gid] = local_buffer[lid];
        }
    );

    // create and build program
    compute::program program =
        compute::program::build_with_source(source, context);

    // create kernel
    compute::kernel kernel = program.create_kernel("foo");

    // setup kernel arguments
    compute::buffer global_buf(context, 128 * sizeof(float));
    kernel.set_arg(0, compute::local_buffer<float>(32));
    kernel.set_arg(1, global_buf);

    // some implementations don't correctly report dynamically-set local buffer sizes
    if(kernel.get_work_group_info<cl_ulong>(device, CL_KERNEL_LOCAL_MEM_SIZE) == 0){
        std::cerr << "skipping checks for local memory size, device reports "
                  << "zero after setting dynamically-sized local buffer size"
                  << std::endl;
        return;
    }

    // check actual memory size
    BOOST_CHECK_GE(
        kernel.get_work_group_info<cl_ulong>(device, CL_KERNEL_LOCAL_MEM_SIZE),
        32 * sizeof(float)
    );

    // increase local buffer size and check new actual local memory size
    kernel.set_arg(0, compute::local_buffer<float>(64));

    BOOST_CHECK_GE(
        kernel.get_work_group_info<cl_ulong>(device, CL_KERNEL_LOCAL_MEM_SIZE),
        64 * sizeof(float)
    );
}

BOOST_AUTO_TEST_SUITE_END()