summaryrefslogtreecommitdiff
path: root/libcuda
diff options
context:
space:
mode:
Diffstat (limited to 'libcuda')
-rw-r--r--libcuda/cuda_runtime_api.cc88
-rw-r--r--libcuda/cuobjdump.l53
-rw-r--r--libcuda/cuobjdump.y12
3 files changed, 109 insertions, 44 deletions
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 27644b3..11dcd5a 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -354,8 +354,8 @@ struct _cuda_device_id *GPGPUSim_Init()
cudaDeviceProp *prop = (cudaDeviceProp *) calloc(sizeof(cudaDeviceProp),1);
snprintf(prop->name,256,"GPGPU-Sim_v%s", g_gpgpusim_version_string );
- prop->major = 5;
- prop->minor = 2;
+ prop->major = the_gpu->compute_capability_major();
+ prop->minor = the_gpu->compute_capability_minor();
prop->totalGlobalMem = 0x80000000 /* 2 GB */;
prop->memPitch = 0;
if(prop->major >= 2) {
@@ -1120,11 +1120,11 @@ __host__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDevic
case 41:
*value= 0;
break;
- case 75:
- *value= 9 ;
+ case 75://cudaDevAttrComputeCapabilityMajor
+ *value= prop->major ;
break;
- case 76:
- *value= 3 ;
+ case 76://cudaDevAttrComputeCapabilityMinor
+ *value= prop->minor ;
break;
case 78:
*value= 0 ; //TODO: as of now, we dont support stream priorities.
@@ -1199,17 +1199,50 @@ __host__ cudaError_t CUDARTAPI cudaGetDevice(int *device)
*device = g_active_device;
return g_last_cudaError = cudaSuccess;
}
+
__host__ cudaError_t CUDARTAPI cudaDeviceGetLimit ( size_t* pValue, cudaLimit limit )
{
if(g_debug_execution >= 3){
announce_call(__my_func__);
}
- cuda_not_implemented(__my_func__,__LINE__);
- return g_last_cudaError = cudaSuccess;
+ _cuda_device_id *dev = GPGPUSim_Init();
+ const struct cudaDeviceProp *prop = dev->get_prop();
+ const gpgpu_sim_config& config=dev->get_gpgpu()->get_config();
+ switch(limit) {
+ case 0: // cudaLimitStackSize
+ *pValue=config.stack_limit();
+ break;
+ case 2: // cudaLimitMallocHeapSize
+ *pValue=config.heap_limit();
+ break;
+#if (CUDART_VERSION > 5050)
+ case 3: // cudaLimitDevRuntimeSyncDepth
+ if(prop->major > 2){
+ *pValue=config.sync_depth_limit();
+ break;
+ }
+ else{
+ printf("ERROR:Limit %s is not supported on this architecture \n",limit);
+ abort();
+ }
+ case 4: // cudaLimitDevRuntimePendingLaunchCount
+ if(prop->major > 2){
+ *pValue=config.pending_launch_count_limit();
+ break;
+ }
+ else{
+ printf("ERROR:Limit %s is not supported on this architecture \n",limit);
+ abort();
+ }
+#endif
+ default:
+ printf("ERROR:Limit %s unimplemented \n",limit);
+ abort();
+ }
+ return g_last_cudaError = cudaSuccess;
}
-
__host__ cudaError_t CUDARTAPI cudaStreamGetPriority ( cudaStream_t hStream, int* priority )
{
if(g_debug_execution >= 3){
@@ -1569,6 +1602,9 @@ __host__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream)
announce_call(__my_func__);
}
#if (CUDART_VERSION >= 3000)
+ //per-stream synchronization required for application using external libraries without explicit synchronization in the code to
+ //avoid the stream_manager from spinning forever to destroy non-empty streams without making any forward progress.
+ stream->synchronize();
g_stream_manager->destroy_stream(stream);
#endif
return g_last_cudaError = cudaSuccess;
@@ -1935,8 +1971,11 @@ void setCuobjdumpsassfilename(const char* filename){
}
(dynamic_cast<cuobjdumpELFSection*>(cuobjdumpSectionList.front()))->setSASSfilename(filename);
}
-extern int cuobjdump_parse();
-extern FILE *cuobjdump_in;
+typedef void * yyscan_t;
+extern int cuobjdump_lex_init(yyscan_t* scanner);
+extern void cuobjdump_set_in (FILE * _in_str ,yyscan_t yyscanner );
+extern int cuobjdump_parse(yyscan_t scanner);
+extern int cuobjdump_lex_destroy(yyscan_t scanner);
//! Return the executable file of the process containing the PTX/SASS code
//!
@@ -2145,9 +2184,14 @@ void extract_code_using_cuobjdump(){
if (parse_output) {
printf("Parsing file %s\n", fname);
+ FILE *cuobjdump_in;
cuobjdump_in = fopen(fname, "r");
- cuobjdump_parse();
+ yyscan_t scanner;
+ cuobjdump_lex_init(&scanner);
+ cuobjdump_set_in(cuobjdump_in, scanner);
+ cuobjdump_parse(scanner);
+ cuobjdump_lex_destroy(scanner);
fclose(cuobjdump_in);
printf("Done parsing!!!\n");
} else {
@@ -2195,8 +2239,13 @@ void extract_code_using_cuobjdump(){
std::cout << "Done" << std::endl;
std::cout << "Trying to parse " << libcodfn.str() << std::endl;
+ FILE *cuobjdump_in;
cuobjdump_in = fopen(libcodfn.str().c_str(), "r");
- cuobjdump_parse();
+ yyscan_t scanner;
+ cuobjdump_lex_init(&scanner);
+ cuobjdump_set_in(cuobjdump_in, scanner);
+ cuobjdump_parse(scanner);
+ cuobjdump_lex_destroy(scanner);
fclose(cuobjdump_in);
std::getline(libsf, line);
}
@@ -2695,14 +2744,14 @@ cudaError_t cudaDeviceReset ( void ) {
return g_last_cudaError = cudaSuccess;
}
cudaError_t CUDARTAPI cudaDeviceSynchronize(void){
- // I don't know what this should do
if(g_debug_execution >= 3){
announce_call(__my_func__);
}
+ //Blocks until the device has completed all preceding requested tasks
+ synchronize();
return g_last_cudaError = cudaSuccess;
}
-
void CUDARTAPI __cudaRegisterFunction(
void **fatCubinHandle,
const char *hostFun,
@@ -3132,6 +3181,8 @@ __host__ cudaError_t CUDARTAPI cudaDeviceSetLimit(enum cudaLimit limit, size_t v
}
return g_last_cudaError = cudaSuccess;
}
+
+
#endif
#endif
@@ -3310,7 +3361,12 @@ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *hostFun,
announce_call(__my_func__);
}
function_info *entry = context->get_kernel(hostFun);
- kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry);
+ gpgpu_t* gpu= context->get_device()->get_gpgpu();
+ /*
+ Passing a snapshot of the GPU's current texture mapping to the kernel's info
+ as kernels should use texture bindings present at the time of their launch.
+ */
+ kernel_info_t *result = new kernel_info_t(gridDim,blockDim,entry,gpu->getNameArrayMapping(),gpu->getNameInfoMapping());
if( entry == NULL ) {
printf("GPGPU-Sim PTX: ERROR launching kernel -- no PTX implementation found for %p\n", hostFun);
abort();
diff --git a/libcuda/cuobjdump.l b/libcuda/cuobjdump.l
index 2b0dac8..9359281 100644
--- a/libcuda/cuobjdump.l
+++ b/libcuda/cuobjdump.l
@@ -38,9 +38,7 @@
#define YYDEBUG 1
-#define yylval cuobjdump_lval
-
-void cuobjdump_error(const char*);
+void cuobjdump_error(yyscan_t yyscanner, const char* msg);
%}
%option stack
@@ -48,6 +46,8 @@ void cuobjdump_error(const char*);
%option yylineno
%option nounput
+%option bison-bridge
+%option reentrant
%s ptxcode
%s sasscode
@@ -69,54 +69,54 @@ newlines {newline}+
"ptxasOptions"{notnewline}*{newline}
-[1-9]{numeric}* yylval.string_value = strdup(yytext); return DECIMAL;
+[1-9]{numeric}* yylval->string_value = strdup(yytext); return DECIMAL;
"has debug info"{newline}
"Fatbin ptx code:"{newline} {
- yy_push_state(ptxcode);
- yy_push_state(identifier);
- yy_push_state(ptxheader);
- yylval.string_value = strdup(yytext);
+ yy_push_state(ptxcode, yyscanner);
+ yy_push_state(identifier, yyscanner);
+ yy_push_state(ptxheader, yyscanner);
+ yylval->string_value = strdup(yytext);
return PTXHEADER;
}
"Fatbin elf code:"{newline} {
- yy_push_state(elfcode);
- yy_push_state(identifier);
- yy_push_state(elfheader);
- yylval.string_value = strdup(yytext);
+ yy_push_state(elfcode, yyscanner);
+ yy_push_state(identifier, yyscanner);
+ yy_push_state(elfheader, yyscanner);
+ yylval->string_value = strdup(yytext);
return ELFHEADER;
}
/*PTX code tokens*/
-<ptxcode>{notnewline}*{newline} yylval.string_value = strdup(yytext); return PTXLINE;
+<ptxcode>{notnewline}*{newline} yylval->string_value = strdup(yytext); return PTXLINE;
/*ELF code tokens*/
<elfcode>{whitespace}*"code for sm_"{numeric}+{newline} {
BEGIN(sasscode);
- yylval.string_value = strdup(yytext);
+ yylval->string_value = strdup(yytext);
return SASSLINE;
}
-<elfcode>{notnewline}*{newline} yylval.string_value = strdup(yytext); return ELFLINE;
+<elfcode>{notnewline}*{newline} yylval->string_value = strdup(yytext); return ELFLINE;
/*SASS code tokens*/
-<sasscode>{notnewline}*{newline} yylval.string_value = strdup(yytext); return SASSLINE;
+<sasscode>{notnewline}*{newline} yylval->string_value = strdup(yytext); return SASSLINE;
<identifier>{newline}"compressed"{newline} BEGIN(conidentifier); return H_COMPRESSED;
<identifier>{newline}"identifier = " BEGIN(endidentifier); return H_IDENTIFIER;
-<identifier>{newline}{newline} yy_pop_state();
+<identifier>{newline}{newline} yy_pop_state(yyscanner);
<conidentifier>"identifier = " BEGIN(endidentifier); return H_IDENTIFIER;
-<conidentifier>{newline} yy_pop_state();
+<conidentifier>{newline} yy_pop_state(yyscanner);
-<endidentifier>{notnewline}+ yylval.string_value = strdup(yytext); yy_pop_state(); return FILENAME;
+<endidentifier>{notnewline}+ yylval->string_value = strdup(yytext); yy_pop_state(yyscanner); return FILENAME;
/*Header tokens*/
-<elfheader>[[:alnum:]_]+ yylval.string_value = strdup(yytext); return IDENTIFIER;
+<elfheader>[[:alnum:]_]+ yylval->string_value = strdup(yytext); return IDENTIFIER;
<elfheader>"================" return H_SEPARATOR;
<elfheader>"arch = " return H_ARCH;
<elfheader>"code version = " return H_CODEVERSION;
@@ -128,7 +128,7 @@ newlines {newline}+
/*Header tokens*/
-<ptxheader>[[:alnum:]_]+ yylval.string_value = strdup(yytext); return IDENTIFIER;
+<ptxheader>[[:alnum:]_]+ yylval->string_value = strdup(yytext); return IDENTIFIER;
<ptxheader>"================" return H_SEPARATOR;
<ptxheader>"arch = " return H_ARCH;
<ptxheader>"code version = " return H_CODEVERSION;
@@ -143,7 +143,7 @@ newlines {newline}+
/* Looking for the identifier (filename) then the header is done */
-<endheader>{notnewline}+ yylval.string_value = strdup(yytext); yy_pop_state(); return IDENTIFIER;
+<endheader>{notnewline}+ yylval->string_value = strdup(yytext); yy_pop_state(yyscanner); return IDENTIFIER;
@@ -153,11 +153,12 @@ newlines {newline}+
<<EOF>> BEGIN(INITIAL);return 0;
/*No other rule matched. Throw an error*/
-. cuobjdump_error("Invalid token");
+. cuobjdump_error(yyscanner, "Invalid token");
%%
-void cuobjdump_error(const char* message)
+void cuobjdump_error(yyscan_t yyscanner, const char* msg)
{
- printf(" %s near \"%s\"",message, yytext);
- printf(" on line %i\n",yylineno);
+ struct yyguts_t * yyg = (struct yyguts_t*)yyscanner;
+ printf(" %s near \"%s\"",msg, yytext);
+ printf(" on line %i\n",yylineno);
}
diff --git a/libcuda/cuobjdump.y b/libcuda/cuobjdump.y
index 31760f7..66cbace 100644
--- a/libcuda/cuobjdump.y
+++ b/libcuda/cuobjdump.y
@@ -29,8 +29,8 @@
%{
#include <stdio.h>
-int yylex(void);
-void yyerror(const char*);
+typedef void * yyscan_t;
+
extern void addCuobjdumpSection(int sectiontype);
void setCuobjdumparch(const char* arch);
void setCuobjdumpidentifier(const char* identifier);
@@ -44,9 +44,17 @@ FILE *elffile;
FILE *sassfile;
char filename [1024];
%}
+%define api.pure full
+%parse-param {yyscan_t scanner}
+%lex-param {yyscan_t scanner}
+
%union {
char* string_value;
}
+%{
+int yylex(YYSTYPE * yylval_param, yyscan_t yyscanner);
+void yyerror(yyscan_t yyscanner, const char* msg);
+%}
%token <string_value> H_SEPARATOR H_ARCH H_CODEVERSION H_PRODUCER H_HOST H_COMPILESIZE H_IDENTIFIER H_UNKNOWN H_COMPRESSED
%token <string_value> CODEVERSION
%token <string_value> STRING
generated by cgit v1.3 (git 2.53.0) at 2026-07-11 07:44:09 +0000