PyCUDA October 2013

pycuda@tiker.net

14 participants
19 discussions

by Andreas Kloeckner

Dear Nicolas, Nicolas LEMERCIER <lemercie(a)igbmc.fr> writes: > I am currently trying to use peer2peer GPU memory access with pycuda and I face with problems regarding the syntax. > My code follow this template: > > dev1=cuda.Device(1) > ctx1=dev1.make_context() > dev0=cuda.Device(0) > ctx0=dev0.make_context() > ctx0.enable_peer_access(ctx1,flags=0) > > and here is the error message returned by python. > > Traceback (most recent call last): > File "C:\Users\solid\Documents\peer2peer.py", line 34, in <module> > ctx0.enable_peer_access(ctx1,flags=0) > LogicError: cuCtxEnablePeerAccess failed: invalid/unknown error code > > I guess the parameter sent to enable_peer_access is not correct but I don't understand the way it works. first, can you please make sure you send emails of this nature are sent to the pycuda list. As far as parameters to this function, there's nothing much to mess up. As a result, have you checked that peer access is supported by these devices? Perhaps add assert dev1.can_access_peer(dev0) to make sure that peer access *should* be possible. Andreas

6 years, 8 months

Re: [PyCUDA] installing pycuda

by Eric Scheffel

Dear Geng, I think the solution is simple. It appears as if you did not have boost installed on your computer, which is a collection of important C++ libraries, some of which (matrix-related?) are used by pycuda. Install the latest boost and you should be okay, fingers crossed! I guess that this is the mistake, although I really would have expected C++ boost to be pre-installed on a standard Ubuntu installation. Greetings from Ningbo, Eric On 10/26/2013 04:49 PM, geng liu wrote: Dear Developer, I am LIU Geng from Chinese Academy of Sciences. I am now trying to install pycuda in my Ubuntu system, but I met some problems. My computer has 2 CUDA enabled GPUs. My system is Lynux Ubuntu 12.04, gcc 4.6.3, python 2.7.3, CUDA 5.5 I followed the instructions from you, but there is an error when I go to "make -j 4". I am sure there is not a compatibility problem, at least for gcc and CUDA. My installation log is in the attachment, and I'm looking forward to a solution from you. Thanks. Best Regards Liu Geng Changchun Institue of Optics, Fine Mechanics and Physics, Chinese Academy of Sciences. Changchun 130033, Jilin, China Mobile: 86-0431-15843126536 Email: liugeng.5202004(a)gmail.com<mailto:liugeng.5202004@gmail.com> -- Dr. Eric Michael Scheffel Assistant Professor in Economics Ningbo Nottingham University Site: http://www.ericscheffel.com This message and any attachment are intended solely for the addressee and may contain confidential information. If you have received this message in error, please send it back to me, and immediately delete it. Please do not use, copy or disclose the information contained in this message or in any attachment. Any views or opinions expressed by the author of this email do not necessarily reflect the views of The University of Nottingham Ningbo China. This message has been checked for viruses but the contents of an attachment may still contain software viruses which could damage your computer system: you are advised to perform your own checks. Email communications with The University of Nottingham Ningbo China may be monitored as permitted by UK and Chinese legislation.

6 years, 8 months

installing pycuda

by geng liu

Dear Developer, I am LIU Geng from Chinese Academy of Sciences. I am now trying to install pycuda in my Ubuntu system, but I met some problems. My computer has 2 CUDA enabled GPUs. My system is Lynux Ubuntu 12.04, gcc 4.6.3, python 2.7.3, CUDA 5.5 I followed the instructions from you, but there is an error when I go to "make -j 4". I am sure there is not a compatibility problem, at least for gcc and CUDA. My installation log is in the attachment, and I'm looking forward to a solution from you. Thanks. Best Regards Liu Geng Changchun Institue of Optics, Fine Mechanics and Physics, Chinese Academy of Sciences. Changchun 130033, Jilin, China Mobile: 86-0431-15843126536 Email: liugeng.5202004(a)gmail.com

6 years, 8 months

cuda-gdb doesn't stop at breakpoint

by Vedran Miletić

Hi, I can't get cuda-gdb to stop at breakpoint. Namely, I get $ cuda-gdb --args python -m pycuda.debug /usr/share/doc/python-pycuda-doc/examples/demo.py NVIDIA (R) CUDA Debugger 5.5 release Portions Copyright (C) 2007-2013 NVIDIA Corporation GNU gdb (GDB) 7.2 Copyright (C) 2010 Free Software Foundation, Inc. License GPLv3+: GNU GPL version 3 or later <http://gnu.org/licenses/gpl.html> This is free software: you are free to change and redistribute it. There is NO WARRANTY, to the extent permitted by law. Type "show copying" and "show warranty" for details. This GDB was configured as "x86_64-unknown-linux-gnu". For bug reporting instructions, please see: <http://www.gnu.org/software/gdb/bugs/>... Reading symbols from /usr/bin/python...(no debugging symbols found)...done. (cuda-gdb) b doublify Function "doublify" not defined. Make breakpoint pending on future shared library load? (y or [n]) y Breakpoint 1 (doublify) pending. (cuda-gdb) r Starting program: /usr/bin/python -m pycuda.debug /usr/share/doc/python-pycuda-doc/examples/demo.py warning: Could not load shared library symbols for linux-vdso.so.1. Do you need "set solib-search-path" or "set sysroot"? [Thread debugging using libthread_db enabled] [New Thread 0x7ffff0001700 (LWP 3395)] [Context Create of context 0xfb38d0 on Device 0] *** compiler output in /tmp/tmpiHClwv [Launch of CUDA Kernel 0 (doublify<<<(1,1,1),(4,4,1)>>>) on Device 0] original array: [[-0.04012666 0.32504103 -0.35359469 -1.33113611] [-0.77782524 1.76319683 1.71368265 1.25119114] [ 1.13762009 -1.3670094 -0.31529811 -0.29683253] [-0.48716831 0.80230588 1.03378654 2.5838871 ]] doubled with kernel: [[-0.08025332 0.65008205 -0.70718938 -2.66227221] [-1.55565047 3.52639365 3.4273653 2.50238228] [ 2.27524018 -2.7340188 -0.63059622 -0.59366506] [-0.97433662 1.60461175 2.06757307 5.1677742 ]] [Launch of CUDA Kernel 1 (doublify<<<(1,1,1),(4,4,1)>>>) on Device 0] doubled with InOut: [[-0.08025332 0.65008205 -0.70718938 -2.66227221] [-1.55565047 3.52639365 3.4273653 2.50238228] [ 2.27524018 -2.7340188 -0.63059622 -0.59366506] [-0.97433662 1.60461175 2.06757307 5.1677742 ]] *** compiler output in /tmp/tmp_v2lny [Launch of CUDA Kernel 2 (axpb<<<(1,1,1),(32,1,1)>>>) on Device 0] original array: [[ 0.43130487 0.56232756 0.78714377 -0.8899222 ] [ 0.12634102 -0.34702763 -0.47055304 -0.35157502] [ 0.3904736 -0.70519412 -0.00533707 -1.11191225] [ 1.07405782 0.69153821 0.42390665 -0.45493162]] doubled with gpuarray: [[ 0.86260974 1.12465513 1.57428753 -1.7798444 ] [ 0.25268203 -0.69405526 -0.94110608 -0.70315003] [ 0.78094721 -1.41038823 -0.01067414 -2.2238245 ] [ 2.14811563 1.38307643 0.84781331 -0.90986323]] [Context Pop of context 0xfb38d0 on Device 0] [Context Push of context 0xfb38d0 on Device 0] [Context Pop of context 0xfb38d0 on Device 0] [Context Push of context 0xfb38d0 on Device 0] [Context Pop of context 0xfb38d0 on Device 0] [Context Push of context 0xfb38d0 on Device 0] [Context Pop of context 0xfb38d0 on Device 0] [Context Push of context 0xfb38d0 on Device 0] [Termination of CUDA Kernel 2 (axpb<<<(1,1,1),(32,1,1)>>>) on Device 0] [Termination of CUDA Kernel 1 (doublify<<<(1,1,1),(4,4,1)>>>) on Device 0] [Termination of CUDA Kernel 0 (doublify<<<(1,1,1),(4,4,1)>>>) on Device 0] [Context Destroy of context 0xfb38d0 on Device 0] [Thread 0x7ffff0001700 (LWP 3395) exited] Program exited normally. I'm using Debian jessie/sid amd64. I have tried both 319.60 and 325.15 drivers (and many other versions previously), and CUDA toolkit 5.0.35 and 5.5.22. No change on this front. I'm using PyCUDA 2013.1.1. Any ideas? Regards, Vedran

6 years, 8 months

Re: [PyCUDA] pycuda on optimus enabled cards

by Graham Mills

---------------------------------------------------------------------- Dorin Niculescu <niculescu_dorin85(a)yahoo.com><mailto:niculescu_dorin85@yahoo.com> writes: > I have a new ASUS laptop with optimus enabled NVIDIA 750M card and i want to install pycuda on it. I've installed Ubuntu 12.04, > Bumblebee+nvidia-319, cuda 5.5 and everything was working great until i've installed pycuda using the ubuntu software center. After that, > instead of the log in screen ubuntu is showing a perfect black screen. I don't know what happend but i think that is because of the libraries > which are installed with pycuda. > > Is pycuda getting along with bumblebee? This chain Ubuntu 12.04 + Bumblebee + nvidia-319 + cuda5.5 + pycuda should work? > > Do you have any solutions ? Sorry to hear that you're having trouble. However, what you're describing sounds more like an issue with your distribution than with PyCUDA. If you could find out somehow what packages were installed and uninstalled along with PyCUDA, someone on the list *may* be able to point you in the right direction. HTH, Andreas Dear Both, I am running pycuda on an up-to-date Slackware distribution on an Optimus laptop (Lenovo Y470) with a gt550m. All works very well here. I am using bumblebee and the "optirun" command. The most common problem people tend to face with their respective Linux distributions is that the cuda drivers (SDK) have traditionally been very picky about which version of GCC they support. Most modern Linux distros will install some version of GCC 4.8 or perhaps even later, but CUDA on Linux afaik currently only supports 4.7 and maybe 4.6, depending on the cuda version installed. HTH, Eric -- Dr. Eric Michael Scheffel Assistant Professor in Economics Ningbo Nottingham University Site: http://www.ericscheffel.com This message and any attachment are intended solely for the addressee and may contain confidential information. If you have received this message in error, please send it back to me, and immediately delete it. Please do not use, copy or disclose the information contained in this message or in any attachment. Any views or opinions expressed by the author of this email do not necessarily reflect the views of The University of Nottingham Ningbo China. This message has been checked for viruses but the contents of an attachment may still contain software viruses which could damage your computer system: you are advised to perform your own checks. Email communications with The University of Nottingham Ningbo China may be monitored as permitted by UK and Chinese legislation. -------------------------------------------- Dorin, I had a similar symptom on a desktop running the same Ubuntu release a few weeks ago, and was able to resolve it by purging and reinstalling my nvidia driver and adding nouveau.modeset=0 to GRUB to ensure that the nvidia driver was the only one being used at boot time. The video driver I use was packaged with the cuda5.5 .run file from Nvidia's download page, version 319.37. The pycuda version which shows up for me in the Ubuntu software center is 2012.1.1, which might be causing you problems if that's the one you installed. I am using 2013.1.1, downloaded from https://pypi.python.org/pypi/pycuda . HTH, Graham

6 years, 8 months

pycuda on optimus enabled cards

by Dorin Niculescu

Hello, I have a new ASUS laptop with optimus enabled NVIDIA 750M card and i want to install pycuda on it. I've installed Ubuntu 12.04, Bumblebee+nvidia-319, cuda 5.5 and everything was working great until i've installed pycuda using the ubuntu software center. After that, instead of the log in screen ubuntu is showing a perfect black screen. I don't know what happend but i think that is because of the libraries which are installed with pycuda. Is pycuda getting along with bumblebee? This chain Ubuntu 12.04 + Bumblebee + nvidia-319 + cuda5.5 + pycuda should work? Do you have any solutions ? Thank you very much for your time, Dorin

6 years, 8 months

Re: [PyCUDA] how do I use module_from_file

by Andreas Kloeckner

Matthias Seise <mseise(a)gmx.de> writes: > Hi, > > I'm using pycuda quite a while. Now I'm supposed to distribute my software > without installing a C++ compiler on all machines. How can I distribute > "compiled" kernels. > > I found the option module_from_file but when I'm using it on the standard > kernel 'multiply_them' I get > logicError: cuModuleLoad failed: invalid context > > I found the Thread > http://stackoverflow.com/questions/5918937/pycudathreading-invalid-handles-… > > But How do I get the correct function name? It's generally easier to declare your functions extern "C" than to guess the (C++-mangled) function names. HTH, Andreas

6 years, 8 months

how do I use module_from_file

by Matthias Seise

Hi, I'm using pycuda quite a while. Now I'm supposed to distribute my software without installing a C++ compiler on all machines. How can I distribute "compiled" kernels. I found the option module_from_file but when I'm using it on the standard kernel 'multiply_them' I get logicError: cuModuleLoad failed: invalid context I found the Thread http://stackoverflow.com/questions/5918937/pycudathreading-invalid-handles-… But How do I get the correct function name? Thanks for your help, Matthias

6 years, 8 months

Bicubic interpolation

by Eric Scheffel

Dear Andreas, Thanks again for making available the cuda wrapper library for python. It's great to use and helps me a lot in my own research. One problem I am facing at the moment is that I am trying to use Danny Ruijters' kernel code for using bicubic interpolation on nvidia GPU's, instead of the "built-in" bilinear ones. The code is available at http://www.dannyruijters.nl/cubicinterpolation/ Attaches is the kernel source code I am trying to run using pycuda. The problem is that even though it compiles using the "SourceModule" function (using the no_extern_c=True option), I cannot fetch the "compnkcross" kernel method using "get_function". I get the error message: In [15]: modo.get_function('compnkcross') --------------------------------------------------------------------------- LogicError Traceback (most recent call last) <ipython-input-15-e2bfb6bd9c2d> in <module>() ----> 1 modo.get_function('compnkcross') /usr/lib/python2.7/site-packages/pycuda-2013.1.1-py2.7-linux-i686.egg/pycuda/compiler.pyc in get_function(self, name) 283 284 def get_function(self, name): --> 285 return self.module.get_function(name) LogicError: cuModuleGetFunction failed: not found I CAN fetch the texture references, but that is about it. The original source code for the kernel is distributed over a number of cu-files as well as at least one header file. I have simply taken them and substituted out to create one big monolithic kernel code without external dependencies. Being paranoid about not exactly knowing what pycuda supports I have also converted template kernel functions into specific cases which I need for running my own code, though I suspect that pycuda of course supports template functions. Thanks, Eric (im Moment im kalten Karlsruhe) //////////////////////////////////////////////////////////////////////////////// typedef unsigned int uint; typedef unsigned short ushort; // float functions //////////////////////////////////////////////////////////////////////////////// // lerp inline __device__ __host__ float lerp(float a, float b, float t) { return a + t*(b-a); } // clamp inline __device__ __host__ float clamp(float f, float a, float b) { return fmaxf(a, fminf(f, b)); } // int2 functions //////////////////////////////////////////////////////////////////////////////// // addition inline __host__ __device__ int2 operator+(int2 a, int2 b) { return make_int2(a.x + b.x, a.y + b.y); } inline __host__ __device__ void operator+=(int2 &a, int2 b) { a.x += b.x; a.y += b.y; } // subtract inline __host__ __device__ int2 operator-(int2 a, int2 b) { return make_int2(a.x - b.x, a.y - b.y); } inline __host__ __device__ void operator-=(int2 &a, int2 b) { a.x -= b.x; a.y -= b.y; } // multiply inline __host__ __device__ int2 operator*(int2 a, int2 b) { return make_int2(a.x * b.x, a.y * b.y); } inline __host__ __device__ int2 operator*(int2 a, int s) { return make_int2(a.x * s, a.y * s); } inline __host__ __device__ int2 operator*(int s, int2 a) { return make_int2(a.x * s, a.y * s); } inline __host__ __device__ void operator*=(int2 &a, int s) { a.x *= s; a.y *= s; } // float2 functions //////////////////////////////////////////////////////////////////////////////// // additional constructors inline __host__ __device__ float2 make_float2(float s) { return make_float2(s, s); } inline __host__ __device__ float2 make_float2(int2 a) { return make_float2(float(a.x), float(a.y)); } // addition inline __host__ __device__ float2 operator+(float2 a, float2 b) { return make_float2(a.x + b.x, a.y + b.y); } inline __host__ __device__ void operator+=(float2 &a, float2 b) { a.x += b.x; a.y += b.y; } // subtract inline __host__ __device__ float2 operator-(float2 a, float2 b) { return make_float2(a.x - b.x, a.y - b.y); } inline __host__ __device__ void operator-=(float2 &a, float2 b) { a.x -= b.x; a.y -= b.y; } // multiply inline __host__ __device__ float2 operator*(float2 a, float2 b) { return make_float2(a.x * b.x, a.y * b.y); } inline __host__ __device__ float2 operator*(float2 a, float s) { return make_float2(a.x * s, a.y * s); } inline __host__ __device__ float2 operator*(float s, float2 a) { return make_float2(a.x * s, a.y * s); } inline __host__ __device__ void operator*=(float2 &a, float s) { a.x *= s; a.y *= s; } // divide inline __host__ __device__ float2 operator/(float2 a, float2 b) { return make_float2(a.x / b.x, a.y / b.y); } inline __host__ __device__ float2 operator/(float2 a, float s) { float inv = 1.0f / s; return a * inv; } inline __host__ __device__ float2 operator/(float s, float2 a) //Danny { // float inv = 1.0f / s; // return a * inv; return make_float2(s / a.x, s / a.y); } inline __host__ __device__ void operator/=(float2 &a, float s) { float inv = 1.0f / s; a *= inv; } // lerp inline __device__ __host__ float2 lerp(float2 a, float2 b, float t) { return a + t*(b-a); } // clamp inline __device__ __host__ float2 clamp(float2 v, float a, float b) { return make_float2(clamp(v.x, a, b), clamp(v.y, a, b)); } inline __device__ __host__ float2 clamp(float2 v, float2 a, float2 b) { return make_float2(clamp(v.x, a.x, b.x), clamp(v.y, a.y, b.y)); } // dot product inline __host__ __device__ float dot(float2 a, float2 b) { return a.x * b.x + a.y * b.y; } // length inline __host__ __device__ float length(float2 v) { return sqrtf(dot(v, v)); } // normalize inline __host__ __device__ float2 normalize(float2 v) { float invLen = 1.0f / sqrtf(dot(v, v)); return v * invLen; } // floor inline __host__ __device__ float2 floor(const float2 v) { return make_float2(floor(v.x), floor(v.y)); } // reflect inline __host__ __device__ float2 reflect(float2 i, float2 n) { return i - 2.0f * n * dot(n,i); } // float3 functions //////////////////////////////////////////////////////////////////////////////// // additional constructors inline __host__ __device__ float3 make_float3(float s) { return make_float3(s, s, s); } inline __host__ __device__ float3 make_float3(float2 a) { return make_float3(a.x, a.y, 0.0f); } inline __host__ __device__ float3 make_float3(float2 a, float s) { return make_float3(a.x, a.y, s); } inline __host__ __device__ float3 make_float3(float4 a) { return make_float3(a.x, a.y, a.z); // discards w } inline __host__ __device__ float3 make_float3(int3 a) { return make_float3(float(a.x), float(a.y), float(a.z)); } // min static __inline__ __host__ __device__ float3 fminf(float3 a, float3 b) { return make_float3(fminf(a.x,b.x), fminf(a.y,b.y), fminf(a.z,b.z)); } // max static __inline__ __host__ __device__ float3 fmaxf(float3 a, float3 b) { return make_float3(fmaxf(a.x,b.x), fmaxf(a.y,b.y), fmaxf(a.z,b.z)); } // addition inline __host__ __device__ float3 operator+(float3 a, float3 b) { return make_float3(a.x + b.x, a.y + b.y, a.z + b.z); } inline __host__ __device__ float3 operator+(float3 a, float b) { return make_float3(a.x + b, a.y + b, a.z + b); } inline __host__ __device__ void operator+=(float3 &a, float3 b) { a.x += b.x; a.y += b.y; a.z += b.z; } // subtract inline __host__ __device__ float3 operator-(float3 a, float3 b) { return make_float3(a.x - b.x, a.y - b.y, a.z - b.z); } inline __host__ __device__ float3 operator-(float3 a, float b) { return make_float3(a.x - b, a.y - b, a.z - b); } inline __host__ __device__ void operator-=(float3 &a, float3 b) { a.x -= b.x; a.y -= b.y; a.z -= b.z; } // multiply inline __host__ __device__ float3 operator*(float3 a, float3 b) { return make_float3(a.x * b.x, a.y * b.y, a.z * b.z); } inline __host__ __device__ float3 operator*(float3 a, float s) { return make_float3(a.x * s, a.y * s, a.z * s); } inline __host__ __device__ float3 operator*(float s, float3 a) { return make_float3(a.x * s, a.y * s, a.z * s); } inline __host__ __device__ void operator*=(float3 &a, float s) { a.x *= s; a.y *= s; a.z *= s; } // divide inline __host__ __device__ float3 operator/(float3 a, float3 b) { return make_float3(a.x / b.x, a.y / b.y, a.z / b.z); } inline __host__ __device__ float3 operator/(float3 a, float s) { float inv = 1.0f / s; return a * inv; } inline __host__ __device__ float3 operator/(float s, float3 a) //Danny { // float inv = 1.0f / s; // return a * inv; return make_float3(s / a.x, s / a.y, s / a.z); } inline __host__ __device__ void operator/=(float3 &a, float s) { float inv = 1.0f / s; a *= inv; } // lerp inline __device__ __host__ float3 lerp(float3 a, float3 b, float t) { return a + t*(b-a); } // clamp inline __device__ __host__ float3 clamp(float3 v, float a, float b) { return make_float3(clamp(v.x, a, b), clamp(v.y, a, b), clamp(v.z, a, b)); } inline __device__ __host__ float3 clamp(float3 v, float3 a, float3 b) { return make_float3(clamp(v.x, a.x, b.x), clamp(v.y, a.y, b.y), clamp(v.z, a.z, b.z)); } // dot product inline __host__ __device__ float dot(float3 a, float3 b) { return a.x * b.x + a.y * b.y + a.z * b.z; } // cross product inline __host__ __device__ float3 cross(float3 a, float3 b) { return make_float3(a.y*b.z - a.z*b.y, a.z*b.x - a.x*b.z, a.x*b.y - a.y*b.x); } // length inline __host__ __device__ float length(float3 v) { return sqrtf(dot(v, v)); } // normalize inline __host__ __device__ float3 normalize(float3 v) { float invLen = 1.0f / sqrtf(dot(v, v)); return v * invLen; } // floor inline __host__ __device__ float3 floor(const float3 v) { return make_float3(floor(v.x), floor(v.y), floor(v.z)); } // reflect inline __host__ __device__ float3 reflect(float3 i, float3 n) { return i - 2.0f * n * dot(n,i); } // float4 functions //////////////////////////////////////////////////////////////////////////////// // additional constructors inline __host__ __device__ float4 make_float4(float s) { return make_float4(s, s, s, s); } inline __host__ __device__ float4 make_float4(float3 a) { return make_float4(a.x, a.y, a.z, 0.0f); } inline __host__ __device__ float4 make_float4(float3 a, float w) { return make_float4(a.x, a.y, a.z, w); } inline __host__ __device__ float4 make_float4(int4 a) { return make_float4(float(a.x), float(a.y), float(a.z), float(a.w)); } // min static __inline__ __host__ __device__ float4 fminf(float4 a, float4 b) { return make_float4(fminf(a.x,b.x), fminf(a.y,b.y), fminf(a.z,b.z), fminf(a.w,b.w)); } // max static __inline__ __host__ __device__ float4 fmaxf(float4 a, float4 b) { return make_float4(fmaxf(a.x,b.x), fmaxf(a.y,b.y), fmaxf(a.z,b.z), fmaxf(a.w,b.w)); } // addition inline __host__ __device__ float4 operator+(float4 a, float4 b) { return make_float4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w); } inline __host__ __device__ void operator+=(float4 &a, float4 b) { a.x += b.x; a.y += b.y; a.z += b.z; a.w += b.w; } // subtract inline __host__ __device__ float4 operator-(float4 a, float4 b) { return make_float4(a.x - b.x, a.y - b.y, a.z - b.z, a.w - b.w); } inline __host__ __device__ void operator-=(float4 &a, float4 b) { a.x -= b.x; a.y -= b.y; a.z -= b.z; a.w -= b.w; } // multiply inline __host__ __device__ float4 operator*(float4 a, float s) { return make_float4(a.x * s, a.y * s, a.z * s, a.w * s); } inline __host__ __device__ float4 operator*(float s, float4 a) { return make_float4(a.x * s, a.y * s, a.z * s, a.w * s); } inline __host__ __device__ void operator*=(float4 &a, float s) { a.x *= s; a.y *= s; a.z *= s; a.w *= s; } // divide inline __host__ __device__ float4 operator/(float4 a, float4 b) { return make_float4(a.x / b.x, a.y / b.y, a.z / b.z, a.w / b.w); } inline __host__ __device__ float4 operator/(float4 a, float s) { float inv = 1.0f / s; return a * inv; } inline __host__ __device__ float4 operator/(float s, float4 a) //Danny { // float inv = 1.0f / s; // return a * inv; return make_float4(s / a.x, s / a.y, s / a.z, s / a.w); } inline __host__ __device__ void operator/=(float4 &a, float s) { float inv = 1.0f / s; a *= inv; } // lerp inline __device__ __host__ float4 lerp(float4 a, float4 b, float t) { return a + t*(b-a); } // clamp inline __device__ __host__ float4 clamp(float4 v, float a, float b) { return make_float4(clamp(v.x, a, b), clamp(v.y, a, b), clamp(v.z, a, b), clamp(v.w, a, b)); } inline __device__ __host__ float4 clamp(float4 v, float4 a, float4 b) { return make_float4(clamp(v.x, a.x, b.x), clamp(v.y, a.y, b.y), clamp(v.z, a.z, b.z), clamp(v.w, a.w, b.w)); } // dot product inline __host__ __device__ float dot(float4 a, float4 b) { return a.x * b.x + a.y * b.y + a.z * b.z + a.w * b.w; } // length inline __host__ __device__ float length(float4 r) { return sqrtf(dot(r, r)); } // normalize inline __host__ __device__ float4 normalize(float4 v) { float invLen = 1.0f / sqrtf(dot(v, v)); return v * invLen; } // floor inline __host__ __device__ float4 floor(const float4 v) { return make_float4(floor(v.x), floor(v.y), floor(v.z), floor(v.w)); } // int3 functions //////////////////////////////////////////////////////////////////////////////// // additional constructors inline __host__ __device__ int3 make_int3(int s) { return make_int3(s, s, s); } inline __host__ __device__ int3 make_int3(float3 a) { return make_int3(int(a.x), int(a.y), int(a.z)); } // min inline __host__ __device__ int3 min(int3 a, int3 b) { return make_int3(min(a.x,b.x), min(a.y,b.y), min(a.z,b.z)); } // max inline __host__ __device__ int3 max(int3 a, int3 b) { return make_int3(max(a.x,b.x), max(a.y,b.y), max(a.z,b.z)); } // addition inline __host__ __device__ int3 operator+(int3 a, int3 b) { return make_int3(a.x + b.x, a.y + b.y, a.z + b.z); } inline __host__ __device__ void operator+=(int3 &a, int3 b) { a.x += b.x; a.y += b.y; a.z += b.z; } // subtract inline __host__ __device__ int3 operator-(int3 a, int3 b) { return make_int3(a.x - b.x, a.y - b.y, a.z - b.z); } inline __host__ __device__ void operator-=(int3 &a, int3 b) { a.x -= b.x; a.y -= b.y; a.z -= b.z; } // multiply inline __host__ __device__ int3 operator*(int3 a, int3 b) { return make_int3(a.x * b.x, a.y * b.y, a.z * b.z); } inline __host__ __device__ int3 operator*(int3 a, int s) { return make_int3(a.x * s, a.y * s, a.z * s); } inline __host__ __device__ int3 operator*(int s, int3 a) { return make_int3(a.x * s, a.y * s, a.z * s); } inline __host__ __device__ void operator*=(int3 &a, int s) { a.x *= s; a.y *= s; a.z *= s; } // divide inline __host__ __device__ int3 operator/(int3 a, int3 b) { return make_int3(a.x / b.x, a.y / b.y, a.z / b.z); } inline __host__ __device__ int3 operator/(int3 a, int s) { return make_int3(a.x / s, a.y / s, a.z / s); } inline __host__ __device__ int3 operator/(int s, int3 a) { return make_int3(a.x / s, a.y / s, a.z / s); } inline __host__ __device__ void operator/=(int3 &a, int s) { a.x /= s; a.y /= s; a.z /= s; } // clamp inline __device__ __host__ int clamp(int f, int a, int b) { return max(a, min(f, b)); } inline __device__ __host__ int3 clamp(int3 v, int a, int b) { return make_int3(clamp(v.x, a, b), clamp(v.y, a, b), clamp(v.z, a, b)); } inline __device__ __host__ int3 clamp(int3 v, int3 a, int3 b) { return make_int3(clamp(v.x, a.x, b.x), clamp(v.y, a.y, b.y), clamp(v.z, a.z, b.z)); } // uint3 functions //////////////////////////////////////////////////////////////////////////////// // additional constructors inline __host__ __device__ uint3 make_uint3(uint s) { return make_uint3(s, s, s); } inline __host__ __device__ uint3 make_uint3(float3 a) { return make_uint3(uint(a.x), uint(a.y), uint(a.z)); } // min inline __host__ __device__ uint3 min(uint3 a, uint3 b) { return make_uint3(min(a.x,b.x), min(a.y,b.y), min(a.z,b.z)); } // max inline __host__ __device__ uint3 max(uint3 a, uint3 b) { return make_uint3(max(a.x,b.x), max(a.y,b.y), max(a.z,b.z)); } // addition inline __host__ __device__ uint3 operator+(uint3 a, uint3 b) { return make_uint3(a.x + b.x, a.y + b.y, a.z + b.z); } inline __host__ __device__ void operator+=(uint3 &a, uint3 b) { a.x += b.x; a.y += b.y; a.z += b.z; } // subtract inline __host__ __device__ uint3 operator-(uint3 a, uint3 b) { return make_uint3(a.x - b.x, a.y - b.y, a.z - b.z); } inline __host__ __device__ void operator-=(uint3 &a, uint3 b) { a.x -= b.x; a.y -= b.y; a.z -= b.z; } // multiply inline __host__ __device__ uint3 operator*(uint3 a, uint3 b) { return make_uint3(a.x * b.x, a.y * b.y, a.z * b.z); } inline __host__ __device__ uint3 operator*(uint3 a, uint s) { return make_uint3(a.x * s, a.y * s, a.z * s); } inline __host__ __device__ uint3 operator*(uint s, uint3 a) { return make_uint3(a.x * s, a.y * s, a.z * s); } inline __host__ __device__ void operator*=(uint3 &a, uint s) { a.x *= s; a.y *= s; a.z *= s; } // divide inline __host__ __device__ uint3 operator/(uint3 a, uint3 b) { return make_uint3(a.x / b.x, a.y / b.y, a.z / b.z); } inline __host__ __device__ uint3 operator/(uint3 a, uint s) { return make_uint3(a.x / s, a.y / s, a.z / s); } inline __host__ __device__ uint3 operator/(uint s, uint3 a) { return make_uint3(a.x / s, a.y / s, a.z / s); } inline __host__ __device__ void operator/=(uint3 &a, uint s) { a.x /= s; a.y /= s; a.z /= s; } // clamp inline __device__ __host__ uint clamp(uint f, uint a, uint b) { return max(a, min(f, b)); } inline __device__ __host__ uint3 clamp(uint3 v, uint a, uint b) { return make_uint3(clamp(v.x, a, b), clamp(v.y, a, b), clamp(v.z, a, b)); } inline __device__ __host__ uint3 clamp(uint3 v, uint3 a, uint3 b) { return make_uint3(clamp(v.x, a.x, b.x), clamp(v.y, a.y, b.y), clamp(v.z, a.z, b.z)); } typedef unsigned int uint; typedef unsigned short ushort; typedef unsigned char uchar; typedef signed char schar; inline __device__ __host__ uint UMIN(uint a, uint b) { return a < b ? a : b; } inline __device__ __host__ uint PowTwoDivider(uint n) { if (n == 0) return 0; uint divider = 1; while ((n & divider) == 0) divider <<= 1; return divider; } inline __host__ __device__ float2 operator-(float a, float2 b) { return make_float2(a - b.x, a - b.y); } inline __host__ __device__ float3 operator-(float a, float3 b) { return make_float3(a - b.x, a - b.y, a - b.z); } // Cubic B-spline function // The 3rd order Maximal Order and Minimum Support function, that it is maximally differentiable. inline __host__ __device__ float bspline(float t) { t = fabs(t); const float a = 2.0f - t; if (t < 1.0f) return 2.0f/3.0f - 0.5f*t*t*a; else if (t < 2.0f) return a*a*a / 6.0f; else return 0.0f; } // The first order derivative of the cubic B-spline inline __host__ __device__ float bspline_1st_derivative(float t) { if (-2.0f < t && t <= -1.0f) return 0.5f*t*t + 2.0f*t + 2.0f; else if (-1.0f < t && t <= 0.0f) return -1.5f*t*t - 2.0f*t; else if ( 0.0f < t && t <= 1.0f) return 1.5f*t*t - 2.0f*t; else if ( 1.0f < t && t < 2.0f) return -0.5f*t*t + 2.0f*t - 2.0f; else return 0.0f; } // The second order derivative of the cubic B-spline inline __host__ __device__ float bspline_2nd_derivative(float t) { t = fabs(t); if (t < 1.0f) return 3.0f*t - 2.0f; else if (t < 2.0f) return 2.0f - t; else return 0.0f; } // Inline calculation of the bspline convolution weights, without conditional statements inline __device__ void bspline_weights(float fraction, float& w0, float& w1, float& w2, float& w3) { const float one_frac = 1.0f - fraction; const float squared = fraction * fraction; const float one_sqd = one_frac * one_frac; w0 = 1.0f/6.0f * one_sqd * one_frac; w1 = 2.0f/3.0f - 0.5f * squared * (2.0f-fraction); w2 = 2.0f/3.0f - 0.5f * one_sqd * (2.0f-one_frac); w3 = 1.0f/6.0f * squared * fraction; } // Inline calculation of the first order derivative bspline convolution weights, without conditional statements inline __device__ void bspline_weights_1st_derivative(float fraction, float& w0, float& w1, float& w2, float& w3) { const float squared = fraction * fraction; w0 = -0.5f * squared + fraction - 0.5f; w1 = 1.5f * squared - 2.0f * fraction; w2 = -1.5f * squared + fraction + 0.5f; w3 = 0.5f * squared; } // Inline calculation of the second order derivative bspline convolution weights, without conditional statements inline __device__ void bspline_weights_2nd_derivative(float fraction, float& w0, float& w1, float& w2, float& w3) { w0 = 1.0f - fraction; w1 = 3.0f * fraction - 2.0f; w2 = -3.0f * fraction + 1.0f; w3 = fraction; } __device__ float cubicTex2DSimple(texture<float, 2> tex, float x, float y) { // transform the coordinate from [0,extent] to [-0.5, extent-0.5] const float2 coord_grid = make_float2(x - 0.5f, y - 0.5f); float2 index = floor(coord_grid); const float2 fraction = coord_grid - index; index.x += 0.5f; //move from [-0.5, extent-0.5] to [0, extent] index.y += 0.5f; //move from [-0.5, extent-0.5] to [0, extent] float result = 0.0f; for (float y=-1; y < 2.5f; y++) { float bsplineY = bspline(y-fraction.y); float v = index.y + y; for (float x=-1; x < 2.5f; x++) { float bsplineXY = bspline(x-fraction.x) * bsplineY; float u = index.x + x; result += bsplineXY * tex2D(tex, u, v); } } return result; } #define CUBICTEX2D cubicTex2DSimple texture<float, 2> polbb, polbg, polgb, polgg; __global__ void compnkcross(int N, int ngrid, int ndeg2, float Theta, float K_min, float K_max, float k_min, float k_max, float K, int asvec, float *y, float *x, int *empstat) { int col = blockIdx.x*blockDim.x + threadIdx.x; if (col < N && empstat[col] == 0 && asvec == 0) { y[col] = CUBICTEX2D(polbb, ((K-K_min)/(K_max-K_min))*ndeg2, powf(((x[col]-k_min)/(k_max-k_min)),(1.0f/Theta))*ngrid); } if (col < N && empstat[col] == 0 && asvec == 1) { y[col] = CUBICTEX2D(polgb, ((K-K_min)/(K_max-K_min))*ndeg2, powf(((x[col]-k_min)/(k_max-k_min)),(1.0f/Theta))*ngrid); } if (col < N && empstat[col] == 1 && asvec == 0) { y[col] = CUBICTEX2D(polbg, ((K-K_min)/(K_max-K_min))*ndeg2, powf(((x[col]-k_min)/(k_max-k_min)),(1.0f/Theta))*ngrid); } if (col < N && empstat[col] == 1 && asvec == 1) { y[col] = CUBICTEX2D(polgg, ((K-K_min)/(K_max-K_min))*ndeg2, powf(((x[col]-k_min)/(k_max-k_min)),(1.0f/Theta))*ngrid); } } #undef CUBICTEX2D -- Dr. Eric Michael Scheffel Assistant Professor in Economics Ningbo Nottingham University Site: http://www.ericscheffel.com --------------------------- Dr. Eric Michael Scheffel Assistant Professor in Economics Ningbo Nottingham University Site: http://www.ericscheffel.com This message and any attachment are intended solely for the addressee and may contain confidential information. If you have received this message in error, please send it back to me, and immediately delete it. Please do not use, copy or disclose the information contained in this message or in any attachment. Any views or opinions expressed by the author of this email do not necessarily reflect the views of The University of Nottingham Ningbo China. This message has been checked for viruses but the contents of an attachment may still contain software viruses which could damage your computer system: you are advised to perform your own checks. Email communications with The University of Nottingham Ningbo China may be monitored as permitted by UK and Chinese legislation.

6 years, 9 months

2020

2019

2018

2017

2016

2015

2014

2013

2012

2011

2010

2009

2008

PyCUDA October 2013