Fix rendering at insane resolutions

This commit is contained in:
Steven Robertson 2011-10-19 14:17:01 -04:00
parent 3466113d64
commit efc2ac23e2
3 changed files with 6 additions and 4 deletions

View File

@ -11,7 +11,7 @@ void colorclip(float4 *pixbuf, float gamma, float vibrancy, float highpow,
float linrange, float lingam, float3 bkgd) { float linrange, float lingam, float3 bkgd) {
// TODO: test if over an edge of the framebuffer - currently gutters are // TODO: test if over an edge of the framebuffer - currently gutters are
// used and up to 256 pixels are ignored, which breaks when width<256 // used and up to 256 pixels are ignored, which breaks when width<256
int i = blockDim.x * blockIdx.x + threadIdx.x; int i = (gridDim.x * blockIdx.y + blockIdx.x) * blockDim.x + threadIdx.x;
float4 pix = pixbuf[i]; float4 pix = pixbuf[i];
if (pix.w <= 0) { if (pix.w <= 0) {

View File

@ -103,7 +103,7 @@ uint32_t trunca(float f) {
__global__ __global__
void zero_dptr(float* dptr, int size) { void zero_dptr(float* dptr, int size) {
int i = blockIdx.x * blockDim.x + threadIdx.x; int i = (gridDim.x * blockIdx.y + blockIdx.x) * blockDim.x + threadIdx.x;
if (i < size) { if (i < size) {
dptr[i] = 0.0f; dptr[i] = 0.0f;
} }
@ -161,8 +161,9 @@ void write_half(float &xy, float x, float y, float den) {
number of 4-byte words in the pointer. number of 4-byte words in the pointer.
""" """
zero = mod.get_function("zero_dptr") zero = mod.get_function("zero_dptr")
blocks = int(np.ceil(np.sqrt(size / 1024 + 1)))
zero(dptr, np.int32(size), stream=stream, zero(dptr, np.int32(size), stream=stream,
block=(1024, 1, 1), grid=(size/1024+1, 1)) block=(1024, 1, 1), grid=(blocks, blocks, 1))
class DataPackerView(object): class DataPackerView(object):
""" """

View File

@ -337,8 +337,9 @@ class Animation(object):
bkgd = vec.make_float3(*(bkgd / n)) bkgd = vec.make_float3(*(bkgd / n))
color_fun = self.mod.get_function("colorclip") color_fun = self.mod.get_function("colorclip")
blocks = int(np.ceil(np.sqrt(nbins / 256)))
color_fun(d_out, gam, vib, hipow, lin, lingam, bkgd, color_fun(d_out, gam, vib, hipow, lin, lingam, bkgd,
block=(256, 1, 1), grid=(nbins / 256, 1), block=(256, 1, 1), grid=(blocks, blocks),
stream=filt_stream) stream=filt_stream)
cuda.memcpy_dtoh_async(h_out, d_out, filt_stream) cuda.memcpy_dtoh_async(h_out, d_out, filt_stream)