I tried to bind a GPUarray to a texture,
to re-use results from previous kernel calls.
Example code which should explain my problem:
# -------------------
import pycuda.autoinit
import pycuda.driver as cuda
import pycuda.reduction
import pycuda.gpuarray
from pycuda.gpuarray import GPUArray
import numpy
#cuda code
cuda_code = """ \
texture<float, 2> texData;
#define IMUL(a, b) __mul24(a, b)
__global__ void copy(
float *res,int res_pitch,
int dataW,int dataH){
const int x = IMUL(blockDim.x, blockIdx.x) + threadIdx.x;
const int y = IMUL(blockDim.y, blockIdx.y) + threadIdx.y;
if(x < dataW && y < dataH){
res[IMUL(y, res_pitch) + x] = tex2D(texData, x , y);
}
}
"""
cu_module = pycuda.compiler.SourceModule(cuda_code)
cu_copy = cu_module.get_function('copy')
cu_tex = cu_module.get_texref("texData");
#configure texture !
cu_tex.set_filter_mode(cuda.filter_mode.POINT)
cu_tex.set_address_mode(0, cuda.address_mode.CLAMP)
cu_tex.set_address_mode(1, cuda.address_mode.CLAMP)
#---------------- Case 1 from numpy Data ------------
#generate Data
A = numpy.random.random((20,10)).astype('float32')
#copy to the gpu & bind to texture
cuda.matrix_to_texref(A, cu_tex, order="C")
B = GPUArray( A.shape, 'float32');
width = numpy.int32( A.shape[1] )
height = numpy.int32( A.shape[0] )
pitch = numpy.int32( width )
#copy A to B
cu_copy(B.gpudata, pitch, width, height,
block=(1,1,1),grid=(10,20),texrefs=[cu_tex])
#A equals B
print numpy.linalg.norm(A-B.get())
# ------------- Case 2 from a GPUArray ----------
C = GPUArray( B.shape, 'float32');
#bint B to cu_tex
B.bind_to_texref_ext(cu_tex,channels=2);
#copy B to C
cu_copy(C.gpudata, pitch, width, height,
block=(1,1,1),grid=(10,20),texrefs=[cu_tex])
#A equals NOT B
print numpy.linalg.norm(B.get()-C.get())
#--------------------------------------
Is the binding of a GPUArray to a Texture not supported,
Or did I miss something ??
Cu
_______________________________________________
PyCUDA mailing list
[email protected]
http://tiker.net/mailman/listinfo/pycuda_tiker.net
