From e9bc2b7b5431194b5cb2c884fa530d39e51128f6 Mon Sep 17 00:00:00 2001 From: Evghenii Date: Mon, 11 Nov 2013 09:18:15 +0100 Subject: [PATCH] added uniform_new/uniform_delete in util_ptx.m4 and __shfl intrinsics --- Makefile | 6 +- builtins/target-nvptx64.ll | 5 ++ builtins/util_ptx.m4 | 112 +++++++++++++----------------- examples_cuda/deferred/kernels.cu | 25 +++++-- 4 files changed, 73 insertions(+), 75 deletions(-) diff --git a/Makefile b/Makefile index 3977fb4e..65f9c265 100644 --- a/Makefile +++ b/Makefile @@ -250,15 +250,15 @@ objs/lex.o: objs/lex.cpp $(HEADERS) objs/parse.cc @echo Compiling $< $(CXX) $(CXXFLAGS) -o $@ -c $< -objs/builtins-dispatch.cpp: builtins/dispatch.ll builtins/util.m4 builtins/svml.m4 $(wildcard builtins/*common.ll) +objs/builtins-dispatch.cpp: builtins/dispatch.ll builtins/util.m4 builtins/util_ptx.m4 builtins/svml.m4 $(wildcard builtins/*common.ll) @echo Creating C++ source from builtins definition file $< m4 -Ibuiltins/ -DLLVM_VERSION=$(LLVM_VERSION) -DBUILD_OS=UNIX $< | python bitcode2cpp.py $< > $@ -objs/builtins-%-32bit.cpp: builtins/%.ll builtins/util.m4 builtins/svml.m4 $(wildcard builtins/*common.ll) +objs/builtins-%-32bit.cpp: builtins/%.ll builtins/util.m4 builtins/util_ptx.m4 builtins/svml.m4 $(wildcard builtins/*common.ll) @echo Creating C++ source from builtins definition file $< \(32 bit version\) m4 -Ibuiltins/ -DLLVM_VERSION=$(LLVM_VERSION) -DBUILD_OS=UNIX -DRUNTIME=32 $< | python bitcode2cpp.py $< 32bit > $@ -objs/builtins-%-64bit.cpp: builtins/%.ll builtins/util.m4 builtins/svml.m4 $(wildcard builtins/*common.ll) +objs/builtins-%-64bit.cpp: builtins/%.ll builtins/util.m4 builtins/util_ptx.m4 builtins/svml.m4 $(wildcard builtins/*common.ll) @echo Creating C++ source from builtins definition file $< \(64 bit version\) m4 -Ibuiltins/ -DLLVM_VERSION=$(LLVM_VERSION) -DBUILD_OS=UNIX -DRUNTIME=64 $< | python bitcode2cpp.py $< 64bit > $@ diff --git a/builtins/target-nvptx64.ll b/builtins/target-nvptx64.ll index 07851279..b0374b46 100644 --- a/builtins/target-nvptx64.ll +++ b/builtins/target-nvptx64.ll @@ -60,6 +60,11 @@ define i32 @__nctaid_z() nounwind readnone alwaysinline %nb = call i32 @llvm.nvvm.read.ptx.sreg.nctaid.z() ret i32 %nb } +define i32 @__shfl_i32(i32, i32) nounwind readnone alwaysinline +{ + %shfl = tail call i32 asm sideeffect "shfl.idx.b32 $0, $1, $2, 0x1f;", "=r,r,r"(i32 %0, i32 %1) nounwind readnone alwaysinline + ret i32 %shfl +} ;;;;;;;;;;;;;; diff --git a/builtins/util_ptx.m4 b/builtins/util_ptx.m4 index 124f896b..c89ffb93 100644 --- a/builtins/util_ptx.m4 +++ b/builtins/util_ptx.m4 @@ -2853,13 +2853,7 @@ ifelse(RUNTIME, `32', ;; - __delete_varying_32rt declare i8* @malloc(i32) -define i32 @posix_memalign(i8**, i32, i32) alwaysinline -{ - %ptr = call i8* @malloc (i32 %2); - store i8* %ptr, i8** %0 - %ret = add i32 0, 0 - ret i32 %ret -} +declare i32 @posix_memalign(i8**, i32, i32) declare void @free(i8 *) define noalias i8 * @__new_uniform_32rt(i64 %size) { @@ -2915,72 +2909,60 @@ RUNTIME, `64', ;; - __delete_varying_64rt declare i8* @malloc(i64) -define i32 @posix_memalign(i8**, i64, i64) alwaysinline -{ - %ptr = call i8* @malloc (i64 %2); - store i8* %ptr, i8** %0 - %ret = add i32 0, 0 - ret i32 %ret -} declare void @free(i8 *) -define noalias i8 * @__new_uniform_64rt(i64 %size) { - %ptr = alloca i8* - %alignment = load i32* @memory_alignment - %alignment64 = sext i32 %alignment to i64 - %call1 = call i32 @posix_memalign(i8** %ptr, i64 %alignment64, i64 %size) - %ptr_val = load i8** %ptr - ret i8* %ptr_val +define noalias i8 * @__new_uniform_64rt(i64 %size) +{ +entry: +;; compute laneIdx = __tid_x() & (__warpsize() - 1) + %call = tail call i32 @__tid_x() + %call1 = tail call i32 @__warpsize() + %sub = add nsw i32 %call1, -1 + %and = and i32 %sub, %call +;; if (laneIdx == 0) + %cmp = icmp eq i32 %and, 0 + br i1 %cmp, label %if.then, label %if.end + +if.then: ; preds = %entry + %call2 = tail call noalias i8* @malloc(i64 %size) #3 + %phitmp = ptrtoint i8* %call2 to i64 + br label %if.end + +if.end: ; preds = %if.then, %entry + %ptr.0 = phi i64 [ %phitmp, %if.then ], [ undef, %entry ] + %val.sroa.0.0.extract.trunc = trunc i64 %ptr.0 to i32 + %call3 = tail call i32 @__shfl_i32(i32 %val.sroa.0.0.extract.trunc, i32 0) + %val.sroa.0.0.insert.ext = zext i32 %call3 to i64 + %val.sroa.0.4.extract.shift = lshr i64 %ptr.0, 32 + %val.sroa.0.4.extract.trunc = trunc i64 %val.sroa.0.4.extract.shift to i32 + %call8 = tail call i32 @__shfl_i32(i32 %val.sroa.0.4.extract.trunc, i32 0) + %val.sroa.0.4.insert.ext = zext i32 %call8 to i64 + %val.sroa.0.4.insert.shift = shl nuw i64 %val.sroa.0.4.insert.ext, 32 + %val.sroa.0.4.insert.insert = or i64 %val.sroa.0.4.insert.shift, %val.sroa.0.0.insert.ext + %0 = inttoptr i64 %val.sroa.0.4.insert.insert to i8* + ret i8* %0 } +define void @__delete_uniform_64rt(i8 * %ptr) +{ +entry: + %call = tail call i32 @__tid_x() + %call1 = tail call i32 @__warpsize() + %sub = add nsw i32 %call1, -1 + %and = and i32 %sub, %call + %cmp = icmp eq i32 %and, 0 + br i1 %cmp, label %if.then, label %if.end -define @__new_varying32_64rt( %size, %mask) { - %ret = alloca - store zeroinitializer, * %ret - %ret64 = bitcast * %ret to i64 * - %alignment = load i32* @memory_alignment - %alignment64 = sext i32 %alignment to i64 +if.then: ; preds = %entry + tail call void @free(i8* %ptr) #3 + br label %if.end - per_lane(WIDTH, %mask, ` - %sz_LANE_ID = extractelement %size, i32 LANE - %sz64_LANE_ID = zext i32 %sz_LANE_ID to i64 - %store_LANE_ID = getelementptr i64 * %ret64, i32 LANE - %ptr_LANE_ID = bitcast i64* %store_LANE_ID to i8** - %call_LANE_ID = call i32 @posix_memalign(i8** %ptr_LANE_ID, i64 %alignment64, i64 %sz64_LANE_ID)') - - %r = load * %ret - ret %r -} - -define @__new_varying64_64rt( %size, %mask) { - %ret = alloca - store zeroinitializer, * %ret - %ret64 = bitcast * %ret to i64 * - %alignment = load i32* @memory_alignment - %alignment64 = sext i32 %alignment to i64 - - per_lane(WIDTH, %mask, ` - %sz64_LANE_ID = extractelement %size, i32 LANE - %store_LANE_ID = getelementptr i64 * %ret64, i32 LANE - %ptr_LANE_ID = bitcast i64* %store_LANE_ID to i8** - %call_LANE_ID = call i32 @posix_memalign(i8** %ptr_LANE_ID, i64 %alignment64, i64 %sz64_LANE_ID)') - - %r = load * %ret - ret %r -} - -define void @__delete_uniform_64rt(i8 * %ptr) { - call void @free(i8 * %ptr) +if.end: ; preds = %if.then, %entry ret void } -define void @__delete_varying_64rt( %ptr, %mask) { - per_lane(WIDTH, %mask, ` - %iptr_LANE_ID = extractelement %ptr, i32 LANE - %ptr_LANE_ID = inttoptr i64 %iptr_LANE_ID to i8 * - call void @free(i8 * %ptr_LANE_ID) - ') - ret void -} +declare @__new_varying32_64rt( %size, %mask); +declare @__new_varying64_64rt( %size, %mask); +declare void @__delete_varying_64rt( %ptr, %mask); ', ` errprint(`RUNTIME should be defined to either 32 or 64 diff --git a/examples_cuda/deferred/kernels.cu b/examples_cuda/deferred/kernels.cu index 4b33852a..fbb6eddc 100644 --- a/examples_cuda/deferred/kernels.cu +++ b/examples_cuda/deferred/kernels.cu @@ -161,6 +161,7 @@ struct Uniform return data[i]; } + __device__ inline T* get_ptr(const int i) {return &data[i]; } __device__ inline void set(const bool active, const int i, T value) { if (active) @@ -219,6 +220,8 @@ static float reduce_max(float value) value = max(value, __shfl_xor(value, 1<= numLights) active = 0; #if 0 - const int idx = tileNumLights + inclusive_scan_warp(active); - const int nactive = reduce_sum(active); -#else const int2 res = warpBinExclusiveScan(active); const int idx = tileNumLights + res.x; const int nactive = res.y; -#endif -// if (active) - tileLightIndices.set(active, idx,lightIndex); + tileLightIndices.set(active, idx,lightIndex); tileNumLights += nactive; +#else + tileNumLights += packed_store_active(active, tileLightIndices.get_ptr(tileNumLights), + lightIndex); #endif }