Skip to content

Commit f8709df

Browse files
author
feiy
committed
Supporting "type-locked". A back-port from VkInline.
1 parent 904dff0 commit f8709df

File tree

9 files changed

+89
-45
lines changed

9 files changed

+89
-45
lines changed

‎Context.cpp

Lines changed: 34 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -30,7 +30,9 @@ namespace CUInline
3030
bool query_struct(const char* name_struct, const std::vector<const char*>& name_members, size_t* offsets);
3131
bool calc_optimal_block_size(const std::vector<CapturedDeviceViewable>& arg_map, const char* code_body, int& sizeBlock, unsigned sharedMemBytes = 0);
3232
bool calc_number_blocks(const std::vector<CapturedDeviceViewable>& arg_map, const char* code_body, int sizeBlock, int& numBlocks, unsigned sharedMemBytes = 0);
33+
bool launch_kernel(KernelId_t kid, dim_type gridDim, dim_type blockDim, size_t num_args, const DeviceViewable** args, unsigned sharedMemBytes);
3334
bool launch_kernel(dim_type gridDim, dim_type blockDim, const std::vector<CapturedDeviceViewable>& arg_map, const char* code_body, unsigned sharedMemBytes = 0);
35+
bool launch_kernel(KernelId_t& kid, dim_type gridDim, dim_type blockDim, const std::vector<CapturedDeviceViewable>& arg_map, const char* code_body, unsigned sharedMemBytes = 0);
3436

3537
void add_include_dir(const char* path);
3638
void add_built_in_header(const char* name, const char* content);
@@ -47,7 +49,6 @@ namespace CUInline
4749
KernelId_t _build_kernel(const std::vector<CapturedDeviceViewable>& arg_map, const char* code_body);
4850
int _launch_calc(KernelId_t kid, unsigned sharedMemBytes);
4951
int _persist_calc(KernelId_t kid, int numBlocks, unsigned sharedMemBytes);
50-
bool _launch_kernel(KernelId_t kid, dim_type gridDim, dim_type blockDim, const std::vector<CapturedDeviceViewable>& arg_map, unsigned sharedMemBytes);
5152

5253
static const char* s_libnvrtc_path;
5354

@@ -156,11 +157,13 @@ namespace CUInline
156157
cuCtxSynchronize();
157158
}
158159

159-
Kernel::Kernel(const std::vector<const char*>& param_names, const char* code_body) :
160-
m_param_names(param_names.size()), m_code_body(code_body)
160+
Kernel::Kernel(const std::vector<const char*>& param_names, const char* code_body, bool type_locked) :
161+
m_param_names(param_names.size()), m_code_body(code_body), m_type_locked(type_locked)
161162
{
162163
for (size_t i = 0; i < param_names.size(); i++)
163164
m_param_names[i] = param_names[i];
165+
166+
m_kid = (unsigned)(-1);
164167
}
165168

166169
bool Kernel::calc_optimal_block_size(const DeviceViewable** args, int& sizeBlock, unsigned sharedMemBytes)
@@ -190,13 +193,36 @@ namespace CUInline
190193
bool Kernel::launch(dim_type gridDim, dim_type blockDim, const DeviceViewable** args, unsigned sharedMemBytes)
191194
{
192195
Context& ctx = Context::get_context();
193-
std::vector<CapturedDeviceViewable> arg_map(m_param_names.size());
194-
for (size_t i = 0; i < m_param_names.size(); i++)
196+
if (!m_type_locked)
195197
{
196-
arg_map[i].obj_name = m_param_names[i].c_str();
197-
arg_map[i].obj = args[i];
198+
std::vector<CapturedDeviceViewable> arg_map(m_param_names.size());
199+
for (size_t i = 0; i < m_param_names.size(); i++)
200+
{
201+
arg_map[i].obj_name = m_param_names[i].c_str();
202+
arg_map[i].obj = args[i];
203+
}
204+
return ctx.launch_kernel(gridDim, blockDim, arg_map, m_code_body.c_str(), sharedMemBytes);
205+
}
206+
else
207+
{
208+
std::unique_lock<std::mutex> locker(m_mu_type_lock);
209+
if (m_kid == (unsigned)(-1))
210+
{
211+
std::vector<CapturedDeviceViewable> arg_map(m_param_names.size());
212+
for (size_t i = 0; i < m_param_names.size(); i++)
213+
{
214+
arg_map[i].obj_name = m_param_names[i].c_str();
215+
arg_map[i].obj = args[i];
216+
}
217+
return ctx.launch_kernel(m_kid, gridDim, blockDim, arg_map, m_code_body.c_str(), sharedMemBytes);
218+
}
219+
else
220+
{
221+
locker.unlock();
222+
return ctx.launch_kernel(m_kid, gridDim, blockDim, m_param_names.size(), args, sharedMemBytes);
223+
}
198224
}
199-
return ctx.launch_kernel(gridDim, blockDim, arg_map, m_code_body.c_str(), sharedMemBytes);
225+
200226
}
201227

202228
}

‎Context.h

Lines changed: 7 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2,6 +2,8 @@
22

33
#include <vector>
44
#include <string>
5+
#include <thread>
6+
#include <mutex>
57
#include "DeviceViewable.h"
68

79
namespace CUInline
@@ -34,7 +36,7 @@ namespace CUInline
3436
public:
3537
size_t num_params() const { return m_param_names.size(); }
3638

37-
Kernel(const std::vector<const char*>& param_names, const char* code_body);
39+
Kernel(const std::vector<const char*>& param_names, const char* code_body, bool type_locked = false);
3840
bool calc_optimal_block_size(const DeviceViewable** args, int& sizeBlock, unsigned sharedMemBytes = 0);
3941
bool calc_number_blocks(const DeviceViewable** args, int sizeBlock, int& numBlocks, unsigned sharedMemBytes = 0);
4042
bool launch(dim_type gridDim, dim_type blockDim, const DeviceViewable** args, unsigned sharedMemBytes = 0);
@@ -43,5 +45,9 @@ namespace CUInline
4345
std::vector<std::string> m_param_names;
4446
std::string m_code_body;
4547

48+
bool m_type_locked;
49+
unsigned m_kid;
50+
std::mutex m_mu_type_lock;
51+
4652
};
4753
}

‎internal/impl_context.inl

Lines changed: 20 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -592,21 +592,20 @@ namespace CUInline
592592
return num;
593593
}
594594

595-
bool Context::_launch_kernel(KernelId_t kid, dim_type gridDim, dim_type blockDim, const std::vector<CapturedDeviceViewable>& arg_map, unsigned sharedMemBytes)
595+
bool Context::launch_kernel(KernelId_t kid, dim_type gridDim, dim_type blockDim, size_t num_args, const DeviceViewable** args, unsigned sharedMemBytes)
596596
{
597597
Kernel *kernel;
598598
{
599599
std::shared_lock<std::shared_mutex> lock(m_mutex_kernels);
600600
kernel = m_kernel_cache[kid];
601601
}
602602

603-
size_t num_params = arg_map.size();
604-
std::vector<ViewBuf> argbufs(num_params);
605-
std::vector<void*> converted_args(num_params);
603+
std::vector<ViewBuf> argbufs(num_args);
604+
std::vector<void*> converted_args(num_args);
606605

607-
for (size_t i = 0; i < num_params; i++)
606+
for (size_t i = 0; i < num_args; i++)
608607
{
609-
argbufs[i] = arg_map[i].obj->view();
608+
argbufs[i] = args[i]->view();
610609
converted_args[i] = argbufs[i].data();
611610
}
612611
CUresult res = cuLaunchKernel(kernel->func, gridDim.x, gridDim.y, gridDim.z, blockDim.x, blockDim.y, blockDim.z, sharedMemBytes, 0, converted_args.data(), 0);
@@ -634,7 +633,21 @@ namespace CUInline
634633
{
635634
KernelId_t kid = _build_kernel(arg_map, code_body);
636635
if (kid == (KernelId_t)(-1)) return false;
637-
return _launch_kernel(kid, gridDim, blockDim, arg_map, sharedMemBytes);
636+
size_t num_params = arg_map.size();
637+
std::vector<const DeviceViewable*> args(num_params);
638+
for (size_t i = 0; i < num_params; i++)
639+
args[i] = arg_map[i].obj;
640+
return launch_kernel(kid, gridDim, blockDim, num_params, args.data(), sharedMemBytes);
641+
}
642+
643+
bool Context::launch_kernel(KernelId_t& kid, dim_type gridDim, dim_type blockDim, const std::vector<CapturedDeviceViewable>& arg_map, const char* code_body, unsigned sharedMemBytes)
644+
{
645+
kid = _build_kernel(arg_map, code_body);
646+
size_t num_params = arg_map.size();
647+
std::vector<const DeviceViewable*> args(num_params);
648+
for (size_t i = 0; i < num_params; i++)
649+
args[i] = arg_map[i].obj;
650+
return launch_kernel(kid, gridDim, blockDim, num_params, args.data(), sharedMemBytes);
638651
}
639652

640653
void Context::add_include_dir(const char* path)

‎python/CUDAInline/Context.py

Lines changed: 21 additions & 22 deletions
Original file line numberDiff line numberDiff line change
@@ -30,9 +30,9 @@ def Wait():
3030
native.n_wait()
3131

3232
class Kernel:
33-
def __init__(self, param_names, body):
33+
def __init__(self, param_names, body, type_locked=False):
3434
o_param_names = StrArray(param_names)
35-
self.m_cptr = native.n_kernel_create(o_param_names.m_cptr, body.encode('utf-8'))
35+
self.m_cptr = native.n_kernel_create(o_param_names.m_cptr, body.encode('utf-8'), type_locked)
3636

3737
def __del__(self):
3838
native.n_kernel_destroy(self.m_cptr)
@@ -65,10 +65,18 @@ def launch(self, gridDim, blockDim, args, sharedMemBytes=0):
6565
sharedMemBytes)
6666

6767
class For:
68-
def __init__(self, param_names, name_iter, body):
68+
def __init__(self, param_names, name_iter, body, type_locked=False):
6969
self.m_param_names = StrArray(param_names)
7070
self.m_name_iter = name_iter
7171
self.m_code_body = body
72+
self.m_kernel = Kernel(['begin', 'end', 'func'],
73+
'''
74+
size_t tid = threadIdx.x + blockIdx.x*blockDim.x + begin;
75+
if(tid>=end) return;
76+
func.inner(tid);
77+
''', type_locked)
78+
self.m_type_locked = type_locked
79+
self.m_sizeBlock = -1
7280

7381
def num_params(self):
7482
return m_param_names.size()
@@ -86,29 +94,20 @@ def launch(self, begin, end, args):
8694
dvend = DVUInt64(end)
8795
o_args = ObjArray(args)
8896
func = self.InnerProcedural(self.m_param_names, o_args, self.m_name_iter, self.m_code_body)
89-
kernel = Kernel(['begin', 'end', 'func'],
90-
'''
91-
size_t tid = threadIdx.x + blockIdx.x*blockDim.x + begin;
92-
if(tid>=end) return;
93-
func.inner(tid);
94-
''')
95-
sizeBlock = kernel.calc_optimal_block_size([dvbegin, dvend, func]);
96-
numBlocks = int((end - begin + sizeBlock - 1) / sizeBlock)
97-
kernel.launch(numBlocks, sizeBlock, [dvbegin, dvend, func])
97+
if not self.m_type_locked or self.m_sizeBlock == -1:
98+
self.m_sizeBlock = self.m_kernel.calc_optimal_block_size([dvbegin, dvend, func]);
99+
numBlocks = int((end - begin + self.m_sizeBlock - 1) / self.m_sizeBlock)
100+
self.m_kernel.launch(numBlocks, self.m_sizeBlock, [dvbegin, dvend, func])
98101

99102
def launch_n(self, n, args):
100-
dv_n = DVUInt64(n)
103+
dvbegin = DVUInt64(0)
104+
dvend = DVUInt64(n)
101105
o_args = ObjArray(args)
102106
func = self.InnerProcedural(self.m_param_names, o_args, self.m_name_iter, self.m_code_body)
103-
kernel = Kernel(['n', 'func'],
104-
'''
105-
size_t tid = threadIdx.x + blockIdx.x*blockDim.x;
106-
if(tid>=n) return;
107-
func.inner(tid);
108-
''')
109-
sizeBlock = kernel.calc_optimal_block_size([dv_n, func]);
110-
numBlocks = int((n + sizeBlock - 1) / sizeBlock)
111-
kernel.launch(numBlocks, sizeBlock, [dv_n, func])
107+
if not self.m_type_locked or self.m_sizeBlock == -1:
108+
self.m_sizeBlock = self.m_kernel.calc_optimal_block_size([dvbegin, dvend, func]);
109+
numBlocks = int((n + self.m_sizeBlock - 1) / self.m_sizeBlock)
110+
self.m_kernel.launch(numBlocks, self.m_sizeBlock, [dvbegin, dvend, func])
112111

113112

114113

‎python/CUDAInline/cffi.py

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -3,6 +3,6 @@
33

44
ffi = _cffi_backend.FFI('CUDAInline.cffi',
55
_version = 0x2601,
6-
_types = b'\x00\x00\x27\x0D\x00\x00\x7F\x03\x00\x00\x00\x0F\x00\x00\x2D\x0D\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x30\x0D\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x0F\x0D\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x0F\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x07\x01\x00\x00\x08\x01\x00\x00\x00\x0F\x00\x00\x0F\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x08\x01\x00\x00\x00\x0F\x00\x00\x0F\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x07\x01\x00\x00\x00\x0F\x00\x00\x0F\x0D\x00\x00\x00\x0F\x00\x00\x36\x0D\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x10\x0D\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x41\x0D\x00\x00\x7D\x03\x00\x00\x00\x0F\x00\x00\x41\x0D\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x0E\x01\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x0D\x01\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x07\x01\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x0B\x01\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x08\x01\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x08\x01\x00\x00\x08\x01\x00\x00\x08\x01\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x0C\x01\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x0C\x01\x00\x00\x27\x03\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x0C\x01\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x0C\x01\x00\x00\x7E\x03\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x01\x11\x00\x00\x27\x11\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x01\x11\x00\x00\x0C\x01\x00\x00\x0C\x01\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x27\x11\x00\x00\x00\x0F\x00\x00\x7F\x0D\x00\x00\x27\x11\x00\x00\x00\x0F\x00\x00\x7F\x0D\x00\x00\x27\x11\x00\x00\x27\x11\x00\x00\x00\x0F\x00\x00\x7F\x0D\x00\x00\x27\x11\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x7F\x0D\x00\x00\x08\x01\x00\x00\x00\x0F\x00\x00\x7F\x0D\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x7F\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x7F\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x0C\x01\x00\x00\x0C\x01\x00\x00\x00\x0F\x00\x00\x7F\x0D\x00\x00\x00\x0F\x00\x00\x02\x01\x00\x00\x7F\x03\x00\x00\x00\x01',
7-
_globals = (b'\x00\x00\x63\x23n_add_built_in_header',0,b'\x00\x00\x60\x23n_add_code_block',0,b'\x00\x00\x67\x23n_add_constant_object',0,b'\x00\x00\x60\x23n_add_include_dir',0,b'\x00\x00\x60\x23n_add_inlcude_filename',0,b'\x00\x00\x1E\x23n_cudainline_try_init',0,b'\x00\x00\x3B\x23n_dim3_create',0,b'\x00\x00\x6E\x23n_dim3_destroy',0,b'\x00\x00\x6E\x23n_dv_destroy',0,b'\x00\x00\x00\x23n_dv_name_view_cls',0,b'\x00\x00\x32\x23n_dvbool_create',0,b'\x00\x00\x09\x23n_dvbool_value',0,b'\x00\x00\x47\x23n_dvbuffer_create',0,b'\x00\x00\x4F\x23n_dvbuffer_from_dvs',0,b'\x00\x00\x47\x23n_dvbuffer_range_create',0,b'\x00\x00\x56\x23n_dvbuffer_range_from_dvbuffer',0,b'\x00\x00\x71\x23n_dvbufferlike_from_host',0,b'\x00\x00\x29\x23n_dvbufferlike_size',0,b'\x00\x00\x75\x23n_dvbufferlike_to_host',0,b'\x00\x00\x5B\x23n_dvcombine_create',0,b'\x00\x00\x2C\x23n_dvdouble_create',0,b'\x00\x00\x03\x23n_dvdouble_value',0,b'\x00\x00\x2F\x23n_dvfloat_create',0,b'\x00\x00\x06\x23n_dvfloat_value',0,b'\x00\x00\x32\x23n_dvint16_create',0,b'\x00\x00\x09\x23n_dvint16_value',0,b'\x00\x00\x32\x23n_dvint32_create',0,b'\x00\x00\x09\x23n_dvint32_value',0,b'\x00\x00\x35\x23n_dvint64_create',0,b'\x00\x00\x20\x23n_dvint64_value',0,b'\x00\x00\x32\x23n_dvint8_create',0,b'\x00\x00\x09\x23n_dvint8_value',0,b'\x00\x00\x38\x23n_dvuint16_create',0,b'\x00\x00\x23\x23n_dvuint16_value',0,b'\x00\x00\x38\x23n_dvuint32_create',0,b'\x00\x00\x23\x23n_dvuint32_value',0,b'\x00\x00\x40\x23n_dvuint64_create',0,b'\x00\x00\x29\x23n_dvuint64_value',0,b'\x00\x00\x38\x23n_dvuint8_create',0,b'\x00\x00\x23\x23n_dvuint8_value',0,b'\x00\x00\x0C\x23n_kernel_calc_number_blocks',0,b'\x00\x00\x12\x23n_kernel_calc_optimal_block_size',0,b'\x00\x00\x52\x23n_kernel_create',0,b'\x00\x00\x6E\x23n_kernel_destroy',0,b'\x00\x00\x17\x23n_kernel_launch',0,b'\x00\x00\x09\x23n_kernel_num_params',0,b'\x00\x00\x4B\x23n_pointer_array_create',0,b'\x00\x00\x6E\x23n_pointer_array_destroy',0,b'\x00\x00\x29\x23n_pointer_array_size',0,b'\x00\x00\x60\x23n_set_libnvrtc_path',0,b'\x00\x00\x6B\x23n_set_verbose',0,b'\x00\x00\x26\x23n_size_of',0,b'\x00\x00\x43\x23n_string_array_create',0,b'\x00\x00\x6E\x23n_string_array_destroy',0,b'\x00\x00\x29\x23n_string_array_size',0,b'\x00\x00\x7B\x23n_wait',0),
6+
_types = b'\x00\x00\x27\x0D\x00\x00\x80\x03\x00\x00\x00\x0F\x00\x00\x2D\x0D\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x30\x0D\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x0F\x0D\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x0F\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x07\x01\x00\x00\x08\x01\x00\x00\x00\x0F\x00\x00\x0F\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x08\x01\x00\x00\x00\x0F\x00\x00\x0F\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x07\x01\x00\x00\x00\x0F\x00\x00\x0F\x0D\x00\x00\x00\x0F\x00\x00\x36\x0D\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x10\x0D\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x41\x0D\x00\x00\x7E\x03\x00\x00\x00\x0F\x00\x00\x41\x0D\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x0E\x01\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x0D\x01\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x07\x01\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x0B\x01\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x08\x01\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x08\x01\x00\x00\x08\x01\x00\x00\x08\x01\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x0C\x01\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x0C\x01\x00\x00\x27\x03\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x0C\x01\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x0C\x01\x00\x00\x7F\x03\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x01\x11\x00\x00\x27\x11\x00\x00\x08\x01\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x01\x11\x00\x00\x0C\x01\x00\x00\x0C\x01\x00\x00\x00\x0F\x00\x00\x01\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x27\x11\x00\x00\x00\x0F\x00\x00\x80\x0D\x00\x00\x27\x11\x00\x00\x00\x0F\x00\x00\x80\x0D\x00\x00\x27\x11\x00\x00\x27\x11\x00\x00\x00\x0F\x00\x00\x80\x0D\x00\x00\x27\x11\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x80\x0D\x00\x00\x08\x01\x00\x00\x00\x0F\x00\x00\x80\x0D\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x80\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x00\x0F\x00\x00\x80\x0D\x00\x00\x01\x11\x00\x00\x01\x11\x00\x00\x0C\x01\x00\x00\x0C\x01\x00\x00\x00\x0F\x00\x00\x80\x0D\x00\x00\x00\x0F\x00\x00\x02\x01\x00\x00\x80\x03\x00\x00\x00\x01',
7+
_globals = (b'\x00\x00\x64\x23n_add_built_in_header',0,b'\x00\x00\x61\x23n_add_code_block',0,b'\x00\x00\x68\x23n_add_constant_object',0,b'\x00\x00\x61\x23n_add_include_dir',0,b'\x00\x00\x61\x23n_add_inlcude_filename',0,b'\x00\x00\x1E\x23n_cudainline_try_init',0,b'\x00\x00\x3B\x23n_dim3_create',0,b'\x00\x00\x6F\x23n_dim3_destroy',0,b'\x00\x00\x6F\x23n_dv_destroy',0,b'\x00\x00\x00\x23n_dv_name_view_cls',0,b'\x00\x00\x32\x23n_dvbool_create',0,b'\x00\x00\x09\x23n_dvbool_value',0,b'\x00\x00\x47\x23n_dvbuffer_create',0,b'\x00\x00\x4F\x23n_dvbuffer_from_dvs',0,b'\x00\x00\x47\x23n_dvbuffer_range_create',0,b'\x00\x00\x57\x23n_dvbuffer_range_from_dvbuffer',0,b'\x00\x00\x72\x23n_dvbufferlike_from_host',0,b'\x00\x00\x29\x23n_dvbufferlike_size',0,b'\x00\x00\x76\x23n_dvbufferlike_to_host',0,b'\x00\x00\x5C\x23n_dvcombine_create',0,b'\x00\x00\x2C\x23n_dvdouble_create',0,b'\x00\x00\x03\x23n_dvdouble_value',0,b'\x00\x00\x2F\x23n_dvfloat_create',0,b'\x00\x00\x06\x23n_dvfloat_value',0,b'\x00\x00\x32\x23n_dvint16_create',0,b'\x00\x00\x09\x23n_dvint16_value',0,b'\x00\x00\x32\x23n_dvint32_create',0,b'\x00\x00\x09\x23n_dvint32_value',0,b'\x00\x00\x35\x23n_dvint64_create',0,b'\x00\x00\x20\x23n_dvint64_value',0,b'\x00\x00\x32\x23n_dvint8_create',0,b'\x00\x00\x09\x23n_dvint8_value',0,b'\x00\x00\x38\x23n_dvuint16_create',0,b'\x00\x00\x23\x23n_dvuint16_value',0,b'\x00\x00\x38\x23n_dvuint32_create',0,b'\x00\x00\x23\x23n_dvuint32_value',0,b'\x00\x00\x40\x23n_dvuint64_create',0,b'\x00\x00\x29\x23n_dvuint64_value',0,b'\x00\x00\x38\x23n_dvuint8_create',0,b'\x00\x00\x23\x23n_dvuint8_value',0,b'\x00\x00\x0C\x23n_kernel_calc_number_blocks',0,b'\x00\x00\x12\x23n_kernel_calc_optimal_block_size',0,b'\x00\x00\x52\x23n_kernel_create',0,b'\x00\x00\x6F\x23n_kernel_destroy',0,b'\x00\x00\x17\x23n_kernel_launch',0,b'\x00\x00\x09\x23n_kernel_num_params',0,b'\x00\x00\x4B\x23n_pointer_array_create',0,b'\x00\x00\x6F\x23n_pointer_array_destroy',0,b'\x00\x00\x29\x23n_pointer_array_size',0,b'\x00\x00\x61\x23n_set_libnvrtc_path',0,b'\x00\x00\x6C\x23n_set_verbose',0,b'\x00\x00\x26\x23n_size_of',0,b'\x00\x00\x43\x23n_string_array_create',0,b'\x00\x00\x6F\x23n_string_array_destroy',0,b'\x00\x00\x29\x23n_string_array_size',0,b'\x00\x00\x7C\x23n_wait',0),
88
)

‎python/CUDAInline/cffi_build.py

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -29,7 +29,7 @@
2929
void n_add_constant_object(const char* name, void* cptr);
3030
void n_wait();
3131
32-
void* n_kernel_create(void* ptr_param_list, const char* body);
32+
void* n_kernel_create(void* ptr_param_list, const char* body, unsigned type_locked);
3333
void n_kernel_destroy(void* cptr);
3434
int n_kernel_num_params(void* cptr);
3535
int n_kernel_calc_optimal_block_size(void* ptr_kernel, void* ptr_arg_list, unsigned sharedMemBytes);

‎python/api.h

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -31,7 +31,7 @@ extern "C"
3131
PY_CUDAInline_API void n_add_constant_object(const char* name, void* cptr);
3232
PY_CUDAInline_API void n_wait();
3333

34-
PY_CUDAInline_API void* n_kernel_create(void* ptr_param_list, const char* body);
34+
PY_CUDAInline_API void* n_kernel_create(void* ptr_param_list, const char* body, unsigned type_locked);
3535
PY_CUDAInline_API void n_kernel_destroy(void* cptr);
3636
PY_CUDAInline_API int n_kernel_num_params(void* cptr);
3737
PY_CUDAInline_API int n_kernel_calc_optimal_block_size(void* ptr_kernel, void* ptr_arg_list, unsigned sharedMemBytes);

‎python/api_Context.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -58,14 +58,14 @@ void n_wait()
5858
Wait();
5959
}
6060

61-
void* n_kernel_create(void* ptr_param_list, const char* body)
61+
void* n_kernel_create(void* ptr_param_list, const char* body, unsigned type_locked)
6262
{
6363
StrArray* param_list = (StrArray*)ptr_param_list;
6464
size_t num_params = param_list->size();
6565
std::vector<const char*> params(num_params);
6666
for (size_t i = 0; i < num_params; i++)
6767
params[i] = (*param_list)[i].c_str();
68-
Kernel* cptr = new Kernel(params, body);
68+
Kernel* cptr = new Kernel(params, body, type_locked!=0);
6969
return cptr;
7070
}
7171

‎python/setup.py

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -9,7 +9,7 @@
99

1010
setup(
1111
name = 'CUDAInline',
12-
version = '0.0.5',
12+
version = '0.0.6',
1313
description = 'A CUDA interface for Python',
1414
long_description=long_description,
1515
long_description_content_type='text/markdown',

0 commit comments

Comments
 (0)