summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--debug_tools/WatchYourStep/ptxjitplus/Makefile27
-rw-r--r--debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp12
-rw-r--r--src/cuda-sim/cuda-sim.cc3
-rw-r--r--src/cuda-sim/instructions.cc29
-rw-r--r--src/cuda-sim/ptx_ir.h1
-rw-r--r--src/cuda-sim/ptx_parser.cc4
6 files changed, 47 insertions, 29 deletions
diff --git a/debug_tools/WatchYourStep/ptxjitplus/Makefile b/debug_tools/WatchYourStep/ptxjitplus/Makefile
index b273ac2..3cca514 100644
--- a/debug_tools/WatchYourStep/ptxjitplus/Makefile
+++ b/debug_tools/WatchYourStep/ptxjitplus/Makefile
@@ -34,7 +34,6 @@
################################################################################
# Location of the CUDA Toolkit
-CUDA_PATH ?= "/home/jonathan/cuda8"
##############################
# start deprecated interface #
@@ -148,7 +147,7 @@ else ifneq ($(TARGET_ARCH),$(HOST_ARCH))
endif
endif
HOST_COMPILER ?= g++
-NVCC := $(CUDA_PATH)/bin/nvcc -ccbin $(HOST_COMPILER)
+NVCC := $(CUDA_INSTALL_PATH)/bin/nvcc -ccbin $(HOST_COMPILER)
# internal flags
NVCCFLAGS := -m${TARGET_SIZE}
@@ -157,7 +156,7 @@ LDFLAGS :=
# build flags
ifeq ($(TARGET_OS),darwin)
- LDFLAGS += -rpath $(CUDA_PATH)/lib
+ LDFLAGS += -rpath $(CUDA_INSTALL_PATH)/lib
CCFLAGS += -arch $(HOST_ARCH)
else ifeq ($(HOST_ARCH)-$(TARGET_ARCH)-$(TARGET_OS),x86_64-armv7l-linux)
LDFLAGS += --dynamic-linker=/lib/ld-linux-armhf.so.3
@@ -241,20 +240,20 @@ else
CUDA_SEARCH_PATH += $(TARGET_FS)/usr/lib
CUDA_SEARCH_PATH += $(TARGET_FS)/usr/lib/arm-linux-gnueabihf
endif
- CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/armv7-linux-gnueabihf/lib/stubs
+ CUDA_SEARCH_PATH += $(CUDA_INSTALL_PATH)/targets/armv7-linux-gnueabihf/lib/stubs
CUDA_SEARCH_PATH += /usr/arm-linux-gnueabihf/lib
else ifeq ($(TARGET_ARCH)-$(TARGET_OS),aarch64-linux)
- CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/aarch64-linux/lib/stubs
+ CUDA_SEARCH_PATH += $(CUDA_INSTALL_PATH)/targets/aarch64-linux/lib/stubs
else ifeq ($(TARGET_ARCH)-$(TARGET_OS),armv7l-android)
- CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/armv7-linux-androideabi/lib/stubs
+ CUDA_SEARCH_PATH += $(CUDA_INSTALL_PATH)/targets/armv7-linux-androideabi/lib/stubs
else ifeq ($(TARGET_ARCH)-$(TARGET_OS),aarch64-android)
- CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/aarch64-linux-androideabi/lib/stubs
+ CUDA_SEARCH_PATH += $(CUDA_INSTALL_PATH)/targets/aarch64-linux-androideabi/lib/stubs
else ifeq ($(TARGET_ARCH)-$(TARGET_OS),armv7l-qnx)
- CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/ARMv7-linux-QNX/lib/stubs
+ CUDA_SEARCH_PATH += $(CUDA_INSTALL_PATH)/targets/ARMv7-linux-QNX/lib/stubs
else ifeq ($(TARGET_ARCH)-$(TARGET_OS),aarch64-qnx)
- CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/aarch64-qnx/lib/stubs
+ CUDA_SEARCH_PATH += $(CUDA_INSTALL_PATH)/targets/aarch64-qnx/lib/stubs
else ifeq ($(TARGET_ARCH)-$(TARGET_OS),ppc64le-linux)
- CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/ppc64le-linux/lib/stubs
+ CUDA_SEARCH_PATH += $(CUDA_INSTALL_PATH)/targets/ppc64le-linux/lib/stubs
endif
else
ifneq ($(UBUNTU),)
@@ -264,21 +263,21 @@ else
endif
ifeq ($(TARGET_ARCH),x86_64)
- CUDA_SEARCH_PATH += $(CUDA_PATH)/lib64/stubs
+ CUDA_SEARCH_PATH += $(CUDA_INSTALL_PATH)/lib64/stubs
endif
ifeq ($(TARGET_ARCH),armv7l)
- CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/armv7-linux-gnueabihf/lib/stubs
+ CUDA_SEARCH_PATH += $(CUDA_INSTALL_PATH)/targets/armv7-linux-gnueabihf/lib/stubs
CUDA_SEARCH_PATH += /usr/lib/arm-linux-gnueabihf
endif
ifeq ($(TARGET_ARCH),aarch64)
CUDA_SEARCH_PATH += /usr/lib
- CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/aarch64-linux/lib/stubs
+ CUDA_SEARCH_PATH += $(CUDA_INSTALL_PATH)/targets/aarch64-linux/lib/stubs
endif
ifeq ($(TARGET_ARCH),ppc64le)
- CUDA_SEARCH_PATH += $(CUDA_PATH)/targets/ppc64le-linux/lib/stubs
+ CUDA_SEARCH_PATH += $(CUDA_INSTALL_PATH)/targets/ppc64le-linux/lib/stubs
CUDA_SEARCH_PATH += /usr/lib/powerpc64le-linux-gnu
endif
endif
diff --git a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
index 1065870..554e831 100644
--- a/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
+++ b/debug_tools/WatchYourStep/ptxjitplus/ptxjitplus.cpp
@@ -61,7 +61,6 @@ const char *sSDKname = "PTX Just In Time (JIT) Compilation Plus";
char *wys_exec_path;
char *wys_exec_name;
char *wys_launch_num;
-bool gpgpusim = false;
dim3 gridDim, blockDim;
std::string kernelName;
@@ -138,10 +137,6 @@ void ptxJIT(int argc, char **argv, CUmodule *phModule, CUfunction *phKernel, CUl
void initializeData(std::vector<param>& v_params)
{
- char *gpgpusim_env = getenv("GPGPUSIM_SETUP_ENVIRONMENT_WAS_RUN");
- if (gpgpusim_env!=NULL&&gpgpusim_env[0] == '1'){
- gpgpusim=true;
- }
wys_exec_path = getenv("WYS_EXEC_PATH");
assert(wys_exec_path!=NULL);
wys_exec_name = getenv("WYS_EXEC_NAME");
@@ -272,23 +267,16 @@ int main(int argc, char **argv)
unsigned char **d_data = (unsigned char **) malloc(sizeof(unsigned char **));
checkCudaErrors(cudaMalloc((void**)d_data, p->size));
checkCudaErrors(cudaMemcpy((void*)*d_data,(void*)p->data,p->size,cudaMemcpyHostToDevice));
- if (gpgpusim){
- checkCudaErrors(cuParamSetv(hKernel, p->offset, d_data, sizeof(*d_data)));
- }
paramKernels[index] = (void*)d_data;
m_device_data[index]=*d_data;
m_cleanup[index]=d_data;
paramOffset = p->offset + 8;
}else{
- if (gpgpusim){
- checkCudaErrors(cuParamSetv(hKernel, p->offset, p->data, p->size));
- }
paramKernels[index] = (void*)p->data;
paramOffset = p->offset + p->size;
}
index ++;
}
- checkCudaErrors(cuParamSetSize(hKernel, paramOffset));
// Launch the kernel (Driver API_)
CUDAAPI cuLaunchKernel(hKernel, gridDim.x, gridDim.y, gridDim.z, blockDim.x, blockDim.y, blockDim.z,
diff --git a/src/cuda-sim/cuda-sim.cc b/src/cuda-sim/cuda-sim.cc
index e0cd3f9..b85ba95 100644
--- a/src/cuda-sim/cuda-sim.cc
+++ b/src/cuda-sim/cuda-sim.cc
@@ -1270,6 +1270,7 @@ void function_info::ptx_jit_config(std::map<unsigned long long, size_t> mallocPt
assert(param.find("<")==std::string::npos);
assert(param.find("*")==std::string::npos);
}
+ printf("%s\n", param.c_str());
if(param.find("*")!=std::string::npos){
paramIsPointer.push_back(true);
}else{
@@ -1288,7 +1289,7 @@ void function_info::ptx_jit_config(std::map<unsigned long long, size_t> mallocPt
param_t param_value = p.get_value();
offsets.push_back((unsigned)p.get_offset());
- if (paramIsPointer[i->first]){
+ if (paramIsPointer[i->first] && (*(unsigned long long*)param_value.pdata != 0)){
//is pointer
assert(param_value.size==sizeof(void*)&&"MisID'd this param as pointer");
size_t array_size = 0;
diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc
index fe36902..31fc434 100644
--- a/src/cuda-sim/instructions.cc
+++ b/src/cuda-sim/instructions.cc
@@ -1459,7 +1459,34 @@ void breakaddr_impl( const ptx_instruction *pI, ptx_thread_info *thread )
assert(pI->has_pred() == false); // pdom analysis cannot handle if this instruction is predicated
}
-void brev_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); }
+void brev_impl( const ptx_instruction *pI, ptx_thread_info *thread )
+{
+ ptx_reg_t src1_data, data;
+ const operand_info &dst = pI->dst();
+ const operand_info &src1 = pI->src1();
+ unsigned i_type = pI->get_type();
+ src1_data = thread->get_operand_value(src1, dst, i_type, thread, 1);
+
+ unsigned msb;
+ switch(i_type){
+ case B32_TYPE:
+ msb = 31;
+ for (unsigned i=0; i<=msb; i++) {
+ if((src1_data.u32 & (1 << i)))
+ data.u32 |= 1 << (msb - i);
+ }
+ break;
+ case B64_TYPE:
+ msb = 63;
+ for (unsigned i=0; i<=msb; i++) {
+ if((src1_data.u64 & (1 << i)))
+ data.u64 |= 1 << (msb - i);
+ }
+ break;
+ default: assert(0);
+ }
+ thread->set_operand_value(dst,data, i_type, thread, pI);
+}
void brkpt_impl( const ptx_instruction *pI, ptx_thread_info *thread ) { inst_not_implemented(pI); }
void call_impl( const ptx_instruction *pI, ptx_thread_info *thread )
diff --git a/src/cuda-sim/ptx_ir.h b/src/cuda-sim/ptx_ir.h
index e726ab9..768b3e7 100644
--- a/src/cuda-sim/ptx_ir.h
+++ b/src/cuda-sim/ptx_ir.h
@@ -1297,6 +1297,7 @@ public:
unsigned offset = 0;
if (m_param_configs.size()>0){
unsigned offset_nom = m_param_configs.back().first + m_param_configs.back().second;
+ //ensure offset matches alignment requirements
offset = offset_nom%alignment ? (offset_nom/alignment + 1) * alignment : offset_nom;
}
m_param_configs.push_back(std::pair<size_t,unsigned>(size, offset));
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index e6d6325..0ba9a5b 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -571,7 +571,9 @@ void add_function_arg()
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 );
+ unsigned alignment = (g_alignment_spec==-1) ? g_size : g_alignment_spec;
+ assert(alignment==1||alignment==2||alignment==4||alignment==8||alignment==16);//known valid alignment values
+ g_func_info->add_config_param( g_size, alignment);
}
}
generated by cgit v1.3 (git 2.53.0) at 2026-07-11 06:17:01 +0000