summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--libcuda/cuda_runtime_api.cc74
-rw-r--r--src/cuda-sim/ptx_ir.h12
-rw-r--r--src/cuda-sim/ptx_parser.cc34
-rw-r--r--src/gpgpu-sim/shader.cc6
4 files changed, 62 insertions, 64 deletions
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index c87c6c3..c12aaeb 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -3033,7 +3033,6 @@ CUresult CUDAAPI cuModuleGetFunction(CUfunction *hfunc, CUmodule hmod, const cha
function_info* f = symtab->lookup_function( std::string(name) );
//just need to add given pointer to map for cudaLaunch
context->register_hostFun_function( (const char*) hfunc, f);
- g_cuda_launch_stack.push_back( kernel_config() );
*hfunc = (CUfunction)hfunc;
return CUDA_SUCCESS;
}
@@ -3047,60 +3046,6 @@ CUresult CUDAAPI cuModuleUnload(CUmodule hmod)
return CUDA_SUCCESS;
}
-CUresult CUDAAPI cuFuncSetBlockShape(CUfunction hfunc, int x, int y, int z)
-{
- if(g_debug_execution >= 3){
- announce_call(__my_func__);
- }
- gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" );
- kernel_config &config = g_cuda_launch_stack.back();
- dim3 *d = new dim3(x,y,z);
- config.set_block_dim(d);
-
- return CUDA_SUCCESS;
-}
-
-CUresult CUDAAPI cuParamSetSize(CUfunction hfunc, unsigned int numbytes)
-{
- if(g_debug_execution >= 3){
- announce_call(__my_func__);
- }
- //check if size matches given args
- gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" );
- kernel_config &config = g_cuda_launch_stack.back();
- gpgpu_ptx_sim_arg_list_t args = config.get_args();
- size_t total_size = 0;
- for( gpgpu_ptx_sim_arg_list_t::iterator a = args.begin(); a != args.end(); a++ ) {
- total_size += a->m_nbytes;
- }
- return (numbytes==total_size) ? CUDA_SUCCESS : CUDA_ERROR_INVALID_VALUE;
-}
-
-CUresult CUDAAPI cuParamSetv(CUfunction hfunc, int offset, void *ptr, unsigned int numbytes)
-{
- if(g_debug_execution >= 3){
- announce_call(__my_func__);
- }
- cudaSetupArgument((const void *) ptr, (size_t) numbytes, (size_t) offset);
- return CUDA_SUCCESS;
-}
-
-CUresult CUDAAPI cuLaunchGrid(CUfunction f, int grid_width, int grid_height)
-{
- if(g_debug_execution >= 3){
- announce_call(__my_func__);
- }
- const char *hostFun = (const char*) f;
- gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" );
- kernel_config &config = g_cuda_launch_stack.back();
- dim3 *d = new dim3(grid_width,grid_height,1);
- config.set_grid_dim(d);
-
- cudaLaunch(hostFun);
- return CUDA_SUCCESS;
-}
-
-
CUresult CUDAAPI cuLaunchKernel(CUfunction f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ,
unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ,
unsigned int sharedMemBytes, CUstream hStream, void **kernelParams, void **extra)
@@ -3108,17 +3053,18 @@ CUresult CUDAAPI cuLaunchKernel(CUfunction f, unsigned int gridDimX, unsigned in
if(g_debug_execution >= 3){
announce_call(__my_func__);
}
- if (sharedMemBytes!=0||hStream!=NULL||kernelParams!=NULL||extra!=NULL){
- printf("GPGPU-Sim CUDA DRIVER API: Warning: Currently do not support \nsharedMemBytes, hStream, kernelParams, and extra parameters.\n");
+ if (extra!=NULL){
+ printf("GPGPU-Sim CUDA DRIVER API: ERROR: Currently do not support void** extra.\n");
+ abort();
}
const char *hostFun = (const char*) f;
- gpgpusim_ptx_assert( !g_cuda_launch_stack.empty(), "empty launch stack" );
- kernel_config &config = g_cuda_launch_stack.back();
- dim3 *d_b = new dim3(blockDimX, blockDimY, blockDimZ);
- config.set_block_dim(d_b);
- dim3 *d_g = new dim3(gridDimX, gridDimY, gridDimZ);
- config.set_grid_dim(d_g);
-
+ CUctx_st *context = GPGPUSim_Context();
+ function_info *entry = context->get_kernel(hostFun);
+ cudaConfigureCall(dim3(gridDimX, gridDimY, gridDimZ), dim3(blockDimX, blockDimY, blockDimZ), sharedMemBytes, (cudaStream_t) hStream);
+ for(unsigned i = 0; i < entry->num_args(); i++){
+ std::pair<size_t, unsigned> p = entry->get_param_config(i);
+ cudaSetupArgument(kernelParams[i], p.first, p.second);
+ }
cudaLaunch(hostFun);
return CUDA_SUCCESS;
}
diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h
index ef4cf48..e726ab9 100644
--- a/src/cuda-sim/ptx_ir.h
+++ b/src/cuda-sim/ptx_ir.h
@@ -1293,6 +1293,17 @@ public:
bool is_pdom_set() const { return pdom_done; } //return pdom flag
void set_pdom() { pdom_done = true; } //set pdom flag
+ void add_config_param( size_t size, unsigned alignment ){
+ unsigned offset = 0;
+ if (m_param_configs.size()>0){
+ unsigned offset_nom = m_param_configs.back().first + m_param_configs.back().second;
+ offset = offset_nom%alignment ? (offset_nom/alignment + 1) * alignment : offset_nom;
+ }
+ m_param_configs.push_back(std::pair<size_t,unsigned>(size, offset));
+ }
+
+ std::pair<size_t, unsigned> get_param_config(unsigned param_num) const { return m_param_configs[param_num]; }
+
private:
unsigned m_uid;
unsigned m_local_mem_framesize;
@@ -1306,6 +1317,7 @@ private:
unsigned m_instr_mem_size;
std::map<std::string,param_t> m_kernel_params;
std::map<unsigned,param_info> m_ptx_kernel_param_info;
+ std::vector< std::pair<size_t, unsigned> > m_param_configs;
const symbol *m_return_var_sym;
std::vector<const symbol*> m_args;
std::list<ptx_instruction*> m_instructions;
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index c418fac..e6d6325 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -61,6 +61,7 @@ memory_space_t g_ptr_spec = undefined_space;
int g_scalar_type_spec = -1;
int g_vector_spec = -1;
int g_alignment_spec = -1;
+int g_size = -1;
int g_extern_spec = 0;
// variable declaration stuff:
@@ -116,6 +117,7 @@ void init_directive_state()
g_vector_spec=-1;
g_opcode=-1;
g_alignment_spec = -1;
+ g_size = -1;
g_extern_spec = 0;
g_scalar_type.clear();
g_operands.clear();
@@ -373,6 +375,9 @@ int pad_address (new_addr_type address, unsigned size, unsigned maxalign) {
void add_identifier( const char *identifier, int array_dim, unsigned array_ident )
{
+ if(array_ident==ARRAY_IDENTIFIER){
+ g_size *= array_dim;
+ }
if( g_func_decl && (g_func_info == NULL) ) {
// return variable decl...
assert( g_add_identifier_cached__identifier == NULL );
@@ -562,10 +567,13 @@ void add_constptr(const char* identifier1, const char* identifier2, int offset)
void add_function_arg()
{
+ assert(g_size>0);
if( g_func_info ) {
PTX_PARSE_DPRINTF("add_function_arg \"%s\"", g_last_symbol->name().c_str() );
g_func_info->add_arg(g_last_symbol);
+ g_func_info->add_config_param( g_size, g_alignment_spec );
}
+
}
void add_extern_spec()
@@ -617,6 +625,32 @@ void add_vector_spec(int spec )
void add_scalar_type_spec( int type_spec )
{
+ //save size of parameter
+ switch ( type_spec ) {
+ case B8_TYPE:
+ case S8_TYPE:
+ case U8_TYPE:
+ g_size = 1; break;
+ case B16_TYPE:
+ case S16_TYPE:
+ case U16_TYPE:
+ case F16_TYPE:
+ g_size = 2; break;
+ case B32_TYPE:
+ case S32_TYPE:
+ case U32_TYPE:
+ case F32_TYPE:
+ g_size = 4; break;
+ case B64_TYPE:
+ case BB64_TYPE:
+ case S64_TYPE:
+ case U64_TYPE:
+ case F64_TYPE:
+ case FF64_TYPE:
+ g_size = 8; break;
+ case BB128_TYPE:
+ g_size = 16; break;
+ }
PTX_PARSE_DPRINTF("add_scalar_type_spec \"%s\"", g_ptx_token_decode[type_spec].c_str());
g_scalar_type.push_back( type_spec );
if ( g_scalar_type.size() > 1 ) {
diff --git a/src/gpgpu-sim/shader.cc b/src/gpgpu-sim/shader.cc
index 4640d65..da85bae 100644
--- a/src/gpgpu-sim/shader.cc
+++ b/src/gpgpu-sim/shader.cc
@@ -3034,6 +3034,12 @@ bool opndcoll_rfu_t::writeback( const warp_inst_t &inst )
for( r=regs.begin(); r!=regs.end();r++,n++ ) {
unsigned reg = *r;
unsigned bank = register_bank(reg,inst.warp_id(),m_num_banks,m_bank_warp_shift);
+ unsigned count = 0;
+ while( !m_arbiter.bank_idle(bank) ) {
+ assert((++count)<m_num_banks);
+ bank++;
+ bank %= m_num_banks;
+ }
if( m_arbiter.bank_idle(bank) ) {
m_arbiter.allocate_bank_for_write(bank,op_t(&inst,reg,m_num_banks,m_bank_warp_shift));
} else {
generated by cgit v1.3 (git 2.53.0) at 2026-07-11 18:49:23 +0000