[PyCUDA] A mini issue in pycuda: pycuda-helpers.hpp

[RA ZA]

Hi!

I've been working with double presicion textures and surfaces, and in
particular with 3D structures.

I found a little mistake in pycuda-helpers.hpp, specifically:

File: pycuda-helpers.hpp
Line: 45
Incorrect: __device__ double fp_tex3D(texture<fp_tex_double, 2, read_mode>
tex, int i, int j, int k)
Correct: __device__ double fp_tex3D(texture<fp_tex_double, *3*, read_mode>
tex, int i, int j, int k)

This change solve a problem when we use this template. Also, i would like
to contribute, but i fail when i try to add in this file
(pycuda-helpers.hpp) the corresponding functions to write and read from
Surfaces, i use a ugly '*If'* in my kernels to correctly use surfaces.
Example:
"""
if (doublePrec == 0){
surf3Dwrite( valueReal, surf_psi0OutReal, t_i*sizeof(float), t_j, t_k,
cudaBoundaryModeClamp);
}
else{
fp_tex_double vReal; // A internal way to use int2<->double, this can be
extended with templates, as in pycuda-helpers.hpp.
vReal.x = __double2loint(valueReal);
vReal.y = __double2hiint(valueReal);
surf3Dwrite( vReal, surf_psi0OutReal, t_i*sizeof(double), t_j, t_k,
cudaBoundaryModeClamp);
"""
Regards,
Roberto (from Mexico, sorry for my English)
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.tiker.net/pipermail/pycuda/attachments/20140428/36215e2c/attachment.html>

[Andreas Kloeckner]

Post by RA ZA
I've been working with double presicion textures and surfaces, and in
particular with 3D structures.
File: pycuda-helpers.hpp
Line: 45
Incorrect: __device__ double fp_tex3D(texture<fp_tex_double, 2, read_mode>
tex, int i, int j, int k)
Correct: __device__ double fp_tex3D(texture<fp_tex_double, *3*, read_mode>
tex, int i, int j, int k)
This change solve a problem when we use this template. Also, i would like
to contribute, but i fail when i try to add in this file
(pycuda-helpers.hpp) the corresponding functions to write and read from
Surfaces, i use a ugly '*If'* in my kernels to correctly use surfaces.
"""
if (doublePrec == 0){
surf3Dwrite( valueReal, surf_psi0OutReal, t_i*sizeof(float), t_j, t_k,
cudaBoundaryModeClamp);
}
else{
fp_tex_double vReal; // A internal way to use int2<->double, this can be
extended with templates, as in pycuda-helpers.hpp.
vReal.x = __double2loint(valueReal);
vReal.y = __double2hiint(valueReal);
surf3Dwrite( vReal, surf_psi0OutReal, t_i*sizeof(double), t_j, t_k,
cudaBoundaryModeClamp);
"""

Fixed in git. Thanks for the report!

Andreas
-------------- next part --------------
A non-text attachment was scrubbed...
Name: not available
Type: application/pgp-signature
Size: 810 bytes
Desc: not available
URL: <http://lists.tiker.net/pipermail/pycuda/attachments/20140428/3c2c9368/attachment.sig>