57 lines
2.7 KiB
Text
57 lines
2.7 KiB
Text
#pragma once
|
|
|
|
// Common rectangle path: use block coordinates directly, avoiding 64-bit tile divisions.
|
|
// One thread keeps its coordinate while scanning all seeds, as in stage 03.
|
|
__global__ void outerRectKernel(DeviceArgs a,show::Hit* hits,unsigned long long* count) {
|
|
__shared__ uint64_t xs[16+show::MAX_N-1],zs[16+show::MAX_N-1];
|
|
const unsigned lane=threadIdx.y*16+threadIdx.x;
|
|
const uint64_t xb=uint64_t(blockIdx.x)*16,zb=uint64_t(blockIdx.y)*16;
|
|
if(lane<unsigned(16+a.n-1)) {
|
|
xs[lane]=show::xTerm(int32_t(int64_t(a.x0)+int64_t(xb+lane)));
|
|
zs[lane]=show::zTerm(int32_t(int64_t(a.z0)+int64_t(zb+lane)));
|
|
}
|
|
__syncthreads();
|
|
const uint64_t xi=xb+threadIdx.x,zi=zb+threadIdx.y;
|
|
if(xi>=a.width || zi>=a.height) return;
|
|
for(uint64_t s=0;s<a.seeds;++s) {
|
|
bool ok=true;
|
|
for(int dx=0;dx<a.n && ok;++dx) for(int dz=0;dz<a.n;++dz) {
|
|
if(!show::original(a.seed+s,xs[threadIdx.x+dx]+zs[threadIdx.y+dz])) {
|
|
ok=false;break;
|
|
}
|
|
}
|
|
if(ok) emit(a,hits,count,a.seed+s,int32_t(int64_t(a.x0)+int64_t(xi)),
|
|
int32_t(int64_t(a.z0)+int64_t(zi)));
|
|
}
|
|
}
|
|
|
|
// One block handles a 16x16 anchor tile; each axis needs 16+n-1 terms.
|
|
// Fallback for unusually tall rectangles exceeding the grid.y limit.
|
|
__global__ void outerKernel(DeviceArgs a, show::Hit* hits, unsigned long long* count) {
|
|
__shared__ uint64_t xs[16+show::MAX_N-1], zs[16+show::MAX_N-1];
|
|
const uint64_t tilesX=(a.width+15)/16, tilesZ=(a.height+15)/16;
|
|
const uint64_t total=tilesX*tilesZ;
|
|
const unsigned lane=threadIdx.y*16+threadIdx.x;
|
|
for(uint64_t tile=blockIdx.x;tile<total;tile+=gridDim.x) {
|
|
const uint64_t tx=tile/tilesZ, tz=tile%tilesZ;
|
|
if(lane<unsigned(16+a.n-1)) {
|
|
xs[lane]=show::xTerm(int32_t(int64_t(a.x0)+int64_t(tx*16+lane)));
|
|
zs[lane]=show::zTerm(int32_t(int64_t(a.z0)+int64_t(tz*16+lane)));
|
|
}
|
|
__syncthreads();
|
|
const uint64_t xi=tx*16+threadIdx.x, zi=tz*16+threadIdx.y;
|
|
for(uint64_t s=0;s<a.seeds;++s) {
|
|
bool ok=xi<a.width && zi<a.height;
|
|
for(int dx=0;dx<a.n && ok;++dx) for(int dz=0;dz<a.n;++dz) {
|
|
if(!show::original(a.seed+s,xs[threadIdx.x+dx]+zs[threadIdx.y+dz])) {
|
|
ok=false; break;
|
|
}
|
|
}
|
|
if(ok) emit(a,hits,count,a.seed+s,int32_t(int64_t(a.x0)+int64_t(xi)),
|
|
int32_t(int64_t(a.z0)+int64_t(zi)));
|
|
}
|
|
// No final barrier is needed for a block assigned only one tile.
|
|
// When a giant workload exceeds the grid limit, synchronize before reuse.
|
|
if(tile+gridDim.x<total) __syncthreads();
|
|
}
|
|
}
|