[Pycuda] LogicError: cuMemcpyDtoH failed: an illegal memory access was encountered
by Jie Liu
Hallo,
I have a piece of very simple code written in Pycuda. Can anyone tell
me, why shouldn't I set the index of array CC as "c = wA *
%(BLOCK_SIZE)d * by + %(BLOCK_SIZE)d * bx"? For example, if I set the
index of CC as 1 or 2 or 3, it can get the right value. Thank you very much!
importpycuda.autoinit
frompycuda importdriver, compiler, gpuarray, tools
frompycuda.compiler importSourceModule
importnumpy asnp
kernel_code_template ="""
#include <cuComplex.h>
__device__ void ExtractVector(cuFloatComplex *mul_a, int &idt,
cuFloatComplex *A){
for (int i = 0; i < %(MATRIX_SIZE_O)d; i ++)
{
A[i] = make_cuFloatComplex(0,0);
A[i] = make_cuFloatComplex(cuCrealf(mul_a[i*%(tSize)d+
idt]),cuCimagf(mul_a[i*%(tSize)d+ idt]));
}
__syncthreads();
}
__global__ void MatrixMulKernel(cuFloatComplex *A, cuFloatComplex *CC)
{
const uint wA = %(MATRIX_SIZE_O)d;
const uint wB = %(MATRIX_SIZE_O)d;
// Block index
const uint bx = blockIdx.x;
const uint by = blockIdx.y;
// Thread index
const uint tx = threadIdx.x;
const uint ty = threadIdx.y;
cuFloatComplex Csub = make_cuFloatComplex(0,0);
__shared__ cuFloatComplex VectorA[%(MATRIX_SIZE_O)d];
__shared__ cuFloatComplex C1[%(MATRIX_SIZE_O)d];
const uint aBegin = wA * %(BLOCK_SIZE)d* by;
const uint aEnd = aBegin + wA - 1;
const uint aStep = 1;
const int bBegin = %(BLOCK_SIZE)d* bx;
const uint bStep = wB;
const uint c = wA * %(BLOCK_SIZE)d* by + %(BLOCK_SIZE)d* bx ;
for (int index_r = 0; index_r < %(rSize)d; index_r ++)
{
for (int index_t = 0; index_t < %(tSize)d; index_t ++)
{
ExtractVector(A, index_t, VectorA);
CC[c] = make_cuFloatComplex(cuCrealf(VectorA[c]),cuCimagf(VectorA[c]));
}
}
}
"""
theta =(np.arange(-80, 80, 40))/360*2*np.pi
tsize =theta.size
rmax =300
r =np.arange(0, rmax, rmax/3)
rsize =r.size
Mt =1
Mr =6
MATRIX_SIZE_O =Mt*Mr
MATRIX_SIZE_I =1
BLOCK_SIZE =1
BLOCK_SIZE_x =1
BLOCK_SIZE_y =1
GRID_SIZE =MATRIX_SIZE_O
kernel_code =kernel_code_template %{
'MATRIX_SIZE_O': MATRIX_SIZE_O,
'BLOCK_SIZE': BLOCK_SIZE,
'tSize': tsize,
'rSize': rsize,
}
a_cpu =1+1j*1+np.arange(0, tsize *MATRIX_SIZE_O)
a_cpu =a_cpu.reshape(MATRIX_SIZE_O,tsize).astype(np.complex64)
b_cpu =1+1j*2+np.arange(0, MATRIX_SIZE_O *MATRIX_SIZE_O)
b_cpu =b_cpu.reshape(MATRIX_SIZE_O,MATRIX_SIZE_O).astype(np.complex64)
bb_cpu =1+1j*3+np.arange(0, tsize *MATRIX_SIZE_O)
bb_cpu =bb_cpu.reshape(MATRIX_SIZE_O, tsize).astype(np.complex64)
f_cpu =1+1j*4+np.arange(0, rsize *MATRIX_SIZE_O)
f_cpu =f_cpu.reshape(MATRIX_SIZE_O,rsize).astype(np.complex64)
Bcon =np.real(a_cpu).astype(np.float32)
+1j*np.imag(a_cpu).astype(np.float32)
a_gpu =gpuarray.to_gpu(Bcon)
Rinv =np.real(b_cpu).astype(np.float32)
+1j*np.imag(b_cpu).astype(np.float32)
b_gpu =gpuarray.to_gpu(Rinv)
FFT_cpu =np.real(f_cpu).astype(np.float32)
+1j*np.imag(f_cpu).astype(np.float32)
f_gpu =gpuarray.to_gpu(FFT_cpu)
B =np.real(bb_cpu).astype(np.float32) +1j*np.imag(bb_cpu).astype(np.float32)
bb_gpu =gpuarray.to_gpu(B)
beta_gpu =gpuarray.empty((rsize,tsize), np.complex64)
c_gpu =gpuarray.empty((MATRIX_SIZE_O), np.complex64)
beta_cpu =np.zeros(shape=(rsize,tsize)).astype(np.complex64)
forindex_r inrange(0, rsize):
forindex_theta inrange(0, theta.size):
c_cpu =np.dot(a_cpu[:, index_theta], b_cpu)
beta_cpu[index_r, index_theta] =np.dot(c_cpu, f_cpu[:, index_r])
#/ (np.dot(c_cpu, bb_cpu[:, index_theta]))
#beta_cpu[index_r, index_theta] = np.dot(c_cpu, bb_cpu[:, index_theta])
mod =compiler.SourceModule(kernel_code)
matrixmul =mod.get_function("MatrixMulKernel")
matrixmul(
# inputs
a_gpu,
# outputs
c_gpu,
# grid of multiple blocks
grid=(GRID_SIZE, GRID_SIZE),
# block of multiple threads
block=(BLOCK_SIZE_x , BLOCK_SIZE_y , 1)
)
# print the results
print("-"*80)
print("Matrix A (GPU): ")
print(a_gpu.get())
print("-"*80)
print("Matrix B (GPU): ")
print(bb_gpu.get())
print("-"*80)
print("Matrix C (GPU): ")
print(c_gpu.get())
print("-"*80)
print("Matrix C (CPU): ")
print(c_cpu)
print("-"*80)
print("Matrix beta (GPU): ")
print(beta_gpu.get())
print("-"*80)
print("Matrix beta (CPU): ")
print(beta_cpu)
print("-"*80)
print("CPU-GPU Difference: ")
print(beta_cpu/beta_gpu.get())
1 month, 3 weeks
Re: Question about PyCUDA
by Andreas Kloeckner
"thierry.moudiki" <thierry.moudiki(a)protonmail.com> writes:
> Hi Andreas,
>
> I'm interested in using your package PyCUDA, and I have one question about it (just to make sure that I understand how it works). In the example presented here: https://documen.tician.de/pycuda/index.html, when you call the sourced function `multiply_them`:
>
> - The option `block` has 3 elements in case it's a 1D, 2D or 3D block,
> and each tuple element is the number of threads per block. E.g: if I
> use threadIdx.x and threadIdx.y with blocks of 400 threads each, will
> I have `block = (400, 400, 1)`?
Yup.
> - For the option `grid`, I'm not sure. Is it: we can have 1D or 2D
> grids? What about this tuple's elements? Tuple element == Number of
> blocks per grid?
Yup.
In the future, please address requests like this to the mailing
list. I've cc'd them for archival.
Andreas
2 months
Re: gpuarray.zeros / to_gpu crash
by Andreas Kloeckner
Dan,
Please make sure the list stays cc'd for archival.
"Guralnik,Dan" <danguralnik(a)ufl.edu> writes:
> Andreas, I'm so sorry, should've done it myself so you have more info. Here is what happens:
>
> ---------------------
> C:\Users\danguralnik\Documents\GitHub\kodlab-uma-sims\mice\smooth>python cuda_test.py
> Traceback (most recent call last):
> File "cuda_test.py", line 23, in <module>
> """)
> File "C:\ProgramData\Anaconda3\lib\site-packages\pycuda\compiler.py", line 291, in __init__
> arch, code, cache_dir, include_dirs)
> File "C:\ProgramData\Anaconda3\lib\site-packages\pycuda\compiler.py", line 254, in compile
> return compile_plain(source, options, keep, nvcc, cache_dir, target)
> File "C:\ProgramData\Anaconda3\lib\site-packages\pycuda\compiler.py", line 137, in compile_plain
> stderr=stderr.decode("utf-8", "replace"))
> pycuda.driver.CompileError: nvcc compilation of C:\Users\DANGUR~1\AppData\Local\Temp\tmppoj82k3c\kernel.cu failed
> [command: nvcc --cubin -arch sm_75 -m64 -Ic:\programdata\anaconda3\lib\site-packages\pycuda\cuda kernel.cu]
> [stdout:
> nvcc fatal : Cannot find compiler 'cl.exe' in PATH
> ]
> ------------------------
> Is this something about environment variables?
Yes. The Visual Studio compilers also need to be on your PATH
environment variable. VS installs a batch file vcvars.bat that should
arrange for that.
Andreas
2 months
gpuarray.zeros / to_gpu crash
by Guralnik,Dan
Hello,
I have just installed pycuda on a new machine running anaconda3 and cuda 10.1.243_426 for Windows 10, and tested the installation by running a program that had been run successfully on another machine, but with anaconda2.
My program breaks on a call for gpuarray.zeros, following a call to gpuarray.to_gpu (no other calls to pycuda have been made before that). Both calls are trying to establish 2D arrays of dtype=np.float32.
The included headers are:
import pycuda.driver as cuda
import pycuda.tools
import pycuda.autoinit
import pycuda.gpuarray as gpuarray
from pycuda.compiler import SourceModule
I've attached the error trace from python below. I'll be grateful if you could offer an explanation/fix.
Many thanks,
-Dan
C:\Users\danguralnik\Documents\GitHub\kodlab-uma-sims\mice\smooth>python mouse_base.py
Traceback (most recent call last):
File "C:\ProgramData\Anaconda3\lib\site-packages\pycuda\tools.py", line 428, in context_dependent_memoize
return ctx_dict[cur_ctx][args]
KeyError: <pycuda._driver.Context object at 0x000001A0CCE991B0>
During handling of the above exception, another exception occurred:
Traceback (most recent call last):
File "mouse_base.py", line 832, in <module>
main()
File "mouse_base.py", line 812, in main
random_state=RS(),
File "mouse_base.py", line 587, in __init__
self.addRepn('global_view',Repn_global_arena_wmouse(self))
File "mouse_base.py", line 641, in __init__
gpuarray.zeros(shape=(self.ysize,self.xsize),dtype=np.float32)
File "C:\ProgramData\Anaconda3\lib\site-packages\pycuda\gpuarray.py", line 1068, in zeros
result.fill(zero)
File "C:\ProgramData\Anaconda3\lib\site-packages\pycuda\gpuarray.py", line 549, in fill
func = elementwise.get_fill_kernel(self.dtype)
File "<C:\ProgramData\Anaconda3\lib\site-packages\decorator.py:decorator-gen-13>", line 2, in get_fill_kernel
File "C:\ProgramData\Anaconda3\lib\site-packages\pycuda\tools.py", line 432, in context_dependent_memoize
result = func(*args)
File "C:\ProgramData\Anaconda3\lib\site-packages\pycuda\elementwise.py", line 496, in get_fill_kernel
"fill")
File "C:\ProgramData\Anaconda3\lib\site-packages\pycuda\elementwise.py", line 161, in get_elwise_kernel
arguments, operation, name, keep, options, **kwargs)
File "C:\ProgramData\Anaconda3\lib\site-packages\pycuda\elementwise.py", line 147, in get_elwise_kernel_and_types
keep, options, **kwargs)
File "C:\ProgramData\Anaconda3\lib\site-packages\pycuda\elementwise.py", line 75, in get_elwise_module
options=options, keep=keep)
File "C:\ProgramData\Anaconda3\lib\site-packages\pycuda\compiler.py", line 291, in __init__
arch, code, cache_dir, include_dirs)
File "C:\ProgramData\Anaconda3\lib\site-packages\pycuda\compiler.py", line 254, in compile
return compile_plain(source, options, keep, nvcc, cache_dir, target)
File "C:\ProgramData\Anaconda3\lib\site-packages\pycuda\compiler.py", line 78, in compile_plain
checksum.update(preprocess_source(source, options, nvcc).encode("utf-8"))
File "C:\ProgramData\Anaconda3\lib\site-packages\pycuda\compiler.py", line 55, in preprocess_source
cmdline, stderr=stderr)
pycuda.driver.CompileError: nvcc preprocessing of C:\Users\DANGUR~1\AppData\Local\Temp\tmphnd7bq1w.cu failed
[command: nvcc --preprocess -arch sm_75 -m64 -Ic:\programdata\anaconda3\lib\site-packages\pycuda\cuda C:\Users\DANGUR~1\AppData\Local\Temp\tmphnd7bq1w.cu --compiler-options -EP]
2 months, 1 week
[Pycuda] LogicError: cuMemcpyDtoH failed: an illegal memory access was encountered
by Jie Liu
Hallo,
I have a Pycuda code, which deals with two kernels. Both kernels run well
separately, but when I put them together, there is a memory problem
"LogicError: cuMemcpyDtoH failed: an illegal memory access was encountered".
In the second kernel "DotKernel", I can't change the values of any shared
array or global array. Could you please have a look at the code? Thank you
very much!
Best regards,
Morvan
2 months, 1 week
pycuda works only from the terminal
by Artur Makhmutov
Hello,
I am not sure if this is the right place to ask for a tech support,
please ignore the message if it is not.
My problem is described as follows: I am trying to run one of the
example scripts (I use ubuntu, pycharm with venv python virtual
environment):
> |import pycuda.autoinit from pycuda.compiler import SourceModule
> import numpy a = numpy.random.randn(4,4) a = a.astype(numpy.float32)
> a_gpu = cuda.mem_alloc(a.size * a.dtype.itemsize)
> cuda.memcpy_htod(a_gpu, a) mod = SourceModule(""" __global__ void
> doublify(float *a) { int idx = threadIdx.x + threadIdx.y*4; a[idx] *=
> 2; } """) func = mod.get_function("doublify") func(a_gpu,
> block=(4,4,1)) a_doubled = numpy.empty_like(a)
> cuda.memcpy_dtoh(a_doubled, a_gpu) print ("original array:") print (a)
> print ("doubled with kernel:") print (a_doubled)|
And I get the following error:
> |/home/art/.pyenv/versions/test380/bin/python
> /home/art/PycharmProjects/pycuda/test.py Traceback (most recent call
> last): File
> "/home/art/.pyenv/versions/test380/lib/python3.8/site-packages/pytools/__init__.py",
> line 536, in _deco return func._memoize_dic[args] # pylint:
> disable=protected-access AttributeError: 'function' object has no
> attribute '_memoize_dic' During handling of the above exception,
> another exception occurred: Traceback (most recent call last): File
> "/home/art/.pyenv/versions/test380/lib/python3.8/site-packages/pytools/prefork.py",
> line 49, in call_capture_output popen = Popen(cmdline, cwd=cwd,
> stdin=PIPE, stdout=PIPE, File
> "/home/art/.pyenv/versions/3.8.0/lib/python3.8/subprocess.py", line
> 854, in __init__ self._execute_child(args, executable, preexec_fn,
> close_fds, File
> "/home/art/.pyenv/versions/3.8.0/lib/python3.8/subprocess.py", line
> 1702, in _execute_child raise child_exception_type(errno_num, err_msg,
> err_filename) FileNotFoundError: [Errno 2] No such file or directory:
> 'nvcc' During handling of the above exception, another exception
> occurred: Traceback (most recent call last): File
> "/home/art/PycharmProjects/pycuda/test.py", line 16, in <module> mod =
> SourceModule(""" File
> "/home/art/.pyenv/versions/test380/lib/python3.8/site-packages/pycuda/compiler.py",
> line 290, in __init__ cubin = compile(source, nvcc, options, keep,
> no_extern_c, File
> "/home/art/.pyenv/versions/test380/lib/python3.8/site-packages/pycuda/compiler.py",
> line 254, in compile return compile_plain(source, options, keep, nvcc,
> cache_dir, target) File
> "/home/art/.pyenv/versions/test380/lib/python3.8/site-packages/pycuda/compiler.py",
> line 84, in compile_plain
> checksum.update(get_nvcc_version(nvcc).encode("utf-8")) File
> "</home/art/.pyenv/versions/test380/lib/python3.8/site-packages/decorator.py:decorator-gen-3>",
> line 2, in get_nvcc_version File
> "/home/art/.pyenv/versions/test380/lib/python3.8/site-packages/pytools/__init__.py",
> line 539, in _deco result = func(*args) File
> "/home/art/.pyenv/versions/test380/lib/python3.8/site-packages/pycuda/compiler.py",
> line 16, in get_nvcc_version result, stdout, stderr =
> call_capture_output(cmdline) File
> "/home/art/.pyenv/versions/test380/lib/python3.8/site-packages/pytools/prefork.py",
> line 227, in call_capture_output return
> forker.call_capture_output(cmdline, cwd, error_on_nonzero) File
> "/home/art/.pyenv/versions/test380/lib/python3.8/site-packages/pytools/prefork.py",
> line 60, in call_capture_output raise ExecError("error invoking '%s':
> %s" pytools.prefork.ExecError: error invoking 'nvcc --version': [Errno
> 2] No such file or directory: 'nvcc' Process finished with exit code 1|
But if I run it with the 'python test.py' from the terminal itself,
then it executes just fine:
> |(test380) art@HP:~/PycharmProjects/pycuda$ python test.py original
> array: [[ 0.63286567 -0.5732655 0.0824481 0.17147891] [-0.24867015
> -2.4119377 -0.41027954 -0.67181575] [-1.2339077 -1.23354 1.0630324
> 0.3807849 ] [-1.5976559 -1.5595584 -0.03161036 -0.50650793]] doubled
> with kernel: [[ 1.2657313 -1.146531 0.1648962 0.34295782] [-0.4973403
> -4.8238754 -0.8205591 -1.3436315 ] [-2.4678154 -2.46708 2.1260648
> 0.7615698 ] [-3.1953118 -3.1191168 -0.06322072 -1.0130159 ]] |
I wonder if it supposed to work like that or am I missing something? I
have checked the .bashrc file, nvcc --version, gcc --version and other
hints from the manual, everything seems to be correct. Would much
appreciate your help.
Kind Regards,
Artur
2 months, 2 weeks
Field Support Engineers jobs
by fieldengineer77＠gmail.com
Field Support Engineers provide on-site support for client networks. The engineer must possess knowledge of network hardware, networking, software, and configuration. They address tasks such as the testing, documenting, setting up, and installation of network hardware. They actively participate with clients in devising solutions to improve services. The engineer must be conversant with the latest technologies. Engineers play a vital role in the success of any organization.
The Field Support Engineer is responsible to create emergency service plans for customers. They must recommend effective solutions and improvements to achieve the targeted goals. It is them who usually take care of the scheduled customer and technical support related trainings to field staffs. The professionals must look into product installations and anything related to maintenance activities.
More info: https://www.fieldengineer.com/skills/field-support-engineer
2 months, 2 weeks