【发布时间】：2014-01-21 13:14:25
【问题描述】：
我在 Tesla C1060 GPU、计算能力 1.3、Ubuntu 12.04 上使用 Cuda 5。在我的内核中，每个线程计算（私有）本地浮点数组locArr[] 的值。
然后，使用locArr[] 计算浮点变量var 的值。
__global__ void gpuKernel
(
float *src, float *out,
size_t memPitchAux, int w
)
{
float locArr[256];
float var=0.0f;
int tidx = blockIdx.x * blockDim.x + threadIdx.x;
int tidy = blockIdx.y * blockDim.y + threadIdx.y;
int idx = tidy * memPitchAux/4 + tidx;
for(int ic=0; ic<256; ic++)
{
locArr[ic] = 0.0f;
}
for(int is=0; is<255; is++)
{
int ic = fabs(src[tidy*w +tidx]);
locArr[ic] += 1.0f;
}
for(int ic=0; ic<255; ic++)
{
var += locArr[ic];
}
out[idx] = var;
}
每个线程没有足够的寄存器，所以locArr[] 溢出到全局内存。执行此内核的时间约为 18 毫秒，但如果我将 locArr[] 乘以零，则时间会下降到 1.1 毫秒。
for(int ic=0; ic<255; ic++)
{
var += locArr[ic] * 0.0f;
}
我不明白为什么，每个线程应该无论如何在全局内存中读取所需的locArr[] 值，然后将其乘以零。时间应该是一样的。相反，就好像线程已经知道没有必要读取数据，因为输出无论如何都是零。
谁能给我解释一下这是怎么回事？
EDIT ：如果我有
for(int ic=0; ic<255; ic++)
{
var += locArr[ic] * locArr2[ic];
}
其中locArr2[] 是零的本地数组（溢出到全局内存），可以在运行时进行优化吗？
编辑 2：我的 makefile
################################################################################
#
# Makefile project only supported on Mac OSX and Linux Platforms)
#
################################################################################
# OS Name (Linux or Darwin)
OSUPPER = $(shell uname -s 2>/dev/null | tr [:lower:] [:upper:])
OSLOWER = $(shell uname -s 2>/dev/null | tr [:upper:] [:lower:])
# Flags to detect 32-bit or 64-bit OS platform
OS_SIZE = $(shell uname -m | sed -e "s/i.86/32/" -e "s/x86_64/64/")
OS_ARCH = $(shell uname -m | sed -e "s/i386/i686/")
# These flags will override any settings
ifeq ($(i386),1)
OS_SIZE = 32
OS_ARCH = i686
endif
ifeq ($(x86_64),1)
OS_SIZE = 64
OS_ARCH = x86_64
endif
# Flags to detect either a Linux system (linux) or Mac OSX (darwin)
DARWIN = $(strip $(findstring DARWIN, $(OSUPPER)))
# Location of the CUDA Toolkit binaries and libraries
CUDA_PATH ?= /usr/local/cuda-5.0
CUDA_INC_PATH ?= $(CUDA_PATH)/include
CUDA_BIN_PATH ?= $(CUDA_PATH)/bin
ifneq ($(DARWIN),)
CUDA_LIB_PATH ?= $(CUDA_PATH)/lib
else
ifeq ($(OS_SIZE),32)
CUDA_LIB_PATH ?= $(CUDA_PATH)/lib
else
CUDA_LIB_PATH ?= $(CUDA_PATH)/lib64
endif
endif
# Common binaries
NVCC ?= $(CUDA_BIN_PATH)/nvcc
GCC ?= g++
# Extra user flags
EXTRA_NVCCFLAGS ?=
EXTRA_LDFLAGS ?=
EXTRA_CCFLAGS ?=
# CUDA code generation flags
GENCODE_SM10 := -gencode arch=compute_10,code=sm_10
GENCODE_SM20 := -gencode arch=compute_20,code=sm_20
GENCODE_SM30 := -gencode arch=compute_30,code=sm_30 -gencode arch=compute_35,code=sm_35
GENCODE_FLAGS := $(GENCODE_SM10) $(GENCODE_SM20) $(GENCODE_SM30)
# OS-specific build flags
ifneq ($(DARWIN),)
LDFLAGS := -Xlinker -rpath $(CUDA_LIB_PATH) -L$(CUDA_LIB_PATH) -lcudart
CCFLAGS := -arch $(OS_ARCH)
else
ifeq ($(OS_SIZE),32)
LDFLAGS := -L$(CUDA_LIB_PATH) -lcudart
CCFLAGS := -m32
else
LDFLAGS := -L$(CUDA_LIB_PATH) -lcudart
CCFLAGS := -m64
endif
endif
# OS-architecture specific flags
ifeq ($(OS_SIZE),32)
NVCCFLAGS := -m32
else
NVCCFLAGS := -m64
endif
# OpenGL specific libraries
ifneq ($(DARWIN),)
# Mac OSX specific libraries and paths to include
LIBPATH_OPENGL := -L../../common/lib/darwin -L/System/Library/Frameworks/OpenGL.framework/Libraries -framework GLUT -lGL -lGLU ../../common/lib/darwin/libGLEW.a
else
# Linux specific libraries and paths to include
LIBPATH_OPENGL := -L../../common/lib/linux/$(OS_ARCH) -L/usr/X11R6/lib -lGL -lGLU -lX11 -lXi -lXmu -lglut -lGLEW -lrt
endif
# Debug build flags
ifeq ($(dbg),1)
CCFLAGS += -g
NVCCFLAGS += -g -G
TARGET := debug
else
TARGET := release
endif
# Common includes and paths for CUDA
INCLUDES := -I$(CUDA_INC_PATH) -I. -I.. -I../../common/inc
LDFLAGS += $(LIBPATH_OPENGL)
# Target rules
all: build
build: stackOverflow
stackOverflow.o: stackOverflow.cu
$(NVCC) $(NVCCFLAGS) $(EXTRA_NVCCFLAGS) $(GENCODE_FLAGS) $(INCLUDES) -o $@ -c $<
stackOverflow: stackOverflow.o
$(GCC) $(CCFLAGS) -o $@ $+ $(LDFLAGS) $(EXTRA_LDFLAGS)
mkdir -p ./bin/$(OSLOWER)/$(TARGET)
cp $@ ./bin/$(OSLOWER)/$(TARGET)
run: build
./stackOverflow
clean:
rm -f stackOverflow.o stackOverflow *.pgm
【问题讨论】：
-
对我来说，这听起来就像您对优化编译器的期望。
-
请注意：您的第二个 for 循环使用了一个名为“is”的计数器，但在其中使用了“ic”：这是一种想要的行为吗？
-
@Andrea：是的，因为我展示的代码只是我真实代码的一部分..
-
你是如何编译你的代码的？
-
与我在上面复制的makefile（编辑2）。