still a bug in texture (image2d)

This commit is contained in:
Jerome Kieffer 2012-10-15 22:27:35 +02:00
parent 83f558befd
commit 6df7575bd2
3 changed files with 1465 additions and 1279 deletions

View File

@ -131,6 +131,29 @@ lut_integrate( const __global float *weights,
};//if bins
};//end kernel
/**
* \brief Performs 1d azimuthal integration with full pixel splitting based on a LUT
*
* An image instensity value is spread across the bins according to the positions stored in the LUT.
* The lut_index contains the positions of the pixel in the input array
* Values of 0 in the mask are processed and values of 1 ignored as per PyFAI
*
* @param weights Float pointer to global memory storing the input image.
* @param bins Unsigned int: number of output bins wanted (and pre-calculated)
* @param lut_size Unsigned int: dimension of the look-up table
* @param lut Pointer to a struct of [("idx",uint32),("coef": float)] where idx is the 1d-index of input pixels and coef is the weight of that pixel
* @param do_dummy Bool/int: shall the dummy pixel be checked. Dummy pixel are pixels marked as bad and ignored
* @param dummy Float: value for bad pixels
* @param delta_dummy Float: precision for bad pixel value
* @param do_dark Bool/int: shall dark-current correction be applied ?
* @param dark Float pointer to global memory storing the dark image.
* @param do_flat Bool/int: shall flat-field correction be applied ? (could contain polarization corrections)
* @param flat Float pointer to global memory storing the flat image.
* @param outData Float pointer to the output 1D array with the weighted histogram
* @param outCount Float pointer to the output 1D array with the unweighted histogram
* @param outMerged Float pointer to the output 1D array with the diffractogram
*/
__kernel void
lut_integrate_single( const __global float *weights,
const uint bins,
@ -181,3 +204,82 @@ lut_integrate_single( const __global float *weights,
outMerge[i] = (float) sum_data / sum_count;
};//if bins
};//end kernel
/**
* \brief Performs 1d azimuthal integration with full pixel splitting based on a LUT
*
* An image instensity value is spread across the bins according to the positions stored in the LUT.
* The lut_index contains the positions of the pixel in the input array
* Values of 0 in the mask are processed and values of 1 ignored as per PyFAI
*
* @param weights Float pointer to global memory storing the input image.
* @param bins Unsigned int: number of output bins wanted (and pre-calculated)
* @param lut_size Unsigned int: dimension of the look-up table
* @param lut Pointer to a struct of [("idx",uint32),("coef": float)] where idx is the 1d-index of input pixels and coef is the weight of that pixel
* @param do_dummy Bool/int: shall the dummy pixel be checked. Dummy pixel are pixels marked as bad and ignored
* @param dummy Float: value for bad pixels
* @param delta_dummy Float: precision for bad pixel value
* @param do_dark Bool/int: shall dark-current correction be applied ?
* @param dark Float pointer to global memory storing the dark image.
* @param do_flat Bool/int: shall flat-field correction be applied ? (could contain polarization corrections)
* @param flat Float pointer to global memory storing the flat image.
* @param outData Float pointer to the output 1D array with the weighted histogram
* @param outCount Float pointer to the output 1D array with the unweighted histogram
* @param outMerged Float pointer to the output 1D array with the diffractogram
*/
__kernel void
lut_integrate_image( __read_only image2d_t weights,
const uint dimX,
const uint dimY,
const uint bins,
const uint lut_size,
const __global struct lut_point_t *lut,
const int do_dummy,
const float dummy,
const float delta_dummy,
const int do_dark,
const __global float *dark,
const int do_flat,
const __global float *flat,
__global float *outData,
__global float *outCount,
__global float *outMerge
)
{
uint idx, k, j, i= get_global_id(0);
bigfloat_t sum_data = 0.0;
bigfloat_t sum_count = 0.0;
const bigfloat_t epsilon = 1e-10;
float coef, data;
const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST;
if(i < bins)
{
for (j=0;j<lut_size;j++)
{
k = i*lut_size+j;
idx = lut[k].idx;
coef = lut[k].coef;
if((idx == 0) && (coef <= 0.0))
break;
data = read_imagef(weights, sampler, (int2)(idx/dimX , idx%dimX)).s0;
//data = weights[idx];
if( (!do_dummy) || (delta_dummy && (fabs(data-dummy) > delta_dummy))|| (data!=dummy) )
{
if(do_dark)
data -= dark[idx];
if(do_flat)
data /= flat[idx];
sum_data += coef * data;
sum_count += coef;
};//test dummy
};//for j
outData[i] = (float) sum_data;
outCount[i] = (float) sum_count;
if (sum_count > epsilon)
outMerge[i] = (float) sum_data / sum_count;
};//if bins
};//end kernel

File diff suppressed because it is too large Load Diff

View File

@ -20,9 +20,9 @@ edf = os.path.join(root, "LaB6_0020.edf")
img = fabio.open(edf)
ai = pyFAI.load(poni)
ai.xrpd(img.data, bins)
tth = ai._ttha.ravel().astype("float32")
dtth = ai._dttha.ravel().astype("float32")
data = img.data.ravel().astype("float32")
tth = ai._ttha.ravel().astype(numpy.float32)
dtth = ai._dttha.ravel().astype(numpy.float32)
data = img.data.ravel().astype(numpy.float32)
import splitBBox
t0 = time.time()
@ -80,19 +80,31 @@ pylab.plot(ra, rb, label="Original")
import pyopencl
mf = pyopencl.mem_flags
ct = pyopencl.channel_type
co = pyopencl.channel_order
ctx = pyopencl.create_some_context()
q = pyopencl.CommandQueue(ctx)
program = pyopencl.Program(ctx, open("../openCL/ocl_azim_LUT.cl").read()).build()
t3 = time.time()
weights_buf = pyopencl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=data)
#weights_img = pyopencl.image_from_array(ctx, ary=img.data.astype(numpy.float32), mode="r", norm_int=False, num_channels=1)
#print co.INTENSITY, ct.FLOAT,
#imf = pyopencl.ImageFormat(numpy.uint32(co.INTENSITY), numpy.uint32(ct.FLOAT))
#weights_img = pyopencl.Image(ctx, flags=mf.READ_ONLY | mf.COPY_HOST_PTR,
# format=imf,
# hostbuf=data,
# pitches=(img.data.shape[-1],))
#image_from_array(ctx, ary=img.data.astype(numpy.float32), mode="r", norm_int=False, num_channels=1)
#lut_idx_buf = pyopencl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=integ.lut_idx.astype(numpy.uint32))
#lut_coef_buf = pyopencl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=integ.lut_coef)
lut_buf = pyopencl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=integ.lut)
None_buf = pyopencl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=numpy.zeros(1, dtype=numpy.float32))
outData_buf = pyopencl.Buffer(ctx, mf.WRITE_ONLY, numpy.dtype("float32").itemsize * bins)
outCount_buf = pyopencl.Buffer(ctx, mf.WRITE_ONLY, numpy.dtype("float32").itemsize * bins)
outMerge_buf = pyopencl.Buffer(ctx, mf.WRITE_ONLY, numpy.dtype("float32").itemsize * bins)
args = (weights_buf,
outData_buf = pyopencl.Buffer(ctx, mf.WRITE_ONLY, numpy.dtype(numpy.float32).itemsize * bins)
outCount_buf = pyopencl.Buffer(ctx, mf.WRITE_ONLY, numpy.dtype(numpy.float32).itemsize * bins)
outMerge_buf = pyopencl.Buffer(ctx, mf.WRITE_ONLY, numpy.dtype(numpy.float32).itemsize * bins)
args = (#weights_img, numpy.uint32(img.dim1), numpy.uint32(img.dim0),
weights_buf,
numpy.uint32(2048),
numpy.uint32(integ.lut_size),
# lut_idx_buf,