Classified TJS-16 spacecraft, companion object emerges alongside earlier TJS-15 satellite
-rough draft
https://spacenews.com/china-launches-classified-tjs-16-spacecraft-companion-object-emerges-alongside-earlier-tjs-15-satellite/
additional object purposed for signal amplification
Feliks Tomasz Konczakowski
similar to:
https://en.wikipedia.org/wiki/Refracting_telescope
https://en.wikipedia.org/wiki/List_of_largest_optical_refracting_telescopes
cm surface resolution
other links:
JUST IN: SPACECOM Prepared to Defend U.S. Assets By Cambrie Eckert https://www.nationaldefensemagazine.org/articles/2023/7/19/advanced-competitor-capabilities-threaten-us-satellites
Russia’s Anti-Satellite Weapons: An Asymmetric Response to U.S. Aerospace Superiority By Jaganath Sankaran https://www.armscontrol.org/act/2022-03/features/russias-anti-satellite-weapons-asymmetric-response-us-aerospace-superiority
___
Beamforming
Satellite group to operate in conjunction in order to maximize resolution
similar to but spaceborne:
Allen Telescope Array Overview https://www.seti.org/ata
BLADE Allen Telescope Array GPU accelerated beamformer backend https://events.gnuradio.org/event/21/contributions/408/attachments/137/316/BLADE%20GNU%20Radio%20Conference.pdf
| Number of antennas | 28 (27 active and 1 spare) |
| Dish size | 25 meters |
| Weight | 230 tons |
| Latitude | 34°04'43.497" North |
| Longitude | 107°37'03.819" West |
| Dish Surface | Aluminum panels accurate up to .5 mm |
| BUS | Steel |
| Frequencies | From 1.0 GHz to 50 GHz |
| Resolution | 0.2 arcseconds to 0.04 arcseconds |
| Reconfigurable Array | North arm is 11 miles long, two other arms are 13 miles long |
| A configuration size | 22.62 miles across |
| B configuration size | 7.08 miles across |
| C configuration size | 2.11 miles across |
| D configuration size | 0.64 miles across |
https://github.com/luigifcruz/blade/blob/aef4bf5230ab64af7b3e22de60e71f4bb52d6aae/benchmarks/pipelines/ata/mode-b/mode_b.cc#L131-L147
#include <memory>
#include <cassert>
#include "blade/base.hh"
#include "blade/logger.hh"
#include "blade/runner.hh"
#include "blade/plan.hh"
#include "blade/pipelines/ata/mode_b.hh"
extern "C" {
#include "mode_b.h"
}
using namespace Blade;
using namespace Blade::Pipelines::ATA;
using TestPipeline = ModeB<BLADE_ATA_MODE_B_OUTPUT_ELEMENT_T>;
static std::unique_ptr<Runner<TestPipeline>> runner;
static Tensor<Device::CPU, F64> dummyJulianDate({1});
static Tensor<Device::CPU, F64> dummyDut1({1});
bool blade_ata_b_initialize(U64 numberOfWorkers) {
if (runner) {
BL_FATAL("Can't initialize because Blade Runner is already initialized.");
BL_CHECK_THROW(Result::ASSERTION_ERROR);
}
dummyJulianDate[0] = (1649366473.0 / 86400) + 2440587.5;
dummyDut1[0] = 0.0;
runner = Runner<TestPipeline>::New(numberOfWorkers, {
.inputShape = ArrayShape({
BLADE_ATA_MODE_B_NANT,
BLADE_ATA_MODE_B_NCHAN,
BLADE_ATA_MODE_B_NTIME,
BLADE_ATA_MODE_B_NPOL,
}),
.preBeamformerChannelizerRate = BLADE_ATA_MODE_B_CHANNELIZER_RATE,
.phasorObservationFrequencyHz = 6500.125*1e6,
.phasorChannelBandwidthHz = 0.5e6,
.phasorTotalBandwidthHz = 1.024e9,
.phasorFrequencyStartIndex = 352,
.phasorReferenceAntennaIndex = 0,
.phasorArrayReferencePosition = {
.LON = BL_DEG_TO_RAD(-121.470733),
.LAT = BL_DEG_TO_RAD(40.815987),
.ALT = 1020.86,
},
.phasorBoresightCoordinate = {
.RA = 0.64169,
.DEC = 1.079896295,
},
.phasorAntennaPositions = {
{-2524041.5388905862, -4123587.965024342, 4147646.4222955606}, // 1c
{-2524068.187873109, -4123558.735413135, 4147656.21282186}, // 1e
{-2524087.2078100787, -4123532.397416349, 4147670.9866770394}, // 1g
{-2524103.384010733, -4123511.111598937, 4147682.4133068994}, // 1h
{-2524056.730228759, -4123515.287949227, 4147706.4850287656}, // 1k
{-2523986.279601761, -4123497.427940991, 4147766.732988923}, // 2a
{-2523970.301363642, -4123515.238502669, 4147758.790023165}, // 2b
{-2523983.5419911123, -4123528.1422073604, 4147737.872218138}, // 2c
{-2523941.5221860334, -4123568.125040547, 4147723.8292249846}, // 2e
{-2524074.096220788, -4123468.5182652213, 4147742.0422435375}, // 2h
{-2524058.6409591637, -4123466.5112451194, 4147753.4513993543}, // 2j
{-2524026.989692545, -4123480.9405167866, 4147758.2356800516}, // 2l
{-2524048.5254066754, -4123468.3463909747, 4147757.835369889}, // 2k
{-2524000.5641107005, -4123498.2984570004, 4147756.815976133}, // 2m
{-2523945.086670364, -4123480.3638816103, 4147808.127865142}, // 3d
{-2523950.6822576034, -4123444.7023326857, 4147839.7474427638}, // 3l
{-2523880.869769226, -4123514.3375464156, 4147813.413426994}, // 4e
{-2523930.3747946257, -4123454.3080821196, 4147842.6449955846}, // 4g
{-2523898.1150373477, -4123456.314794732, 4147860.3045849088}, // 4j
{-2523824.598229116, -4123527.93080514, 4147833.98936114} // 5b
},
.phasorAntennaCalibrations = ArrayTensor<Device::CPU, CF64>({
BLADE_ATA_MODE_B_NANT,
BLADE_ATA_MODE_B_NCHAN * BLADE_ATA_MODE_B_CHANNELIZER_RATE,
1,
BLADE_ATA_MODE_B_NPOL,
}),
.phasorBeamCoordinates = {
{0.63722, 1.07552424},
{0.64169, 1.079896295},
{0.64169, 1.079896295},
{0.64169, 1.079896295},
{0.64169, 1.079896295},
{0.64169, 1.079896295},
{0.64169, 1.079896295},
{0.64169, 1.079896295},
},
.beamformerIncoherentBeam = BLADE_ATA_MODE_B_ENABLE_INCOHERENT_BEAM,
.detectorEnable = BLADE_ATA_MODE_B_DETECTOR_ENABLED,
.detectorIntegrationSize = BLADE_ATA_MODE_B_DETECTOR_INTEGRATION,
.detectorNumberOfOutputPolarizations = BLADE_ATA_MODE_B_DETECTOR_POLS,
});
// Terminate if profiling.
if (Memory::Profiler::IsCapturing()) {
blade_ata_b_terminate();
}
return true;
}
void blade_ata_b_terminate() {
if (!runner) {
BL_FATAL("Can't terminate because Blade Runner isn't initialized.");
BL_CHECK_THROW(Result::ASSERTION_ERROR);
}
runner.reset();
}
U64 blade_ata_b_get_input_size() {
assert(runner);
return runner->getWorker().getInputBuffer().size();
}
U64 blade_ata_b_get_output_size() {
assert(runner);
return runner->getWorker().getOutputBuffer().size();
}
bool blade_ata_b_enqueue(void* input_ptr, void* output_ptr, U64 id) {
assert(runner);
return runner->enqueue([&](auto& worker) {
// Convert C pointers to Blade::Vector.
auto input = ArrayTensor<Device::CPU, CI8>(input_ptr, worker.getInputBuffer().shape());
auto output = ArrayTensor<Device::CPU, BLADE_ATA_MODE_B_OUTPUT_ELEMENT_T>(output_ptr,
worker.getOutputBuffer().shape());
// Transfer input data from CPU memory to the worker.
Plan::TransferIn(worker, dummyJulianDate, dummyDut1, input);
// Compute block.
Plan::Compute(worker);
// Transfer output data from the worker to the CPU memory.
Plan::TransferOut(output, worker.getOutputBuffer(), worker);
return id;
});
}
bool blade_ata_b_dequeue(U64* id) {
assert(runner);
return runner->dequeue(id);
}
https://github.com/luigifcruz/blade/blob/aef4bf5230ab64af7b3e22de60e71f4bb52d6aae/include/blade/runner.hh
#ifndef BLADE_RUNNER_HH
#define BLADE_RUNNER_HH
#include <deque>
#include <vector>
#include <memory>
#include "blade/logger.hh"
#include "blade/module.hh"
#include "blade/pipeline.hh"
#include "blade/macros.hh"
namespace Blade {
template<class Pipeline>
class BLADE_API Runner {
public:
static std::unique_ptr<Runner<Pipeline>> New(const U64& numberOfWorkers,
const typename Pipeline::Config& config,
const BOOL& printET = true) {
return std::make_unique<Runner<Pipeline>>(numberOfWorkers, config, printET);
}
explicit Runner(const U64& numberOfWorkers,
const typename Pipeline::Config& config,
const BOOL& printET = true) {
if (printET) {
BL_LOG_PRINT_ET();
}
BL_INFO("Instantiating new runner.");
if (numberOfWorkers == 0) {
BL_FATAL("Number of worker has to be larger than zero.");
BL_CHECK_THROW(Result::ASSERTION_ERROR);
}
for (U64 i = 0; i < numberOfWorkers; i++) {
BL_DEBUG("Initializing new worker.");
workers.push_back(std::make_unique<Pipeline>(config));
}
}
constexpr Pipeline& getWorker(const U64& index = 0) const {
return *workers[index];
}
constexpr const U64& getHead() const {
return head;
}
constexpr const bool slotAvailable() const {
return jobs.size() != workers.size();
}
constexpr const bool empty() const {
return jobs.size() == 0;
}
constexpr Pipeline& getNextWorker() {
return *workers[head];
}
const Result applyToAllWorkers(const std::function<const Result(Pipeline&)>& modifier,
const bool block = false) {
for (auto& worker : workers) {
BL_CHECK(modifier(*worker));
}
if (block) {
for (auto& worker : workers) {
BL_CHECK(worker->synchronize());
}
}
return Result::SUCCESS;
}
bool enqueue(const std::function<const U64(Pipeline&)>& jobFunc) {
// Return if there are no workers available.
if (jobs.size() == workers.size()) {
return false;
}
try {
jobs.push_back({
.id = jobFunc(*workers[head]),
.worker = workers[head],
});
} catch (const Result& err) {
// Print user friendly error and issue fatal error.
if (err == Result::PLAN_ERROR_ACCUMULATION_COMPLETE) {
BL_FATAL("Can't accumulate block because buffer is full.");
BL_CHECK_THROW(err);
}
if (err == Result::PLAN_ERROR_DESTINATION_NOT_SYNCHRONIZED) {
BL_FATAL("Can't transfer data because destination is not synchronized.");
BL_CHECK_THROW(err);
}
if (err == Result::PLAN_ERROR_NO_ACCUMULATOR) {
BL_FATAL("This mode doesn't support accumulation.");
BL_CHECK_THROW(err);
}
if (err == Result::PLAN_ERROR_NO_SLOT) {
BL_FATAL("No slot available after compute. Data has nowhere to go.")
BL_CHECK_THROW(err);
}
// Ignore if throw was a skip operation.
if (err == Result::PLAN_SKIP_ACCUMULATION_INCOMPLETE ||
err == Result::PLAN_SKIP_COMPUTE_INCOMPLETE ||
err == Result::PLAN_SKIP_USER_INITIATED ||
err == Result::PLAN_SKIP_NO_DEQUEUE ||
err == Result::PLAN_SKIP_NO_SLOT) {
return false;
}
// Ignore if throw originates from exhaustion.
if (err == Result::EXHAUSTED) {
return false;
}
BL_FATAL("Unknown error.");
// Fatal error otherwise.
BL_CHECK_THROW(err);
}
// Bump job queue head index.
head = (head + 1) % workers.size();
return true;
}
bool dequeue(U64* id) {
// Return if there are no jobs.
if (jobs.size() == 0) {
return false;
}
const auto& job = jobs.front();
// Synchronize front if all workers have jobs.
if (jobs.size() == workers.size()) {
job.worker->synchronize();
}
// Return if front isn't synchronized.
if (!job.worker->isSynchronized()) {
return false;
}
if (id != nullptr) {
*id = job.id;
}
jobs.pop_front();
return true;
}
private:
struct Job {
U64 id;
std::unique_ptr<Pipeline>& worker;
};
U64 head = 0;
std::deque<Job> jobs;
std::vector<std::unique_ptr<Pipeline>> workers;
};
} // namespace Blade
#endif
https://github.com/luigifcruz/blade/blob/aef4bf5230ab64af7b3e22de60e71f4bb52d6aae/src/pipeline.cc
#define BL_LOG_DOMAIN "PIPELINE"
#include "blade/pipeline.hh"
namespace Blade {
Pipeline::Pipeline(const U64& numberOfAccumulationSteps,
const U64& numberOfComputeSteps)
: state(State::IDLE),
numberOfAccumulationSteps(numberOfAccumulationSteps),
numberOfComputeSteps(numberOfComputeSteps),
accumulationStepCounter(0),
computeStepCounter(0),
currentComputeCount(0) {
BL_INFO("Pipeline with {} accumulation and {} compute steps.",
numberOfAccumulationSteps, numberOfComputeSteps);
BL_CUDA_CHECK_THROW(cudaStreamCreateWithFlags(&this->stream,
cudaStreamNonBlocking), [&]{
BL_FATAL("Failed to create stream for CUDA steam: {}", err);
});
}
Pipeline::~Pipeline() {
this->synchronize();
if (this->state == State::GRAPH) {
cudaGraphDestroy(this->graph);
}
cudaStreamDestroy(this->stream);
}
const U64 Pipeline::incrementAccumulatorStep() {
return ++accumulationStepCounter;
}
const U64 Pipeline::resetAccumulatorSteps() {
const auto& previous = accumulationStepCounter;
accumulationStepCounter = 0;
return previous;
}
const U64 Pipeline::incrementComputeStep() {
return ++computeStepCounter;
}
const U64 Pipeline::resetComputeSteps() {
const auto& previous = computeStepCounter;
computeStepCounter = 0;
return previous;
}
const Result Pipeline::synchronize() {
BL_CUDA_CHECK(cudaStreamSynchronize(this->stream), [&]{
BL_FATAL("Failed to synchronize stream: {}", err);
});
return Result::SUCCESS;
}
bool Pipeline::isSynchronized() {
return cudaStreamQuery(this->stream) == cudaSuccess;
}
const Result Pipeline::compute() {
for (auto& module : this->modules) {
BL_CHECK(module->preprocess(this->stream, this->currentComputeCount));
}
switch (state) {
case State::GRAPH:
BL_CUDA_CHECK(cudaGraphLaunch(this->instance, this->stream), [&]{
BL_FATAL("Failed launch CUDA graph: {}", err);
});
break;
case State::CACHED:
BL_DEBUG("Creating CUDA Graph.");
BL_CUDA_CHECK(cudaStreamBeginCapture(this->stream,
cudaStreamCaptureModeGlobal), [&]{
BL_FATAL("Failed to begin the capture of CUDA Graph: {}", err);
});
for (auto& module : this->modules) {
BL_CHECK(module->process(this->stream));
}
BL_CUDA_CHECK(cudaStreamEndCapture(this->stream, &this->graph), [&]{
BL_FATAL("Failed to end the capture of CUDA Graph: {}", err);
});
BL_CUDA_CHECK(cudaGraphInstantiate(&this->instance, this->graph,
NULL, NULL, 0), [&]{
BL_FATAL("Failed to instantiate CUDA Graph: {}", err);
});
BL_CUDA_CHECK(cudaGraphLaunch(this->instance, this->stream), [&]{
BL_FATAL("Failed launch CUDA graph: {}", err);
});
this->state = State::GRAPH;
break;
case State::IDLE:
BL_DEBUG("Caching kernels ahead of CUDA Graph instantiation.");
for (auto& module : this->modules) {
BL_CHECK(module->process(this->stream));
}
this->state = State::CACHED;
break;
default:
BL_FATAL("Internal error.");
return Result::ERROR;
}
BL_CUDA_CHECK_KERNEL([&]{
BL_FATAL("Failed to process: {}", err);
return Result::CUDA_ERROR;
});
this->currentComputeCount += 1;
return Result::SUCCESS;
}
} // namespace Blade
https://github.com/luigifcruz/blade/blob/aef4bf5230ab64af7b3e22de60e71f4bb52d6aae/src/modules/beamformer/beamformer.cu
#include "cuComplex.h"
#include <stdint.h>
__device__ cuFloatComplex detect(const cuFloatComplex a) {
return make_cuFloatComplex((a.x * a.x) + (a.y * a.y), 0.0f);
}
template<uint64_t NBEAMS, uint64_t NANTS, uint64_t NCHANS,
uint64_t NTIME, uint64_t NPOLS, uint64_t TBLOCK,
bool EnableIncoherentBeam, bool EnableIncoherentBeamSqrt>
__global__ void ATA(const cuFloatComplex* input,
const cuFloatComplex* phasor,
cuFloatComplex* out) {
int bi = threadIdx.x;
int ti = bi + (blockIdx.y * TBLOCK);
int ch = blockIdx.x;
// Load the phasors to shared memory.
__shared__ cuFloatComplex phr_cache[NBEAMS][NANTS][NPOLS];
int iy = (ch * NPOLS) + (bi * NPOLS * NCHANS * NANTS);
const int dy = NPOLS * NCHANS;
if (bi < NBEAMS) {
for (int a = 0; a < NANTS; a++, iy += dy) {
phr_cache[bi][a][0] = phasor[iy+0];
phr_cache[bi][a][1] = phasor[iy+1];
}
}
__syncthreads();
// Load the antenna values to registers.
cuFloatComplex ant_cache[NANTS][NPOLS];
int ix = (ch * NTIME * NPOLS) + (ti * NPOLS);
const int dx = NTIME * NCHANS * NPOLS;
for (int a = 0; a < NANTS; a++, ix += dx) {
ant_cache[a][0] = input[ix+0];
ant_cache[a][1] = input[ix+1];
}
// Multiply and accumulate.
int iz = (ch * NTIME) + ti;
const int dz = NTIME * NCHANS;
for (int b = 0; b < NBEAMS; b++, iz += dz) {
cuFloatComplex acc[NPOLS] = {{0.0, 0.0}};
for (int a = 0; a < NANTS; a++) {
acc[0] = cuCaddf(acc[0], cuCmulf(ant_cache[a][0], phr_cache[b][a][0]));
acc[1] = cuCaddf(acc[1], cuCmulf(ant_cache[a][1], phr_cache[b][a][1]));
}
reinterpret_cast<float4*>(out)[iz] = *reinterpret_cast<float4*>(acc);
}
if (EnableIncoherentBeam) {
cuFloatComplex acc[NPOLS] = {{0.0, 0.0}};
for (int a = 0; a < NANTS; a++) {
acc[0] = cuCaddf(acc[0], detect(cuCmulf(ant_cache[a][0], phr_cache[0][a][0])));
acc[1] = cuCaddf(acc[1], detect(cuCmulf(ant_cache[a][1], phr_cache[0][a][1])));
}
if (EnableIncoherentBeamSqrt) {
acc[0] = make_cuFloatComplex(sqrt(acc[0].x), acc[0].y);
acc[1] = make_cuFloatComplex(sqrt(acc[1].x), acc[1].y);
}
reinterpret_cast<float4*>(out)[iz] = *reinterpret_cast<float4*>(acc);
}
}
template<uint64_t NBEAMS, uint64_t NANTS, uint64_t NCHANS,
uint64_t NTIME, uint64_t NPOLS, uint64_t TBLOCK,
bool EnableIncoherentBeam, bool EnableIncoherentBeamSqrt>
__global__ void MeerKAT(const cuFloatComplex* input,
const cuFloatComplex* phasor,
cuFloatComplex* out) {
int bi = threadIdx.x;
int ti = bi + (blockIdx.y * TBLOCK);
int ch = blockIdx.x;
// Load the antenna values to registers.
cuFloatComplex ant_cache[NANTS][NPOLS];
int ix = (ch * NTIME * NPOLS) + (ti * NPOLS);
const int dx = NTIME * NCHANS * NPOLS;
for (int a = 0; a < NANTS; a++, ix += dx) {
ant_cache[a][0] = input[ix+0];
ant_cache[a][1] = input[ix+1];
}
// Multiply and accumulate.
int iy = 0;
int iz = (ch * NTIME) + ti;
const int dz = NTIME * NCHANS;
for (int b = 0; b < NBEAMS; b++, iz += dz) {
cuFloatComplex acc[NPOLS] = {{0.0, 0.0}};
for (int a = 0, x = ix; a < NANTS; a++, iy += 1, x += dx) {
acc[0] = cuCaddf(acc[0], cuCmulf(ant_cache[a][0], phasor[iy]));
acc[1] = cuCaddf(acc[1], cuCmulf(ant_cache[a][1], phasor[iy]));
}
reinterpret_cast<float4*>(out)[iz] = *reinterpret_cast<float4*>(acc);
}
}
.jpg)
.jpg)
.jpg)
