Skip to content

Fix cudaDeviceGetAttribute(81) and cudart version for cudaPopConfiguration #173

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 1 commit into
base: dev
Choose a base branch
from
Open

Fix cudaDeviceGetAttribute(81) and cudart version for cudaPopConfiguration #173

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions libcuda/cuda_runtime_api.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1449,7 +1449,7 @@ __host__ cudaError_t CUDARTAPI cudaConfigureCall(dim3 gridDim, dim3 blockDim, si
}


#if CUDART_VERSION >= 10000
#if CUDART_VERSION >= 9020
/*
* CUDA 10 requires a new CUDA kernel launch sequence
* A call to __cudaPushCallConfiguration() preceeds any call to cudaLaunchKernel()
Expand Down Expand Up @@ -1596,7 +1596,7 @@ __host__ cudaError_t CUDARTAPI cudaLaunchKernel ( const char* hostFun, dim3 grid
CUctx_st *context = GPGPUSim_Context();
function_info *entry = context->get_kernel(hostFun);

#if CUDART_VERSION >= 10000
#if CUDART_VERSION >= 9020
assert(g_cudaPushArgsBuffer::g_is_initialized == false);
cudaConfigureCall(g_cudaPushArgsBuffer::g_gridDim, g_cudaPushArgsBuffer::g_blockDim, g_cudaPushArgsBuffer::g_sharedMem, g_cudaPushArgsBuffer::g_stream);
#else
Expand Down
5 changes: 1 addition & 4 deletions src/gpgpu-sim/gpu-sim.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -320,9 +320,6 @@ void shader_core_config::reg_options(class OptionParser * opp)
option_parser_register(opp, "-adaptive_volta_cache_config", OPT_BOOL, &adaptive_volta_cache_config,
"adaptive_volta_cache_config",
"0");
option_parser_register(opp, "-gpgpu_shmem_size", OPT_UINT32, &gpgpu_shmem_sizeDefault,
"Size of shared memory per shader core (default 16kB)",
"16384");
option_parser_register(opp, "-gpgpu_shmem_size_PrefL1", OPT_UINT32, &gpgpu_shmem_sizePrefL1,
"Size of shared memory per shader core (default 16kB)",
"16384");
Expand Down Expand Up @@ -778,7 +775,7 @@ gpgpu_sim::gpgpu_sim( const gpgpu_sim_config &config )

int gpgpu_sim::shared_mem_size() const
{
return m_shader_config->gpgpu_shmem_size;
return m_shader_config->gpgpu_shmem_size * m_shader_config->n_simt_cores_per_cluster;
}

int gpgpu_sim::shared_mem_per_block() const
Expand Down
2 changes: 1 addition & 1 deletion src/gpgpu-sim/shader.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1399,7 +1399,7 @@ struct shader_core_config : public core_config
}

delete[] tokd;

gpgpu_shmem_sizeDefault = gpgpu_shmem_size;
if (n_thread_per_shader > MAX_THREAD_PER_SM) {
printf("GPGPU-Sim uArch: Error ** increase MAX_THREAD_PER_SM in abstract_hardware_model.h from %u to %u\n",
MAX_THREAD_PER_SM, n_thread_per_shader);
Expand Down